diff --git a/docs/markdown/hip_bugs.md b/docs/markdown/hip_bugs.md index 0bece74974..834cd9f8ce 100644 --- a/docs/markdown/hip_bugs.md +++ b/docs/markdown/hip_bugs.md @@ -2,7 +2,8 @@ -- [Errors related to undefined reference to `__hcLaunchKernel__***__grid_launch_parm**](#error-undefined-reference) +- [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) @@ -22,3 +23,11 @@ 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. + +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