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

* first version after cleaning unnecessary code

* update _program.pyi

* refactor(cuda.core): move NVRTC debug source helpers onto Program

Keep materialize/unlink next to the Program lifetime that owns the temp file,
instead of as module-level functions.

* test(cuda.core): verify cuda-gdb can list materialized NVRTC debug source

Skip when cuda-gdb is missing, and use shorter {caller}_{kernel}_ temp names.

* test(cuda.core): cover NVRTC debug fallback when temp is unwritable

* test(cuda.core): cover concurrent NVRTC debug temp-file uniqueness

Author: lijinf2

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,18 @@ 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}_{kernel}_XXXXXXXX.cu`` temp file for cuda-gdb.
+
+        ``caller`` is the Python file stem (no ``.py``). ``kernel`` is the first
+        ``__global__`` function name, or ``kernel`` if none is found.
+
+        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,49 @@ 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}_{kernel}_XXXXXXXX.cu`` temp file for cuda-gdb.
+
+        ``caller`` is the Python file stem (no ``.py``). ``kernel`` is the first
+        ``__global__`` function name, or ``kernel`` if none is found.
+
+        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 = "program"
+        if frame:
+            caller = os.path.splitext(os.path.basename(frame.f_code.co_filename))[0]
+        match = re.search(r"__global__.*?(\w+)\s*\(", code, re.DOTALL)
+        prefix = re.sub(r"\W", "_", f"{caller}_{match.group(1) if match 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 +270,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 +835,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 +1022,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 +1050,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/helpers/cuda_gdb_src.py

@@ -0,0 +1,27 @@
+# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: Apache-2.0
+"""Inferior for cuda-gdb: compile with debug=True and launch, keep Program alive."""
+
+from cuda.core import Device, LaunchConfig, Program, ProgramOptions, launch
+
+CODE = """
+extern "C" __global__ void kernel() {
+    int x = 0;
+    x += 1;  // ISSUE_2422_SOURCE_LINE
+}
+"""
+
+
+def main() -> None:
+    dev = Device()
+    dev.set_current()
+    stream = dev.create_stream()
+    prog = Program(CODE, "c++", ProgramOptions(debug=True, arch=f"sm_{dev.arch}"))
+    mod = prog.compile("cubin")
+    k = mod.get_kernel("kernel")
+    launch(stream, LaunchConfig(grid=1, block=1), k)
+    stream.sync()
+
+
+if __name__ == "__main__":
+    main()

cuda_core/tests/test_program.py

@@ -4,6 +4,9 @@
 
 import contextlib
 import re
+import shutil
+import subprocess
+import sys
 import warnings
 
 import pytest
@@ -989,6 +992,131 @@ 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_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)
+
+
+@pytest.mark.agent_authored(model="cursor-grok-4.6")
+@pytest.mark.thread_unsafe(reason="monkeypatches tempfile.mkstemp on the Program module")
+def test_nvrtc_debug_falls_back_when_tmp_not_writable(init_cuda, monkeypatch):
+    """debug=True still compiles if the temp dir cannot be written (issue #2422)."""
+    from cuda.core import _program
+
+    def _denied(*_args, **_kwargs):
+        raise OSError(30, "Read-only file system")
+
+    monkeypatch.setattr(_program.tempfile, "mkstemp", _denied)
+
+    code = 'extern "C" __global__ void matmul() {}'
+    prog = Program(code, "c++", ProgramOptions(debug=True, arch="sm_80"))
+    try:
+        assert prog.compile("ptx").name == "default_program"
+    finally:
+        prog.close()
+
+
+@pytest.mark.agent_authored(model="cursor-grok-4.6")
+def test_nvrtc_debug_concurrent_compile_uses_unique_temp_files(init_cuda):
+    """Same kernel compiled concurrently gets distinct mkstemp paths (issue #2422)."""
+    import os
+    import threading
+    from concurrent.futures import ThreadPoolExecutor
+
+    code = 'extern "C" __global__ void matmul() {}'
+    n = 2
+    barrier = threading.Barrier(n)
+
+    def _compile_one():
+        Device().set_current()
+        barrier.wait()
+        prog = Program(code, "c++", ProgramOptions(debug=True, arch="sm_80"))
+        name = prog.compile("ptx").name
+        return prog, name
+
+    with ThreadPoolExecutor(max_workers=n) as pool:
+        futures = [pool.submit(_compile_one) for _ in range(n)]
+        results = [fut.result() for fut in futures]
+
+    progs, names = zip(*results)
+    try:
+        assert len(set(names)) == n
+        for name in names:
+            assert os.path.isfile(name)
+            assert re.fullmatch(r"test_program_matmul_[a-z0-9_]{8}\.cu", os.path.basename(name))
+            with open(name, encoding="utf-8") as fh:
+                assert fh.read() == code
+    finally:
+        for prog in progs:
+            prog.close()
+    for name in names:
+        assert not os.path.isfile(name)
+
+
+@pytest.mark.agent_authored(model="cursor-grok-4.6")
+def test_cuda_gdb_shows_nvrtc_debug_source_lines(init_cuda):
+    import pathlib
+
+    cuda_gdb = shutil.which("cuda-gdb")
+    if cuda_gdb is None:
+        pytest.skip("cuda-gdb is not on PATH")
+
+    child = pathlib.Path(__file__).resolve().parent / "helpers" / "cuda_gdb_src.py"
+    proc = subprocess.run(  # noqa: S603 - trusted argv: cuda-gdb + this interpreter + in-tree helper
+        [
+            cuda_gdb,
+            "--batch",
+            "-ex",
+            "set cuda break_on_launch application",
+            "-ex",
+            "run",
+            "-ex",
+            "list",
+            "--args",
+            sys.executable,
+            "-u",
+            str(child),
+        ],
+        check=False,
+        capture_output=True,
+        text=True,
+        timeout=120,
+    )
+    output = (proc.stdout or "") + (proc.stderr or "")
+    if "Operation not permitted" in output or "ptrace" in output.lower():
+        pytest.skip("cuda-gdb cannot trace this process")
+    assert re.search(r"cuda_gdb_src_kernel_\w+\.cu", output), output
+    assert "ISSUE_2422_SOURCE_LINE" in output, output
+    assert "No such file or directory" not in output, output
+
+
 def test_nvrtc_compile_with_logs_capture(init_cuda):
     """Program.compile with logs= exercises the NVRTC program-log reading path."""
     import io