From f2bfc8a405d8ff49dff0063a203f2a02e4c7623c Mon Sep 17 00:00:00 2001 From: pensun Date: Wed, 8 Feb 2017 16:19:28 -0600 Subject: [PATCH] Add pseudo code example for hip_bugs.md Change-Id: Ia2af8e6165faeb3fbb81428e20d4dc5b19b2fa9e --- hipamd/docs/markdown/hip_bugs.md | 15 ++++++++++++++- 1 file changed, 14 insertions(+), 1 deletion(-) diff --git a/hipamd/docs/markdown/hip_bugs.md b/hipamd/docs/markdown/hip_bugs.md index 834cd9f8ce..3274ed5b5c 100644 --- a/hipamd/docs/markdown/hip_bugs.md +++ b/hipamd/docs/markdown/hip_bugs.md @@ -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. \ No newline at end of file +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); + +``` \ No newline at end of file