From 67e14da511d74e4a6fcf29200f5e772f8e51249a Mon Sep 17 00:00:00 2001 From: "Sun, Peng" Date: Thu, 30 Mar 2017 15:03:09 -0500 Subject: [PATCH] update hip_faq.md on using GGL as default Change-Id: I6ce1112eedeac3b377fe55ad8445f3c465c2eed4 --- hipamd/docs/markdown/hip_faq.md | 13 ++++++++----- 1 file changed, 8 insertions(+), 5 deletions(-) diff --git a/hipamd/docs/markdown/hip_faq.md b/hipamd/docs/markdown/hip_faq.md index 5ad8d9e8a9..ee7e5dd347 100644 --- a/hipamd/docs/markdown/hip_faq.md +++ b/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) @@ -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. \ No newline at end of file