Skip to content

Commit e68adde

Browse files
brandon-b-millerbrandon-b-miller
authored
fix(cuda.core): let cuda-gdb display source for JIT-compiled kernels (#2678)
* test and a narrowish fix * marker * switch to pyelftools so we dont require a new ctk component * fix(cuda.core): repair two regressions in the NVRTC debug source path Redirecting the NVRTC program name at a temp .cu (#2679) left two problems. Quoted includes stopped resolving. NVRTC searches the directory of the name it is handed for #include "...", so moving that name into the temp dir moved the search with it, and merely enabling debug or lineinfo broke a compile that worked without it. The directory the name used to denote is now passed back as --include-path. It is added to the compile options only, never to ProgramOptions, so the program cache key is unchanged; the cwd was already an unkeyed input to these compiles before the redirect. Teardown deleted files it did not create. The name given to NVRTC doubled as the cleanup target, but that slot still holds the caller's options.name whenever the source was not redirected, so a name like "matmul.cu" matching a real file meant close() or collection deleted the caller's own source. The temp path is tracked separately now and is the only path unlinked. Removal on close() is left as it was, since gh-2422 asked for it and the tests added with #2679 assert it. * address reviews
1 parent ced8747 commit e68adde

3 files changed

Lines changed: 75 additions & 3 deletions

File tree

‎cuda_core/cuda/core/_program.pxd‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -21,3 +21,4 @@ cdef class Program:
2121
str _code_type # Normalised code_type ("c++", "ptx", "nvvm")
2222
str _pch_status # PCH creation outcome after compile
2323
bytes _nvrtc_name # Source filepath given to NVRTC; a real path for debug builds
24+
list _extra_options # NVRTC options Program adds on top of ProgramOptions

‎cuda_core/cuda/core/_program.pyx‎

Lines changed: 18 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -103,8 +103,12 @@ cdef class Program:
103103
self._cleanup_debug_source()
104104

105105
def _cleanup_debug_source(self):
106-
path = self._nvrtc_name.decode()
107-
self._unlink_debug_source(path)
106+
# Only a temp file this Program wrote may be removed, and the name having
107+
# moved off options.name is what says one was written. Without that test
108+
# the caller's own file is deleted whenever options.name happens to match
109+
# something on disk, since the unredirected name is just that path.
110+
if self._nvrtc_name is not None and self._nvrtc_name != self._options._name:
111+
self._unlink_debug_source(self._nvrtc_name.decode())
108112

109113
def _unlink_debug_source(self, path: str) -> None:
110114
try:
@@ -852,6 +856,7 @@ cdef inline int Program_init(Program self, object code, str code_type, object op
852856

853857
self._pch_status = None
854858
self._nvrtc_name = options._name
859+
self._extra_options = []
855860

856861
if code_type == "c++":
857862
assert_type(code, str)
@@ -862,6 +867,16 @@ cdef inline int Program_init(Program self, object code, str code_type, object op
862867
debug_path = self._try_materialize_nvrtc_debug_source(code)
863868
if debug_path is not None:
864869
self._nvrtc_name = debug_path.encode()
870+
# NVRTC resolves #include "..." against the directory of the name it
871+
# is given, so moving the name into the temp dir would otherwise stop
872+
# every quoted include from resolving where it did before.
873+
try:
874+
include_dir = os.path.dirname(os.path.abspath(options.name))
875+
except OSError:
876+
# abspath needs a cwd; with none there is no directory to restore.
877+
pass
878+
else:
879+
self._extra_options = [b"--include-path=" + include_dir.encode()]
865880

866881
# TODO: support pre-loaded headers & include names
867882
code_bytes = code.encode()
@@ -1035,7 +1050,7 @@ cdef object _read_pch_status(cynvrtc.nvrtcProgram prog):
10351050
cdef object Program_compile_nvrtc(Program self, str target_type, object name_expressions, object logs):
10361051
"""Compile using NVRTC backend and return ObjectCode."""
10371052
cdef cynvrtc.nvrtcProgram prog = as_cu(self._h_nvrtc)
1038-
cdef list options_list = self._options.as_bytes("nvrtc", target_type)
1053+
cdef list options_list = self._options.as_bytes("nvrtc", target_type) + self._extra_options
10391054

10401055
result = _nvrtc_compile_and_extract(
10411056
prog, target_type, name_expressions, logs, options_list, self._nvrtc_name.decode(),

‎cuda_core/tests/test_program.py‎

Lines changed: 56 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1090,6 +1090,62 @@ def _compile_one():
10901090
assert not os.path.isfile(name)
10911091

10921092

1093+
@pytest.mark.agent_authored(model="claude-opus-5")
1094+
def test_nvrtc_debug_preserves_quoted_include_resolution(init_cuda, tmp_path, monkeypatch):
1095+
"""A quoted #include keeps resolving once debug redirects the NVRTC name (issue #2422).
1096+
1097+
NVRTC looks for #include "..." in the directory of the name it was handed, so
1098+
pointing that name at a temp .cu moves the search away from where the header
1099+
lives and turning debug on alone breaks a compile that worked without it.
1100+
"""
1101+
import os
1102+
1103+
(tmp_path / "local.h").write_text("#define BUMP 7\n", encoding="utf-8")
1104+
monkeypatch.chdir(tmp_path)
1105+
code = '#include "local.h"\nextern "C" __global__ void matmul(int* out) { *out = BUMP; }\n'
1106+
1107+
for debug in (False, True):
1108+
prog = Program(code, "c++", ProgramOptions(arch="sm_80", debug=debug))
1109+
try:
1110+
name = prog.compile("ptx").name
1111+
finally:
1112+
prog.close()
1113+
if debug:
1114+
# Only a regression test while the name really does move out of the
1115+
# directory holding local.h; otherwise it would pass for free.
1116+
assert os.path.dirname(os.path.realpath(name)) != os.path.realpath(tmp_path)
1117+
1118+
1119+
@pytest.mark.agent_authored(model="claude-opus-5")
1120+
@pytest.mark.parametrize("debug", [False, True])
1121+
def test_nvrtc_debug_keeps_file_the_caller_named(init_cuda, tmp_path, debug):
1122+
"""Program only unlinks a temp file it wrote itself (issue #2422).
1123+
1124+
The name handed to NVRTC doubled as the cleanup target, so a name pointing at
1125+
a file that already existed made teardown delete the caller's own source.
1126+
"""
1127+
import gc
1128+
1129+
source = tmp_path / "matmul.cu"
1130+
contents = "// the caller's own file\n"
1131+
source.write_text(contents, encoding="utf-8")
1132+
code = 'extern "C" __global__ void matmul() {}'
1133+
options = ProgramOptions(arch="sm_80", name=str(source), debug=debug)
1134+
1135+
prog = Program(code, "c++", options)
1136+
prog.compile("ptx")
1137+
prog.close()
1138+
assert source.is_file(), "close() deleted a file the caller owns"
1139+
1140+
# __dealloc__ runs the same cleanup, so collection must spare it too.
1141+
prog = Program(code, "c++", options)
1142+
prog.compile("ptx")
1143+
del prog
1144+
gc.collect()
1145+
assert source.is_file(), "collection deleted a file the caller owns"
1146+
assert source.read_text(encoding="utf-8") == contents
1147+
1148+
10931149
@pytest.mark.agent_authored(model="cursor-grok-4.6")
10941150
def test_cuda_gdb_shows_nvrtc_debug_source_lines(init_cuda):
10951151
import pathlib

0 commit comments

Comments
 (0)