diff --git a/projects/clr/hipamd/docs/markdown/hip_bugs.md b/projects/clr/hipamd/docs/markdown/hip_bugs.md index 3274ed5b5c..e15c37fc54 100644 --- a/projects/clr/hipamd/docs/markdown/hip_bugs.md +++ b/projects/clr/hipamd/docs/markdown/hip_bugs.md @@ -4,6 +4,7 @@ - [Errors related to undefined reference to `__hcLaunchKernel__***__grid_launch_parm**](#errors-related-to-undefined-reference-to-hclaunchkernel__grid_launch_parm) - [Application hangs after a hipLaunchKernel call](#what-if-i-see-application-hangs-after-a-hiplaunchkernel-call) +- [What is the current limitation of HIP Generic Grid Launch method?](#what-is-the-current-limitation-of-hip-generic-grid-launch-method) @@ -23,8 +24,6 @@ namespace { __global__ MyKernel … - Avoid calling member functions - -### What if I see application hangs after a hipLaunchKernel call? If hipLaunchKernel takes parameters that request explicitly memcpy, then it will cause application hang. Reason is that the hipLaunchKernel macro locks the stream. If kernel paramters are actually function calls which invoke other hip apis (i.e. memcpy) to the same stream, then deadlock occurs. @@ -43,4 +42,8 @@ auto bot_gpu_data = bottom[0]->gpu_data(); hipLaunchKernel( LRNComputeDiff, dim3(CAFFE_GET_BLOCKS(n_threads)), dim3(CAFFE_HIP_NUM_THREADS), 0, 0, n_threads, bot_gpu_data); -``` \ No newline at end of file +``` + +### 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 diff --git a/projects/clr/hipamd/docs/markdown/hip_faq.md b/projects/clr/hipamd/docs/markdown/hip_faq.md index ee7e5dd347..8ccb458103 100644 --- a/projects/clr/hipamd/docs/markdown/hip_faq.md +++ b/projects/clr/hipamd/docs/markdown/hip_faq.md @@ -27,7 +27,6 @@ * [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) @@ -243,7 +242,3 @@ To disable it and use the legancy grid launch method, please either change the d $HIP/include/hip/hcc_detail/hip_runtime_api.h $HIP/include/hip/hcc_detail/host_defines.h Or pass "-DGENERIC_GRID_LAUNCH=0" to hipcc at application compilation time. - -### 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