cuda.core: source-level cuda-gdb debugging broken for JIT kernels — NVRTC program name defaults to "default_program"

[rparolin]

Summary

When a cuda.core.Program is compiled from an inline source string with debug=True / lineinfo=True but no name= in ProgramOptions, cuda.core passes an empty program name to nvrtcCreateProgram. NVRTC then embeds its hardcoded fallback default_program as the DWARF source filename. Because no file named default_program exists on disk, cuda-gdb cannot resolve source — every kernel frame maps to default_program:10 and source-level GPU debugging is unusable.

Reproduction

Compile an inline kernel and break in it under cuda-gdb:

program = Program(MATMUL_KERNEL, code_type="c++",
                  options=ProgramOptions(arch=arch, debug=True, lineinfo=True, device_code_optimize=True))
kernel = program.compile(target_type="cubin").get_kernel("matmul_shared")

CUDA thread hit Breakpoint 1, matmul_shared<<<...>>> ()
    at .../default_program:10
warning: 10  .../default_program: No such file or directory

Isolated to the program name: nvrtcCreateProgram with name "" → cubin embeds default_program; with "matmul.cu" → embeds matmul.cu and line mapping resolves correctly.

Current workaround (confirmed by the reporting customer and NVIDIA QA)

Persist the kernel source to a real .cu file and pass its absolute path as name=:

filename = os.path.abspath("matmul.cu")
with open(filename, "w") as f:
    f.write(MATMUL_KERNEL)
program = Program(MATMUL_KERNEL, code_type="c++",
                  options=ProgramOptions(arch=arch, debug=True, lineinfo=True,
                                         device_code_optimize=True, name=filename))

With the source on disk and name= pointing at it, cuda-gdb resolves the source and line mapping is correct.

Proposed enhancement

When debug or lineinfo is enabled and name is unset:

1. Quick win — emit a warning that source-level debugging requires a name that points at an on-disk source file, so users aren't silently left with default_program.
2. Full fix — auto-populate a deterministic program name and persist the JIT source to a cache/temp path, so cuda-gdb resolves source lines with no manual steps.

Context

• Origin: internal NVBugs 6420444 (Linaro, CUDA-GDB Python debugging). cuda-gdb ruled itself out; root cause is cuda.core / NVRTC program-name handling.
• Not a blocker today (Python GPU debugging is nascent) — this is a developer-experience improvement for future usability.

[leofang]

Does cuda-gdb actually need a file to exist on-disk? 👀 I thought it just needs the debug symbols and line info?