From 407edd77444755af066ed32e5416bba0ae15d4f0 Mon Sep 17 00:00:00 2001 From: Jinfeng Date: Thu, 20 Aug 2026 23:31:49 +0000 Subject: [PATCH 1/3] first version after cleaning unnecessary code --- cuda_core/cuda/core/_program.pxd | 1 + cuda_core/cuda/core/_program.pyx | 60 +++++++++++++++++++++++++++++--- cuda_core/tests/test_program.py | 32 +++++++++++++++++ 3 files changed, 88 insertions(+), 5 deletions(-) diff --git a/cuda_core/cuda/core/_program.pxd b/cuda_core/cuda/core/_program.pxd index cea430c3f20..7c82119bcbd 100644 --- a/cuda_core/cuda/core/_program.pxd +++ b/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 diff --git a/cuda_core/cuda/core/_program.pyx b/cuda_core/cuda/core/_program.pyx index 1d07b88bbf4..91853de8868 100644 --- a/cuda_core/cuda/core/_program.pyx +++ b/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,14 @@ 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() + _unlink_debug_source(path) def compile( self, @@ -223,7 +235,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 @@ -768,6 +780,38 @@ cdef inline object _translate_program_options(object options): ) +def _unlink_debug_source(path: str) -> None: + try: + os.unlink(path) + except OSError: + pass + + +def _try_materialize_nvrtc_debug_source(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: + _unlink_debug_source(path) + return None + + 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,16 +832,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 = _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 = code_bytes - name_ptr = options._name + name_ptr = self._nvrtc_name with nogil: HANDLE_RETURN_NVRTC(NULL, cynvrtc.nvrtcCreateProgram( @@ -969,7 +1019,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 +1047,14 @@ 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._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) diff --git a/cuda_core/tests/test_program.py b/cuda_core/tests/test_program.py index 72f8b8f5942..c6d2cf939ef 100644 --- a/cuda_core/tests/test_program.py +++ b/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 From a67bbb69b20e9937d630983c4330d8aa0dd1b21c Mon Sep 17 00:00:00 2001 From: Jinfeng Date: Thu, 20 Aug 2026 23:34:24 +0000 Subject: [PATCH 2/3] update _program.pyi --- cuda_core/cuda/core/_program.pyi | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/cuda_core/cuda/core/_program.pyi b/cuda_core/cuda/core/_program.pyi index 40d2f6c7dee..b73be7791ad 100644 --- a/cuda_core/cuda/core/_program.pyi +++ b/cuda_core/cuda/core/_program.pyi @@ -40,6 +40,8 @@ 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 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. @@ -424,6 +426,13 @@ def _find_libdevice_path() -> object: """Find libdevice*.bc for NVVM compilation using cuda.pathfinder.""" def _can_load_generated_ptx() -> bool: """Check if the driver can load PTX generated by the current NVRTC version.""" +def _unlink_debug_source(path: str) -> None: ... +def _try_materialize_nvrtc_debug_source(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 _assert_single_dashed_nvvm_options(options: list[str]) -> None: """Guard against emitting a double-dashed option to libNVVM. From 485de34b09e5a2edea13882c7244d9f27d92db0c Mon Sep 17 00:00:00 2001 From: Jinfeng Date: Thu, 20 Aug 2026 23:53:45 +0000 Subject: [PATCH 3/3] 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. --- cuda_core/cuda/core/_program.pyi | 14 +++---- cuda_core/cuda/core/_program.pyx | 66 ++++++++++++++++---------------- 2 files changed, 39 insertions(+), 41 deletions(-) diff --git a/cuda_core/cuda/core/_program.pyi b/cuda_core/cuda/core/_program.pyi index b73be7791ad..a1f62562253 100644 --- a/cuda_core/cuda/core/_program.pyi +++ b/cuda_core/cuda/core/_program.pyi @@ -42,6 +42,13 @@ class Program: """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. @@ -426,13 +433,6 @@ def _find_libdevice_path() -> object: """Find libdevice*.bc for NVVM compilation using cuda.pathfinder.""" def _can_load_generated_ptx() -> bool: """Check if the driver can load PTX generated by the current NVRTC version.""" -def _unlink_debug_source(path: str) -> None: ... -def _try_materialize_nvrtc_debug_source(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 _assert_single_dashed_nvvm_options(options: list[str]) -> None: """Guard against emitting a double-dashed option to libNVVM. diff --git a/cuda_core/cuda/core/_program.pyx b/cuda_core/cuda/core/_program.pyx index 91853de8868..dad49b50edb 100644 --- a/cuda_core/cuda/core/_program.pyx +++ b/cuda_core/cuda/core/_program.pyx @@ -98,7 +98,37 @@ cdef class Program: def _cleanup_debug_source(self): path = self._nvrtc_name.decode() - _unlink_debug_source(path) + 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, @@ -780,38 +810,6 @@ cdef inline object _translate_program_options(object options): ) -def _unlink_debug_source(path: str) -> None: - try: - os.unlink(path) - except OSError: - pass - - -def _try_materialize_nvrtc_debug_source(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: - _unlink_debug_source(path) - return None - - 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 @@ -840,7 +838,7 @@ cdef inline int Program_init(Program self, object code, str code_type, object op 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 = _try_materialize_nvrtc_debug_source(code) + debug_path = self._try_materialize_nvrtc_debug_source(code) if debug_path is not None: self._nvrtc_name = debug_path.encode()