Add pseudo code example for hip_bugs.md
Change-Id: Ia2af8e6165faeb3fbb81428e20d4dc5b19b2fa9e
This commit is contained in:
@@ -30,4 +30,17 @@ 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.
|
||||
|
||||
To workaround, try:
|
||||
Move the function calls so they occur outside the hipLaunchKernel macro, store results in temps, then use the tems inside the kernel.
|
||||
Move the function calls so they occur outside the hipLaunchKernel macro, store results in temps, then use the tems inside the kernel.
|
||||
|
||||
```
|
||||
// Example pseudo code causing system hang:
|
||||
// "bottom[0]->gpu_data()" calls hipMemcpy() implicitly and using the same stream, cause deadlock condition.
|
||||
hipLaunchKernel(HIP_KERNEL_NAME(LRNComputeDiff),dim3(CAFFE_GET_BLOCKS(n_threads)), dim3(CAFFE_HIP_NUM_THREADS), 0, 0, n_threads,
|
||||
bottom[0]->gpu_data());
|
||||
|
||||
// Move "gpu_data()" ouside of hipLaunchKernel to avoid hang.
|
||||
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);
|
||||
|
||||
```
|
||||
Reference in New Issue
Block a user