From 95677edabbe4e1fca92ab762b451c2207bc6ff70 Mon Sep 17 00:00:00 2001 From: pensun Date: Tue, 24 Jan 2017 22:43:25 -0600 Subject: [PATCH] Add more hip_bug.md entry, regarding hang after hipLaunchKernel Change-Id: I5800cb627179ec0e913cd36d332fb8c2994ab71e --- docs/markdown/hip_bugs.md | 11 ++++++++++- 1 file changed, 10 insertions(+), 1 deletion(-) 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