Trivial GPU-printf hello causes hang on gfx803

[BryantLam]

I'm running a trivial hello program on gfx803 with a printf statement inside the GPU kernel. The host code appears to get stuck in hipDeviceSynchronize and the GPU kernel never prints (the program hangs). rocm-smi reports 100% GPU.

The expected behavior / no problem occurs on gfx900: the GPU prints 8 lines and the code returns.

// hello.cpp
#include <hip/hip_runtime.h>
#include <stdio.h>

__global__ void hello(void)
{
  int id = threadIdx.x;
  printf("Hello world from GPU thread %d!\n", id);
  return;
}

int main(void)
{
  hipLaunchKernelGGL(hello, dim3(1), dim3(8), 0, 0);
  hipDeviceSynchronize();
  return 0;
}

$  /opt/rocm/bin/hipcc  hello.cpp  -o hello

Environment

• ROCm 3.9.0
• HIP version: 3.9.20412-6d111f85
• gfx803: AMD Radeon Fury Nano
  • repro'd on three different devices

[francis0407]

Hi, I have the same problem on MI50 (gfx906). Have you solved this?

[BryantLam]

I am still observing the same behavior with ROCm 3.9.0 and 4.2. Hang on gfx803; okay on gfx900.

[lahwaacz]

Still hangs on gfx803 with ROCm 5.6...

Interestingly, only /opt/rocm/bin/hipcc hello.cpp hangs. /opt/rocm/llvm/bin/clang++ -x hip -l amdhip64 -L /opt/rocm/lib hello.cpp does not hang, but nothing is printed to stdout.

[bdenhollander]

ROCm 5.7.0 includes a fallback to OpenCL-based printf that may resolve this issue. https://rocm.docs.amd.com/en/docs-5.7.0/release.html#non-hostcall-hip-printf

[ppanchad-amd]

@BryantLam, Sorry for the lack of response. Please try latest ROCm 6.0.2 (HIP 6.0.32831) to see if your issue still exists? If resolved, please close the ticket. Thanks.

[lahwaacz]

@ppanchad-amd Still hangs on gfx803 with HIP 6.0.32831.

[ppanchad-amd]

@lahwaacz Unfortunately, gfx803 had been dropped official support from ROCm-4.0. Closing bug. Thanks!