Add pseudo code example for hip_bugs.md

Change-Id: Ia2af8e6165faeb3fbb81428e20d4dc5b19b2fa9e


[ROCm/clr commit: f2bfc8a405]
Этот коммит содержится в:
pensun
2017-02-08 16:19:28 -06:00
родитель 76a58521f5
Коммит b411549397
+14 -1
Просмотреть файл
@@ -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);
```