From 08d729895a325be361f92aa3722ed94c10eb4737 Mon Sep 17 00:00:00 2001 From: brandon-b-miller Date: Thu, 20 Aug 2026 12:58:16 -0700 Subject: [PATCH 1/4] test and a narrowish fix --- cuda_core/cuda/core/_jit_source.py | 61 ++++++++++++++++++++++++++++++ cuda_core/cuda/core/_program.pxd | 2 + cuda_core/cuda/core/_program.pyx | 40 ++++++++++++++++++-- cuda_core/tests/test_program.py | 45 ++++++++++++++++++++++ 4 files changed, 144 insertions(+), 4 deletions(-) create mode 100644 cuda_core/cuda/core/_jit_source.py diff --git a/cuda_core/cuda/core/_jit_source.py b/cuda_core/cuda/core/_jit_source.py new file mode 100644 index 00000000000..acf0ed9ba78 --- /dev/null +++ b/cuda_core/cuda/core/_jit_source.py @@ -0,0 +1,61 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# +# SPDX-License-Identifier: Apache-2.0 + + +from __future__ import annotations + +import contextlib +import hashlib +import os +import tempfile +import threading +from pathlib import Path + +_DIGEST_CHARS = 32 + + +_lock = threading.Lock() +_source_dir: tempfile.TemporaryDirectory[str] | None = None + + +def _ensure_source_dir() -> Path: + """Return the store directory, creating it if needed. Caller holds ``_lock``.""" + global _source_dir + if _source_dir is None: + _source_dir = tempfile.TemporaryDirectory(prefix="cuda-core-jit-") + root = Path(_source_dir.name) + # Re-created rather than assumed: a /tmp reaper can delete the tree out from + # under a long-running process. + root.mkdir(parents=True, exist_ok=True) + return root + + +def source_dir() -> Path: + """The process-scoped directory holding materialized JIT source.""" + with _lock: + return _ensure_source_dir() + + +def materialize(code: bytes, suffix: str = ".cu") -> str | None: + + digest = hashlib.sha256(code).hexdigest()[:_DIGEST_CHARS] + try: + with _lock: + target = _ensure_source_dir() / f"{digest}{suffix}" + # Writers are serialized and the directory is ours alone, so an + # entry that exists is one a previous caller finished writing. + if not target.exists(): + try: + target.write_bytes(code) + except BaseException: + # Otherwise a half-written entry survives to be mistaken + # for a complete one. + with contextlib.suppress(OSError): + target.unlink() + raise + return os.fspath(target) + except OSError: + # A read only or full filesystem, or a sandbox that forbids the temp + # dir should pass + return None diff --git a/cuda_core/cuda/core/_program.pxd b/cuda_core/cuda/core/_program.pxd index cea430c3f20..aa49bb5ad5e 100644 --- a/cuda_core/cuda/core/_program.pxd +++ b/cuda_core/cuda/core/_program.pxd @@ -20,3 +20,5 @@ 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 _source_name # Name handed to NVRTC/NVVM as the DWARF source path + list _extra_options # Compiler options Program adds on top of ProgramOptions diff --git a/cuda_core/cuda/core/_program.pyx b/cuda_core/cuda/core/_program.pyx index 1d07b88bbf4..3e3144ecc3b 100644 --- a/cuda_core/cuda/core/_program.pyx +++ b/cuda_core/cuda/core/_program.pyx @@ -10,6 +10,7 @@ This module provides :class:`Program` for compiling source code into from __future__ import annotations from dataclasses import dataclass +import os import threading from typing import TYPE_CHECKING from warnings import warn @@ -30,6 +31,7 @@ from ._resource_handles cimport ( ) from cuda.bindings cimport cynvrtc, cynvvm from cuda.core._utils.cuda_utils cimport HANDLE_RETURN_NVRTC, HANDLE_RETURN_NVVM +from cuda.core import _jit_source from cuda.core._device import Device from cuda.core._linker import Linker, LinkerHandleT, LinkerOptions from cuda.core._module import ObjectCode @@ -768,6 +770,31 @@ cdef inline object _translate_program_options(object options): ) +cdef inline int _program_setup_debug_source(Program self, object options, bytes code_bytes) except -1: + """Point the NVRTC source name at an on-disk copy of the source + """ + if not (options.debug or options.lineinfo): + return 0 + + # Resolved before anything is written, so a failure here leaves the original + # name in place rather than a redirected name with broken includes. + try: + include_dir = os.path.dirname(os.path.abspath(options.name)) + except OSError: + return 0 + + source_path = _jit_source.materialize(code_bytes) + if source_path is None: + return 0 + + self._source_name = source_path.encode() + + # NVRTC searches the dir of the name it's given for quoted includes, so + # hand the original one back or every #include "foo.h" stops resolving. + self._extra_options = [b"--include-path=" + include_dir.encode()] + return 0 + + cdef inline int Program_init(Program self, object code, str code_type, object options) except -1: """Initialize a Program instance.""" cdef cynvrtc.nvrtcProgram nvrtc_prog @@ -788,6 +815,10 @@ cdef inline int Program_init(Program self, object code, str code_type, object op self._libdevice_added = False self._pch_status = None + # ProgramOptions may be shared across Programs and is never mutated here, + # so any name override or added compiler option is held per-Program. + self._source_name = options._name + self._extra_options = [] if code_type == "c++": assert_type(code, str) @@ -796,8 +827,9 @@ cdef inline int Program_init(Program self, object code, str code_type, object op # TODO: support pre-loaded headers & include names code_bytes = code.encode() + _program_setup_debug_source(self, options, code_bytes) code_ptr = code_bytes - name_ptr = options._name + name_ptr = self._source_name with nogil: HANDLE_RETURN_NVRTC(NULL, cynvrtc.nvrtcCreateProgram( @@ -829,7 +861,7 @@ cdef inline int Program_init(Program self, object code, str code_type, object op # Use self._code (strictly bytes) for the C pointer so a bytearray # input doesn't trip the `code` cast at runtime. code_ptr = self._code - name_ptr = options._name + name_ptr = self._source_name code_len = len(self._code) with nogil: @@ -966,7 +998,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._options.name, @@ -997,7 +1029,7 @@ cdef object Program_compile_nvrtc(Program self, str target_type, object name_exp cdef cynvrtc.nvrtcProgram retry_prog cdef const char* code_ptr = self._code - cdef const char* name_ptr = self._options._name + cdef const char* name_ptr = self._source_name with nogil: HANDLE_RETURN_NVRTC(NULL, cynvrtc.nvrtcCreateProgram( &retry_prog, code_ptr, name_ptr, 0, NULL, NULL)) diff --git a/cuda_core/tests/test_program.py b/cuda_core/tests/test_program.py index 72f8b8f5942..a846d69abd6 100644 --- a/cuda_core/tests/test_program.py +++ b/cuda_core/tests/test_program.py @@ -3,7 +3,10 @@ # SPDX-License-Identifier: Apache-2.0 import contextlib +import os import re +import shutil +import subprocess import warnings import pytest @@ -367,6 +370,48 @@ def test_program_options_name_accepts_none(name): assert options._name == expected.encode() +@pytest.mark.parametrize("name", [None, "my_program"]) +def test_program_string_source_debug(name, tmp_path): + # when a file doesn't exist on disk, nvrtc + # still ends up referencing a file it thinks is there + # in the dwarf table. As a WAR, we put the file there + # ourselves. This test verifies its there and that its + # contents match the passed in source code. + code = 'extern "C" __global__ void my_kernel() {}' + program = Program(code, "c++", options={"name": name, "debug": True, "lineinfo": True}) + cubin = program.compile("cubin") + + # read dwarf table + nvdisasm = shutil.which("nvdisasm") + + cubin_path = tmp_path / "program.cubin" + cubin_path.write_bytes(bytes(cubin.code)) + + result = subprocess.run( # noqa: S603 + [nvdisasm, "-g", str(cubin_path)], + capture_output=True, + text=True, + errors="replace", + ) + if result.returncode != 0: + pytest.fail(f"nvdisasm -g failed with exit code {result.returncode}\n{result.stderr}", pytrace=False) + + # -g annotates each instruction with the line-table entry that covers it: + # //## File "/abs/dir/name.cu", line 12 + # The directory comes from the line table's directory table, so this is the + # fully resolved path cuda-gdb will try to open. + paths = set(re.findall(r'//## File "([^"]+)", line \d+', result.stdout)) + assert len(paths) == 1, f"expected exactly one source file in the line table, got {sorted(paths)}" + dwarf_path = paths.pop() + + # cuda-gdb opens this path literally, so the source has to actually be + # sitting there for source-level debugging to work. + assert os.path.isfile(dwarf_path) + + with open(dwarf_path, encoding="utf-8") as source_file: + assert source_file.read().splitlines() == code.splitlines() + + # This is tested against the current device's arch def test_program_compile_valid_target_type(init_cuda): code = 'extern "C" __global__ void my_kernel() {}' From ec60dcff712b53dbeff1967dd8c8479c68d9a024 Mon Sep 17 00:00:00 2001 From: brandon-b-miller Date: Thu, 20 Aug 2026 13:25:27 -0700 Subject: [PATCH 2/4] marker --- cuda_core/tests/test_program.py | 1 + 1 file changed, 1 insertion(+) diff --git a/cuda_core/tests/test_program.py b/cuda_core/tests/test_program.py index a846d69abd6..f19b0d34da5 100644 --- a/cuda_core/tests/test_program.py +++ b/cuda_core/tests/test_program.py @@ -370,6 +370,7 @@ def test_program_options_name_accepts_none(name): assert options._name == expected.encode() +@pytest.mark.human_reviewed @pytest.mark.parametrize("name", [None, "my_program"]) def test_program_string_source_debug(name, tmp_path): # when a file doesn't exist on disk, nvrtc From 219dc35972a81a5b068b7d0c842f88ed9bb44742 Mon Sep 17 00:00:00 2001 From: brandon-b-miller Date: Fri, 21 Aug 2026 07:21:55 -0700 Subject: [PATCH 3/4] switch to pyelftools so we dont require a new ctk component --- cuda_core/pixi.toml | 1 + cuda_core/pyproject.toml | 3 ++ cuda_core/tests/test_program.py | 61 +++++++++++++++++---------------- 3 files changed, 36 insertions(+), 29 deletions(-) diff --git a/cuda_core/pixi.toml b/cuda_core/pixi.toml index b2c6a3389c3..374dcc46cd5 100644 --- a/cuda_core/pixi.toml +++ b/cuda_core/pixi.toml @@ -23,6 +23,7 @@ pytest-rerunfailures = "*" cloudpickle = "*" docutils = "*" psutil = "*" +pyelftools = "*" pyglet = "*" [feature.test.pypi-dependencies] diff --git a/cuda_core/pyproject.toml b/cuda_core/pyproject.toml index 4c2c65e9b3f..a197379e9f3 100644 --- a/cuda_core/pyproject.toml +++ b/cuda_core/pyproject.toml @@ -70,6 +70,9 @@ test = [ "cloudpickle==3.1.2", "psutil==7.2.2", "docutils==0.23", + # Reads the DWARF line table out of a compiled cubin; pure Python, so it + # works on every platform and free-threaded build in the test matrix. + "pyelftools==0.33", # TODO: remove the Python 3.15 guard once 3.15 is officially supported "cffi==2.0.0; python_version < '3.15'", ] diff --git a/cuda_core/tests/test_program.py b/cuda_core/tests/test_program.py index f19b0d34da5..f68b2c52aff 100644 --- a/cuda_core/tests/test_program.py +++ b/cuda_core/tests/test_program.py @@ -3,13 +3,13 @@ # SPDX-License-Identifier: Apache-2.0 import contextlib +import io import os import re -import shutil -import subprocess import warnings import pytest +from elftools.elf.elffile import ELFFile from cuda.core import _linker from cuda.core._device import Device @@ -370,9 +370,36 @@ def test_program_options_name_accepts_none(name): assert options._name == expected.encode() +def _dwarf_source_path(cubin: bytes) -> str: + """The source path a debugger resolves for ``cubin``. + + A cubin is an ELF64 image carrying ordinary DWARF, so the compilation + unit's ``DW_AT_name`` is the path cuda-gdb opens to display source. + """ + with io.BytesIO(cubin) as stream: + dwarf = ELFFile(stream).get_dwarf_info(relocate_dwarf_sections=False) + units = list(dwarf.iter_CUs()) + assert len(units) == 1, f"expected one compilation unit, got {len(units)}" + + top = units[0].get_top_DIE() + path = top.attributes["DW_AT_name"].value.decode() + comp_dir = top.attributes.get("DW_AT_comp_dir") + if not os.path.isabs(path) and comp_dir is not None: + path = os.path.join(comp_dir.value.decode(), path) + + # The line table is what actually maps addresses to lines, so confirm it + # names the same file instead of trusting DW_AT_name alone. Comparing + # basenames avoids the dir_index encoding, which differs across DWARF + # versions. + listed = {entry.name.decode() for entry in dwarf.line_program_for_CU(units[0]).header["file_entry"]} + assert os.path.basename(path) in listed, f"{path!r} absent from line table {sorted(listed)}" + + return path + + @pytest.mark.human_reviewed @pytest.mark.parametrize("name", [None, "my_program"]) -def test_program_string_source_debug(name, tmp_path): +def test_program_string_source_debug(name): # when a file doesn't exist on disk, nvrtc # still ends up referencing a file it thinks is there # in the dwarf table. As a WAR, we put the file there @@ -382,33 +409,9 @@ def test_program_string_source_debug(name, tmp_path): program = Program(code, "c++", options={"name": name, "debug": True, "lineinfo": True}) cubin = program.compile("cubin") - # read dwarf table - nvdisasm = shutil.which("nvdisasm") - - cubin_path = tmp_path / "program.cubin" - cubin_path.write_bytes(bytes(cubin.code)) - - result = subprocess.run( # noqa: S603 - [nvdisasm, "-g", str(cubin_path)], - capture_output=True, - text=True, - errors="replace", - ) - if result.returncode != 0: - pytest.fail(f"nvdisasm -g failed with exit code {result.returncode}\n{result.stderr}", pytrace=False) - - # -g annotates each instruction with the line-table entry that covers it: - # //## File "/abs/dir/name.cu", line 12 - # The directory comes from the line table's directory table, so this is the - # fully resolved path cuda-gdb will try to open. - paths = set(re.findall(r'//## File "([^"]+)", line \d+', result.stdout)) - assert len(paths) == 1, f"expected exactly one source file in the line table, got {sorted(paths)}" - dwarf_path = paths.pop() - - # cuda-gdb opens this path literally, so the source has to actually be - # sitting there for source-level debugging to work. - assert os.path.isfile(dwarf_path) + dwarf_path = _dwarf_source_path(bytes(cubin.code)) + assert os.path.isfile(dwarf_path), f"no source file on disk at the DWARF path {dwarf_path!r}" with open(dwarf_path, encoding="utf-8") as source_file: assert source_file.read().splitlines() == code.splitlines() From 88b1cf600d794bd08e598c0d5040b8172f63473d Mon Sep 17 00:00:00 2001 From: brandon-b-miller Date: Mon, 31 Aug 2026 06:46:33 -0700 Subject: [PATCH 4/4] 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. --- cuda_core/cuda/core/_program.pxd | 2 ++ cuda_core/cuda/core/_program.pyx | 22 +++++++++++-- cuda_core/tests/test_program.py | 56 ++++++++++++++++++++++++++++++++ 3 files changed, 77 insertions(+), 3 deletions(-) diff --git a/cuda_core/cuda/core/_program.pxd b/cuda_core/cuda/core/_program.pxd index 7c82119bcbd..9464a5ef28e 100644 --- a/cuda_core/cuda/core/_program.pxd +++ b/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 diff --git a/cuda_core/cuda/core/_program.pyx b/cuda_core/cuda/core/_program.pyx index 2d946730233..8307436b34d 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. 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(), 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