Move current GGL limitations to hip_bugs.md

Change-Id: I77d0eae0a67eccef7dd2bea0f402736642c96554


[ROCm/clr commit: 5826e53644]
Esse commit está contido em:
Sun, Peng
2017-04-01 08:21:06 -05:00
commit a4901cdcd3
2 arquivos alterados com 6 adições e 8 exclusões
+6 -3
Ver Arquivo
@@ -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)
<!-- tocstop -->
@@ -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);
```
```
### 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.
@@ -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)
<!-- tocstop -->
@@ -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.