ROCgdb unable to properly debug OpenCL GPU kernels on MI300X (ROCm 6.2.2) — breakpoint hits but stepping fails

[AshayaRamteke98]

I am trying to debug an OpenCL kernel running on an AMD GPU using roc-gdb and want to understand whether this workflow is officially supported or if I am missing some required debug setup.

Environment
GPU: AMD Instinct MI300X (gfx942)
CPU: AMD EPYC 9534
ROCm version: 6.2.2
Debugger: roc-gdb
OpenCL runtime: ROCm OpenCL
OS: Linux

Test setup :-
I have a simple OpenCL vector addition example:
Kernel:

__kernel void vecadd(__global const float *a,
__global const float *b,
__global float *c)
{
int gid = get_global_id(0);
c[gid] = a[gid] + b[gid];
}

Host code:

• creates OpenCL context
• selects GPU device
• creates queue with:
  clCreateCommandQueueWithProperties(...)
  (not deprecated clCreateCommandQueue())

builds program with:
clBuildProgram(program,
1,
&device,
"-g -cl-opt-disable",
NULL,
NULL);

launches kernel with:
clEnqueueNDRangeKernel(...)

What works
Host-side debugging works fine:

break main
break clBuildProgram
break clEnqueueNDRangeKernel

Kernel execution works correctly.

I can also set a breakpoint on the kernel symbol:

break vecadd

and roc-gdb does stop at kernel entry.

So kernel dispatch is definitely happening.

Problem

When I try to step inside the OpenCL GPU kernel:

step

roc-gdb fails with:
fatal error:
Per-queue memory reserved for the debugger is missing.
amd_dbgapi_displaced_stepping_start failed

Additional investigation

I extracted the generated OpenCL program binary using:
clGetProgramInfo(program, CL_PROGRAM_BINARIES, ...)
and inspected it:
llvm-readelf -S kernel.bin

The binary contains:

.text
.symtab
.dynsym

but no debug sections, e.g. no:

.debug_info
.debug_line
.debug_abbrev
.debug_str

So although I pass:

-g -cl-opt-disable

to clBuildProgram(), the generated GPU kernel binary does not appear to contain DWARF debug information.

Questions
-Is source-level OpenCL GPU kernel debugging with roc-gdb officially supported on ROCm 6.2.2?
-Does ROCm OpenCL actually support generating debug info for kernels via:
clBuildProgram(..., "-g")
or is this ignored / unsupported?

-What causes:
Per-queue memory reserved for the debugger is missing
amd_dbgapi_displaced_stepping_start failed
Is this an OpenCL runtime limitation?

-Is the OpenCL runtime expected to create debugger-compatible GPU queues automatically when launched under roc-gdb?
-Is roc-gdb intended primarily for HIP kernels, with OpenCL support being partial / limited?

[lancesix]

Hi, thanks for reaching out.

First things first, ROCgdb is primarily supporting HIP kernels, OpenCL is not fully supported as described in the documentation:

• https://github.com/ROCm/ROCgdb/blob/amd-staging/gdb/doc/gdb.texinfo#L29220-L29224
• https://github.com/ROCm/ROCgdb/blob/amd-staging/gdb/doc/gdb.texinfo#L30486-L30488

As of now, only assembly debugging is expected to fully work.

-Is source-level OpenCL GPU kernel debugging with roc-gdb officially supported on ROCm 6.2.2?

Unfortunately no. That being said, even if source level is not fully supported, there are pieces available. See below.

-Is roc-gdb intended primarily for HIP kernels, with OpenCL support being partial / limited?

Yes, the primarily supported framework is HIP.

-Does ROCm OpenCL actually support generating debug info for kernels via clBuildProgram(..., "-g") or is this ignored / unsupported?

As far as I can see, clBuildProgram's -g flag in the options parameter is not the way to produce debug info. You can pass this as an environment variable with AMD_OCL_BUILD_OPTIONS_APPEND="-O0 -g", which should produce debug info.

Note that due to the nature of online compilation of OpenCL, there is no source file the debugger can find to present to the user. It is possible to ask the OpenCL runtime to keep temporary files around so the debugger can find them: AMD_COMGR_SAVE_TEMPS=1 AMD_COMGR_SAVE_LLVM_TEMPS=1.

