Dateien
rocm-systems/docs/markdown/hip_bugs.md
T
pensun 95677edabb Add more hip_bug.md entry, regarding hang after hipLaunchKernel
Change-Id: I5800cb627179ec0e913cd36d332fb8c2994ab71e
2017-01-24 22:43:25 -06:00

1.4 KiB

HIP Bugs

Some common code practices may lead to hipcc generating a error with the form : undefined reference to `_hcLaunchKernel__ZN15vecAddNamespace6vecAddIidEEv16grid_launch_parmPT0_S3_S3_T

To workaround, try:

  • Avoid calling hcLaunchKernel from a function with the host attribute host MyFunc(…) { hipLaunchKernel(myKernel, …)
  • Avoid use of static with kernel definition: static global MyKernel
  • Avoid defining kernels in anonymous namespace 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.