fix(cuda.core): write NVRTC source to a temp file when debug or lineinfo is enabled

cuda_core/cuda/core/_program.pxd

@@ -20,3 +20,4 @@ cdef class Program:
         bytes _code             # Source code as bytes: used for key derivation and NVRTC PCH retry
         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

cuda_core/cuda/core/_program.pyi

@@ -40,6 +40,15 @@ class Program:
     def __init__(self, code: str | bytes | bytearray, code_type: SourceCodeType | str, options: ProgramOptions | None=None): ...
     def close(self) -> None:
         """Destroy this program."""
+    def __dealloc__(self): ...
+    def _cleanup_debug_source(self): ...
+    def _unlink_debug_source(self, path: str) -> None: ...
+    def _try_materialize_nvrtc_debug_source(self, code: str) -> str | None:
+        """Write *code* to a ``caller_py__kernel_XXXXXXXX.cu`` temp file for cuda-gdb.
+
+        Returns None if the filesystem is not writable, so the caller can fall back
+        to the label-only behavior instead of failing the compile.
+        """
     def compile(self, target_type: ObjectCodeFormatType | str, name_expressions: tuple[str, ...] | list[str]=(), logs: object | None=None, *, cache: ProgramCacheResource | None=None) -> ObjectCode:
         """Compile the program to the specified target type.
 

cuda_core/cuda/core/_program.pyx

@@ -10,6 +10,10 @@ This module provides :class:`Program` for compiling source code into
 from __future__ import annotations
 
 from dataclasses import dataclass
+import os
+import re
+import sys
+import tempfile
 import threading
 from typing import TYPE_CHECKING
 from warnings import warn
@@ -87,6 +91,44 @@ cdef class Program:
         # Reset handles - the C++ shared_ptr destructor handles cleanup
         self._h_nvrtc.reset()
         self._h_nvvm.reset()
+        self._cleanup_debug_source()
+
+    def __dealloc__(self):
+        self._cleanup_debug_source()
+
+    def _cleanup_debug_source(self):
+        path = self._nvrtc_name.decode()
+        self._unlink_debug_source(path)
+
+    def _unlink_debug_source(self, path: str) -> None:
+        try:
+            os.unlink(path)
+        except OSError:
+            pass
+
+    def _try_materialize_nvrtc_debug_source(self, code: str) -> str | None:
+        """Write *code* to a ``caller_py__kernel_XXXXXXXX.cu`` temp file for cuda-gdb.
+
+        Returns None if the filesystem is not writable, so the caller can fall back
+        to the label-only behavior instead of failing the compile.
+        """
+        frame = sys._getframe()
+        while frame and frame.f_globals.get("__name__", "").startswith("cuda.core"):
+            frame = frame.f_back
+        caller = os.path.basename(frame.f_code.co_filename) if frame else "program"
+        kernel = re.search(r"__global__.*?(\w+)\s*\(", code, re.DOTALL)
+        prefix = re.sub(r"\W", "_", f"{caller}__{kernel.group(1) if kernel else 'kernel'}_")
+        try:
+            fd, path = tempfile.mkstemp(prefix=prefix, suffix=".cu")
+        except OSError:
+            return None
+        try:
+            with os.fdopen(fd, "w", encoding="utf-8") as f:
+                f.write(code)
+            return path
+        except OSError:
+            self._unlink_debug_source(path)
+            return None
 
     def compile(
         self,
@@ -223,7 +265,7 @@ cdef class Program:
                     stacklevel=2,
                     category=RuntimeWarning,
                 )
-            return ObjectCode._init(hit_bytes, target_type, name=self._options.name)
+            return ObjectCode._init(hit_bytes, target_type, name=self._nvrtc_name.decode())
         compiled = _program_compile_uncached(self, target_type, name_expressions, logs)
         cache[key] = compiled
         return compiled
@@ -788,16 +830,22 @@ cdef inline int Program_init(Program self, object code, str code_type, object op
     self._libdevice_added = False
 
     self._pch_status = None
+    self._nvrtc_name = options._name
 
     if code_type == "c++":
         assert_type(code, str)
         if options.extra_sources is not None:
             raise ValueError("extra_sources is not supported by the NVRTC backend (C++ code_type)")
 
+        if (options.debug or options.lineinfo) and options.name == "default_program":
+            debug_path = self._try_materialize_nvrtc_debug_source(code)
+            if debug_path is not None:
+                self._nvrtc_name = debug_path.encode()
+
         # TODO: support pre-loaded headers & include names
         code_bytes = code.encode()
         code_ptr = <const char*>code_bytes
-        name_ptr = <const char*>options._name
+        name_ptr = <const char*>self._nvrtc_name
 
         with nogil:
             HANDLE_RETURN_NVRTC(NULL, cynvrtc.nvrtcCreateProgram(
@@ -969,7 +1017,7 @@ cdef object Program_compile_nvrtc(Program self, str target_type, object name_exp
     cdef list options_list = self._options.as_bytes("nvrtc", target_type)
 
     result = _nvrtc_compile_and_extract(
-        prog, target_type, name_expressions, logs, options_list, self._options.name,
+        prog, target_type, name_expressions, logs, options_list, self._nvrtc_name.decode(),
     )
 
     cdef bint pch_creation_possible = self._options.create_pch or self._options.pch
@@ -997,14 +1045,14 @@ cdef object Program_compile_nvrtc(Program self, str target_type, object name_exp
 
     cdef cynvrtc.nvrtcProgram retry_prog
     cdef const char* code_ptr = <const char*>self._code
-    cdef const char* name_ptr = <const char*>self._options._name
+    cdef const char* name_ptr = <const char*>self._nvrtc_name
     with nogil:
         HANDLE_RETURN_NVRTC(NULL, cynvrtc.nvrtcCreateProgram(
             &retry_prog, code_ptr, name_ptr, 0, NULL, NULL))
     self._h_nvrtc = create_nvrtc_program_handle(retry_prog)
 
     result = _nvrtc_compile_and_extract(
-        retry_prog, target_type, name_expressions, logs, options_list, self._options.name,
+        retry_prog, target_type, name_expressions, logs, options_list, self._nvrtc_name.decode(),
     )
 
     status = _read_pch_status(retry_prog)

cuda_core/tests/test_program.py

@@ -989,6 +989,38 @@ def fake_find(name):
     assert captured == ["device"]
 
 
+@pytest.mark.agent_authored(model="cursor-grok-4.6")
+def test_nvrtc_debug_materializes_source_to_temp_file(init_cuda, tmp_path):
+    """debug/lineinfo writes NVRTC source to a real path; off and explicit name= do not."""
+    import os
+
+    code = 'extern "C" __global__ void matmul() {}'
+
+    # case 1: (debug=False, lineinfo=False)
+    off = Program(code, "c++", ProgramOptions(arch="sm_80"))
+    assert off.compile("ptx").name == "default_program"
+    off.close()
+
+    # case 2: (debug=True or lineinfo=True) and explicit_name is provided
+    explicit_name = str(tmp_path / "user_kernel.cu")
+    named = Program(code, "c++", ProgramOptions(name=explicit_name, debug=True, arch="sm_80"))
+    assert named.compile("ptx").name == explicit_name
+    assert not os.path.isfile(explicit_name)
+    named.close()
+
+    # case 3: (debug=True or lineinfo=True) and explicit_name is not provided
+    default_named = Program(code, "c++", ProgramOptions(debug=True, arch="sm_80"))
+    implicit_name = default_named.compile("ptx").name
+    try:
+        assert os.path.isfile(implicit_name)
+        assert re.fullmatch(r"test_program_py__matmul_[a-z0-9_]{8}\.cu", os.path.basename(implicit_name))
+        with open(implicit_name, encoding="utf-8") as fh:
+            assert fh.read() == code
+    finally:
+        default_named.close()
+    assert not os.path.isfile(implicit_name)
+
+
 def test_nvrtc_compile_with_logs_capture(init_cuda):
     """Program.compile with logs= exercises the NVRTC program-log reading path."""
     import io