With those 3 environment variables, you should have some level of debug-ability*.

For the "Per-queue memory reserved for the debugger is missing" failure, this looks like a runtime initialisation issue which does not reserve the necessary buffers. If the issue is observed after a fresh install / reboot, you can try to use set displaced-stepping off in ROCgdb. This will disable ROCgdb's feature which relies on this runtime provided buffer. But really, the low level runtime should allocate this buffer for all queues, this should not be opt-in.

I hope those elements help you, please let me know otherwise.

*) I did not test those environment variables on the 6.2.2 release, only on current development tree. I however expect their effect should not have changed and should be applicable to 6.2 as well.

[AshayaRamteke98]

Thank you for the detailed explanation and suggestions.

After setting:

export AMD_OCL_BUILD_OPTIONS_APPEND="-O0 -g"
export AMD_COMGR_SAVE_TEMPS=1
export AMD_COMGR_SAVE_LLVM_TEMPS=1

and using:
set displaced-stepping off

I am now able to debug OpenCL kernels to some extent.

I was also able to inspect some COMGR-generated artifacts such as:

/tmp/comgr-/input/linked.bc.o
/tmp/comgr-/output/a.so

which helped me better understand the compilation flow.

I still have a few questions regarding OpenCL debugging support in ROCgdb:

1. OpenCL compilation model and pending breakpoints

When you refer to "online compilation" of OpenCL kernels, does this mean JIT compilation performed at clBuildProgram() time, or some other compilation mechanism?

I have noticed that OpenCL kernel breakpoints remain pending until the kernel is built. However, I also see similar pending breakpoint behavior with HIP kernels. Could you explain how ROCgdb discovers and resolves GPU kernel symbols, and whether pending breakpoints necessarily indicate JIT compilation?

2. Meaning of "assembly-level debugging"

You mentioned that only assembly-level debugging is expected to fully work.

Could you explain precisely what "assembly-level debugging" means in ROCgdb?

With the environment variables above, I am often able to step through source lines inside kernels, so I would like to better understand what "assembly-level debugging" means in practice.

3. Wave/thread switching during stepping

This is currently a major issue i am facing.

When multiple waves are active, ROCgdb often switches to another wave when I issue a step command.

For example:

• execution stops on wave 7
• I issue step
• ROCgdb switches to wave 8
• wave 7 appears to continue execution

Using:

set scheduler-locking on

makes stepping much more reliable.

Could you explain:

• why this wave switching occurs
• whether it is expected behavior
• When a different wave becomes active after I issue a step command, which part of the system is making that decision: ROCgdb, ROCdbgapi, the ROCm runtime, or the GPU hardware scheduler?
• whether there is a recommended long-term solution other than manually enabling scheduler-locking

4. Per-queue debugger memory

Using:

set displaced-stepping off

works around the "Per-queue memory reserved for the debugger is missing" error.

Could you explain the root cause of this issue and whether there is a permanent fix or runtime configuration that should enable the required debugger memory allocation automatically?

I would also like to understand the "per-queue debugger memory" better. Which component allocates and owns this memory (runtime, queue, ROCdbgapi, etc.), and where does it actually reside?

5. COMGR intermediate files

Using:

export AMD_COMGR_SAVE_TEMPS=1
export AMD_COMGR_SAVE_LLVM_TEMPS=1

I am able to save artifacts such as:

linked.bc.o
a.so

However, I am not consistently seeing or preserving earlier compilation stages such as:
.ll
.bc
.s
.o

My questions are:

• Are these intermediate files expected to be generated for OpenCL kernel compilation?
• If so, what is the recommended way to preserve them?
• Are there additional environment variables required to save every compilation stage?
• Is there a documented OpenCL compilation pipeline showing which COMGR artifacts correspond to each stage?

6. ROCgdb internal representation of GPU execution

How does ROCgdb internally distinguish:

• host CPU threads
• GPU waves
• GPU lanes

Does ROCgdb use PTIDs for GPU execution contexts as well, or is there a different representation used internally through ROCdbgapi?

Thank you again for the help and clarification.