diff --git a/cuda_core/cuda/core/_program.pxd b/cuda_core/cuda/core/_program.pxd index 7c82119bcbd..e1cbaa6fbda 100644 --- a/cuda_core/cuda/core/_program.pxd +++ b/cuda_core/cuda/core/_program.pxd @@ -21,3 +21,4 @@ cdef class Program: str _code_type # Normalised code_type ("c++", "ptx", "nvvm") str _pch_status # PCH creation outcome after compile bytes _nvrtc_name # Source filepath given to NVRTC; a real path for debug builds + list _extra_options # NVRTC options Program adds on top of ProgramOptions diff --git a/cuda_core/cuda/core/_program.pyx b/cuda_core/cuda/core/_program.pyx index 2d946730233..b937ded7548 100644 --- a/cuda_core/cuda/core/_program.pyx +++ b/cuda_core/cuda/core/_program.pyx @@ -103,8 +103,12 @@ cdef class Program: self._cleanup_debug_source() def _cleanup_debug_source(self): - path = self._nvrtc_name.decode() - self._unlink_debug_source(path) + # Only a temp file this Program wrote may be removed, and the name having + # moved off options.name is what says one was written. Without that test + # the caller's own file is deleted whenever options.name happens to match + # something on disk, since the unredirected name is just that path. + if self._nvrtc_name is not None and self._nvrtc_name != self._options._name: + self._unlink_debug_source(self._nvrtc_name.decode()) def _unlink_debug_source(self, path: str) -> None: try: @@ -852,6 +856,7 @@ cdef inline int Program_init(Program self, object code, str code_type, object op self._pch_status = None self._nvrtc_name = options._name + self._extra_options = [] if code_type == "c++": assert_type(code, str) @@ -862,6 +867,16 @@ cdef inline int Program_init(Program self, object code, str code_type, object op debug_path = self._try_materialize_nvrtc_debug_source(code) if debug_path is not None: self._nvrtc_name = debug_path.encode() + # NVRTC resolves #include "..." against the directory of the name it + # is given, so moving the name into the temp dir would otherwise stop + # every quoted include from resolving where it did before. + try: + include_dir = os.path.dirname(os.path.abspath(options.name)) + except OSError: + # abspath needs a cwd; with none there is no directory to restore. + pass + else: + self._extra_options = [b"--include-path=" + include_dir.encode()] # TODO: support pre-loaded headers & include names code_bytes = code.encode() @@ -1035,7 +1050,7 @@ cdef object _read_pch_status(cynvrtc.nvrtcProgram prog): cdef object Program_compile_nvrtc(Program self, str target_type, object name_expressions, object logs): """Compile using NVRTC backend and return ObjectCode.""" cdef cynvrtc.nvrtcProgram prog = as_cu(self._h_nvrtc) - cdef list options_list = self._options.as_bytes("nvrtc", target_type) + cdef list options_list = self._options.as_bytes("nvrtc", target_type) + self._extra_options result = _nvrtc_compile_and_extract( prog, target_type, name_expressions, logs, options_list, self._nvrtc_name.decode(), diff --git a/cuda_core/tests/test_program.py b/cuda_core/tests/test_program.py index 8c6dfae8467..4874c6184c7 100644 --- a/cuda_core/tests/test_program.py +++ b/cuda_core/tests/test_program.py @@ -1090,6 +1090,62 @@ def _compile_one(): assert not os.path.isfile(name) +@pytest.mark.agent_authored(model="claude-opus-5") +def test_nvrtc_debug_preserves_quoted_include_resolution(init_cuda, tmp_path, monkeypatch): + """A quoted #include keeps resolving once debug redirects the NVRTC name (issue #2422). + + NVRTC looks for #include "..." in the directory of the name it was handed, so + pointing that name at a temp .cu moves the search away from where the header + lives and turning debug on alone breaks a compile that worked without it. + """ + import os + + (tmp_path / "local.h").write_text("#define BUMP 7\n", encoding="utf-8") + monkeypatch.chdir(tmp_path) + code = '#include "local.h"\nextern "C" __global__ void matmul(int* out) { *out = BUMP; }\n' + + for debug in (False, True): + prog = Program(code, "c++", ProgramOptions(arch="sm_80", debug=debug)) + try: + name = prog.compile("ptx").name + finally: + prog.close() + if debug: + # Only a regression test while the name really does move out of the + # directory holding local.h; otherwise it would pass for free. + assert os.path.dirname(os.path.realpath(name)) != os.path.realpath(tmp_path) + + +@pytest.mark.agent_authored(model="claude-opus-5") +@pytest.mark.parametrize("debug", [False, True]) +def test_nvrtc_debug_keeps_file_the_caller_named(init_cuda, tmp_path, debug): + """Program only unlinks a temp file it wrote itself (issue #2422). + + The name handed to NVRTC doubled as the cleanup target, so a name pointing at + a file that already existed made teardown delete the caller's own source. + """ + import gc + + source = tmp_path / "matmul.cu" + contents = "// the caller's own file\n" + source.write_text(contents, encoding="utf-8") + code = 'extern "C" __global__ void matmul() {}' + options = ProgramOptions(arch="sm_80", name=str(source), debug=debug) + + prog = Program(code, "c++", options) + prog.compile("ptx") + prog.close() + assert source.is_file(), "close() deleted a file the caller owns" + + # __dealloc__ runs the same cleanup, so collection must spare it too. + prog = Program(code, "c++", options) + prog.compile("ptx") + del prog + gc.collect() + assert source.is_file(), "collection deleted a file the caller owns" + assert source.read_text(encoding="utf-8") == contents + + @pytest.mark.agent_authored(model="cursor-grok-4.6") def test_cuda_gdb_shows_nvrtc_debug_source_lines(init_cuda): import pathlib