fix(cuda.core): let cuda-gdb display source for JIT-compiled kernels

cuda_core/cuda/core/_program.pxd

@@ -21,3 +21,5 @@ 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
+        str _debug_source_path  # Temp source this Program wrote, or None; the only path it may unlink
+        list _extra_options     # NVRTC options Program adds on top of ProgramOptions

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. self._nvrtc_name is
+        # still the caller's options.name whenever the source was not redirected,
+        # and a name like "matmul.cu" can match a file the caller owns.
+        if self._debug_source_path is not None:
+            self._unlink_debug_source(self._debug_source_path)
+            self._debug_source_path = None
 
     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,17 @@ 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()
+                self._debug_source_path = debug_path
+                # 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 +1051,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(),

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