update hip_faq.md on using GGL as default

Change-Id: I6ce1112eedeac3b377fe55ad8445f3c465c2eed4

hipamd/docs/markdown/hip_faq.md

@@ -27,6 +27,7 @@
   * [Using CodeXL markers for HIP Functions](#using-codexl-markers-for-hip-functions)
   * [Using HIP_TRACE_API](#using-hip_trace_api)
 - [How do I enable HIP Generic Grid Launch option?](#how-do-i-enable-hip-generic-grid-launch-option)
+- [What is the current limitation of HIP Generic Grid Launch method?](#what-is-the-current-limitation-of-hip-generic-grid-launch-method)
 
 <!-- tocstop -->
 
@@ -236,11 +237,13 @@ Unlike CUDA, in HCC, for functions defined in the header files, the keyword of "
 Thus, if failed to define "static" keyword, you might see a lot of "symbol multiply defined!" errors at compilation.
 The workaround is to explicitly add the keyword of "static" before any functions that were defined as "__forceinline__".
 
-### How do I enable HIP Generic Grid Launch option? 
-Generic Grid Launch(GGL) provide a second choice for kernel launch.
-To enable it, either change the default value of GENERIC_GRID_LAUNCH to 1 in the following to header files and rebuild HIP:
+### How do I disable HIP Generic Grid Launch option? 
+Generic Grid Launch(GGL) is currently the default method for hip kernel launch.
+To disable it and use the legancy grid launch method, please either change the default value of GENERIC_GRID_LAUNCH to 0 in the following to header files and rebuild HIP:
 $HIP/include/hip/hcc_detail/hip_runtime_api.h 
 $HIP/include/hip/hcc_detail/host_defines.h
-Or pass "-DGENERIC_GRID_LAUNCH=1" to hipcc at application compilation time.
+Or pass "-DGENERIC_GRID_LAUNCH=0" to hipcc at application compilation time.
 
-GGL was only tested with Ubuntu16.04, assuming HCC and HIP only link against libstdc++ but not libc++.
+### What is the current limitation of HIP Generic Grid Launch method?
+1. __global__ functions cannot be marked as static or put in an unnamed namespace i.e. they cannot be given internal linkage (this would clash with __attribute__((weak)));
+2. using the macro based dispatch mechanism i.e. hipLaunchKernel* only works for functions that take no more than 20 arguments (this limit can be increased up to 126, and is temporary until we can enable C++14 mode and use variadic generic lambdas); no such limitation applies do dispatching directly through grid_launch.