From 268dd9b6e101cfedc22f183a917c47844e7441b6 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Tue, 7 Jul 2026 17:26:23 -0700 Subject: [PATCH 1/6] Check nvJitLink driver major compatibility --- cuda_core/cuda/core/_linker.pyi | 3 +- cuda_core/cuda/core/_linker.pyx | 30 ++++++--- .../tests/test_optional_dependency_imports.py | 63 +++++++++++++++++++ 3 files changed, 88 insertions(+), 8 deletions(-) diff --git a/cuda_core/cuda/core/_linker.pyi b/cuda_core/cuda/core/_linker.pyi index 42b08313f78..8b1fde49ee3 100644 --- a/cuda_core/cuda/core/_linker.pyi +++ b/cuda_core/cuda/core/_linker.pyi @@ -112,7 +112,8 @@ class LinkerOptions: Since the linker may choose to use nvJitLink or the driver APIs as the linking backend, not all options are applicable. When the system's installed nvJitLink is too old (<12.3), - or not installed, the driver APIs (cuLink) will be used instead. + not installed, or does not match the CUDA driver major version, the driver APIs + (cuLink) will be used instead. Attributes ---------- diff --git a/cuda_core/cuda/core/_linker.pyx b/cuda_core/cuda/core/_linker.pyx index 10a85b998a1..fed568ceb7b 100644 --- a/cuda_core/cuda/core/_linker.pyx +++ b/cuda_core/cuda/core/_linker.pyx @@ -39,6 +39,7 @@ from cuda.core._utils.cuda_utils import ( driver, is_sequence, ) +from cuda.core._utils.version import driver_version from cuda.core.typing import CompilerBackendType, ObjectCodeFormatType if TYPE_CHECKING: @@ -207,7 +208,8 @@ class LinkerOptions: Since the linker may choose to use nvJitLink or the driver APIs as the linking backend, not all options are applicable. When the system's installed nvJitLink is too old (<12.3), - or not installed, the driver APIs (cuLink) will be used instead. + not installed, or does not match the CUDA driver major version, the driver APIs + (cuLink) will be used instead. Attributes ---------- @@ -680,12 +682,26 @@ def _decide_nvjitlink_or_driver() -> bool: from cuda.bindings._internal import nvjitlink if _nvjitlink_has_version_symbol(nvjitlink): - _use_nvjitlink_backend = True - return False # Use nvjitlink - warn_txt = ( - f"{'nvJitLink*.dll' if sys.platform == 'win32' else 'libnvJitLink.so*'} is too old (<12.3)." - f" Therefore cuda.bindings.nvjitlink is not usable and {warn_txt_common} nvJitLink." - ) + nvjitlink_version = nvjitlink_module.version() + try: + driver_major = driver_version()[0] + except CUDAError: + driver_major = None + + if driver_major is None or driver_major == nvjitlink_version[0]: + _use_nvjitlink_backend = True + return False # Use nvjitlink + + warn_txt = ( + f"CUDA driver major version {driver_major} does not match " + f"nvJitLink major version {nvjitlink_version[0]}; therefore " + f"{warn_txt_common} nvJitLink." + ) + else: + warn_txt = ( + f"{'nvJitLink*.dll' if sys.platform == 'win32' else 'libnvJitLink.so*'} is too old (<12.3)." + f" Therefore cuda.bindings.nvjitlink is not usable and {warn_txt_common} nvJitLink." + ) warn(warn_txt, stacklevel=2, category=RuntimeWarning) _use_nvjitlink_backend = False diff --git a/cuda_core/tests/test_optional_dependency_imports.py b/cuda_core/tests/test_optional_dependency_imports.py index 02edcc9839a..48494cb5c05 100644 --- a/cuda_core/tests/test_optional_dependency_imports.py +++ b/cuda_core/tests/test_optional_dependency_imports.py @@ -7,6 +7,14 @@ from cuda.core import _linker, _program +class FakeNvJitLinkModule: + def __init__(self, version): + self._version = version + + def version(self): + return self._version + + @pytest.fixture(autouse=True) def restore_optional_import_state(): saved_nvvm_module = _program._nvvm_module @@ -103,3 +111,58 @@ def fake__optional_cuda_import(modname, probe_function=None): assert use_driver_backend is True assert _linker._use_nvjitlink_backend is False + + +@pytest.mark.agent_authored(model="gpt-5") +def test_decide_nvjitlink_or_driver_uses_nvjitlink_when_majors_match(monkeypatch): + monkeypatch.setattr( + _linker, + "_optional_cuda_import", + lambda *_args, **_kwargs: FakeNvJitLinkModule((12, 9)), + ) + monkeypatch.setattr(_linker, "_nvjitlink_has_version_symbol", lambda _nvjitlink: True) + monkeypatch.setattr(_linker, "driver_version", lambda: (12, 8, 0)) + + assert _linker._decide_nvjitlink_or_driver() is False + assert _linker._use_nvjitlink_backend is True + + +@pytest.mark.agent_authored(model="gpt-5") +@pytest.mark.parametrize( + ("driver_version", "nvjitlink_version"), + [ + ((13, 0, 0), (12, 9)), + ((12, 9, 0), (13, 0)), + ], +) +def test_decide_nvjitlink_or_driver_falls_back_when_major_mismatches(monkeypatch, driver_version, nvjitlink_version): + monkeypatch.setattr( + _linker, + "_optional_cuda_import", + lambda *_args, **_kwargs: FakeNvJitLinkModule(nvjitlink_version), + ) + monkeypatch.setattr(_linker, "_nvjitlink_has_version_symbol", lambda _nvjitlink: True) + monkeypatch.setattr(_linker, "driver_version", lambda: driver_version) + + with pytest.warns(RuntimeWarning, match="does not match nvJitLink major version"): + use_driver_backend = _linker._decide_nvjitlink_or_driver() + + assert use_driver_backend is True + assert _linker._use_nvjitlink_backend is False + + +@pytest.mark.agent_authored(model="gpt-5") +def test_decide_nvjitlink_or_driver_uses_nvjitlink_when_driver_version_unknown(monkeypatch): + def raise_cuda_error(): + raise _linker.CUDAError("driver unavailable") + + monkeypatch.setattr( + _linker, + "_optional_cuda_import", + lambda *_args, **_kwargs: FakeNvJitLinkModule((12, 9)), + ) + monkeypatch.setattr(_linker, "_nvjitlink_has_version_symbol", lambda _nvjitlink: True) + monkeypatch.setattr(_linker, "driver_version", raise_cuda_error) + + assert _linker._decide_nvjitlink_or_driver() is False + assert _linker._use_nvjitlink_backend is True From 9a91b81c93aabdf385e3f639463cfd8d4e3f95c9 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Wed, 8 Jul 2026 15:01:20 -0700 Subject: [PATCH 2/6] Narrow nvJitLink driver compatibility fallback --- cuda_core/cuda/core/_linker.pyi | 4 +- cuda_core/cuda/core/_linker.pyx | 14 +++---- .../tests/test_optional_dependency_imports.py | 41 ++++++------------- 3 files changed, 20 insertions(+), 39 deletions(-) diff --git a/cuda_core/cuda/core/_linker.pyi b/cuda_core/cuda/core/_linker.pyi index 8b1fde49ee3..37e05fc7fe6 100644 --- a/cuda_core/cuda/core/_linker.pyi +++ b/cuda_core/cuda/core/_linker.pyi @@ -112,8 +112,8 @@ class LinkerOptions: Since the linker may choose to use nvJitLink or the driver APIs as the linking backend, not all options are applicable. When the system's installed nvJitLink is too old (<12.3), - not installed, or does not match the CUDA driver major version, the driver APIs - (cuLink) will be used instead. + not installed, or older than the CUDA driver major version, the driver APIs (cuLink) + will be used instead. Attributes ---------- diff --git a/cuda_core/cuda/core/_linker.pyx b/cuda_core/cuda/core/_linker.pyx index fed568ceb7b..4cadaf0d8dc 100644 --- a/cuda_core/cuda/core/_linker.pyx +++ b/cuda_core/cuda/core/_linker.pyx @@ -208,8 +208,8 @@ class LinkerOptions: Since the linker may choose to use nvJitLink or the driver APIs as the linking backend, not all options are applicable. When the system's installed nvJitLink is too old (<12.3), - not installed, or does not match the CUDA driver major version, the driver APIs - (cuLink) will be used instead. + not installed, or older than the CUDA driver major version, the driver APIs (cuLink) + will be used instead. Attributes ---------- @@ -683,17 +683,13 @@ def _decide_nvjitlink_or_driver() -> bool: if _nvjitlink_has_version_symbol(nvjitlink): nvjitlink_version = nvjitlink_module.version() - try: - driver_major = driver_version()[0] - except CUDAError: - driver_major = None - - if driver_major is None or driver_major == nvjitlink_version[0]: + driver_major = driver_version()[0] + if driver_major <= nvjitlink_version[0]: _use_nvjitlink_backend = True return False # Use nvjitlink warn_txt = ( - f"CUDA driver major version {driver_major} does not match " + f"CUDA driver major version {driver_major} is newer than " f"nvJitLink major version {nvjitlink_version[0]}; therefore " f"{warn_txt_common} nvJitLink." ) diff --git a/cuda_core/tests/test_optional_dependency_imports.py b/cuda_core/tests/test_optional_dependency_imports.py index 48494cb5c05..6759fd3cfa4 100644 --- a/cuda_core/tests/test_optional_dependency_imports.py +++ b/cuda_core/tests/test_optional_dependency_imports.py @@ -113,29 +113,17 @@ def fake__optional_cuda_import(modname, probe_function=None): assert _linker._use_nvjitlink_backend is False -@pytest.mark.agent_authored(model="gpt-5") -def test_decide_nvjitlink_or_driver_uses_nvjitlink_when_majors_match(monkeypatch): - monkeypatch.setattr( - _linker, - "_optional_cuda_import", - lambda *_args, **_kwargs: FakeNvJitLinkModule((12, 9)), - ) - monkeypatch.setattr(_linker, "_nvjitlink_has_version_symbol", lambda _nvjitlink: True) - monkeypatch.setattr(_linker, "driver_version", lambda: (12, 8, 0)) - - assert _linker._decide_nvjitlink_or_driver() is False - assert _linker._use_nvjitlink_backend is True - - @pytest.mark.agent_authored(model="gpt-5") @pytest.mark.parametrize( ("driver_version", "nvjitlink_version"), [ - ((13, 0, 0), (12, 9)), + ((12, 8, 0), (12, 9)), ((12, 9, 0), (13, 0)), ], ) -def test_decide_nvjitlink_or_driver_falls_back_when_major_mismatches(monkeypatch, driver_version, nvjitlink_version): +def test_decide_nvjitlink_or_driver_uses_nvjitlink_when_driver_is_not_newer( + monkeypatch, driver_version, nvjitlink_version +): monkeypatch.setattr( _linker, "_optional_cuda_import", @@ -144,25 +132,22 @@ def test_decide_nvjitlink_or_driver_falls_back_when_major_mismatches(monkeypatch monkeypatch.setattr(_linker, "_nvjitlink_has_version_symbol", lambda _nvjitlink: True) monkeypatch.setattr(_linker, "driver_version", lambda: driver_version) - with pytest.warns(RuntimeWarning, match="does not match nvJitLink major version"): - use_driver_backend = _linker._decide_nvjitlink_or_driver() - - assert use_driver_backend is True - assert _linker._use_nvjitlink_backend is False + assert _linker._decide_nvjitlink_or_driver() is False + assert _linker._use_nvjitlink_backend is True @pytest.mark.agent_authored(model="gpt-5") -def test_decide_nvjitlink_or_driver_uses_nvjitlink_when_driver_version_unknown(monkeypatch): - def raise_cuda_error(): - raise _linker.CUDAError("driver unavailable") - +def test_decide_nvjitlink_or_driver_falls_back_when_driver_is_newer(monkeypatch): monkeypatch.setattr( _linker, "_optional_cuda_import", lambda *_args, **_kwargs: FakeNvJitLinkModule((12, 9)), ) monkeypatch.setattr(_linker, "_nvjitlink_has_version_symbol", lambda _nvjitlink: True) - monkeypatch.setattr(_linker, "driver_version", raise_cuda_error) + monkeypatch.setattr(_linker, "driver_version", lambda: (13, 0, 0)) - assert _linker._decide_nvjitlink_or_driver() is False - assert _linker._use_nvjitlink_backend is True + with pytest.warns(RuntimeWarning, match="is newer than nvJitLink major version"): + use_driver_backend = _linker._decide_nvjitlink_or_driver() + + assert use_driver_backend is True + assert _linker._use_nvjitlink_backend is False From 71e7cbcb784c6fe05980615b2eac4fb21630c4c4 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Wed, 8 Jul 2026 15:56:15 -0700 Subject: [PATCH 3/6] Fix cuLink state storage lifetimes --- cuda_core/cuda/core/_linker.pxd | 10 +++++++- cuda_core/cuda/core/_linker.pyx | 38 +++++++++++++++++++---------- cuda_core/tests/test_linker.py | 42 +++++++++++++++++++++++++++++++++ 3 files changed, 77 insertions(+), 13 deletions(-) diff --git a/cuda_core/cuda/core/_linker.pxd b/cuda_core/cuda/core/_linker.pxd index e50ebb97705..1b7d39fd1d4 100644 --- a/cuda_core/cuda/core/_linker.pxd +++ b/cuda_core/cuda/core/_linker.pxd @@ -2,6 +2,10 @@ # # SPDX-License-Identifier: Apache-2.0 +from libcpp.vector cimport vector + +from cuda.bindings cimport cydriver + from ._resource_handles cimport NvJitLinkHandle, CuLinkHandle @@ -9,8 +13,12 @@ cdef class Linker: cdef: NvJitLinkHandle _nvjitlink_handle CuLinkHandle _culink_handle + # The driver retains these arrays until cuLinkDestroy. Declare them + # after the handle so their destructors run after cuLinkDestroy. + vector[cydriver.CUjit_option] _drv_jit_keys + vector[void*] _drv_jit_values bint _use_nvjitlink - object _drv_log_bufs # formatted_options list (driver); None for nvjitlink; cleared in link() + object _drv_log_bufs # formatted_options list (driver); None for nvjitlink str _info_log # decoded log; None until link() or pre-link get_*_log() str _error_log # decoded log; None until link() or pre-link get_*_log() object _options # LinkerOptions diff --git a/cuda_core/cuda/core/_linker.pyx b/cuda_core/cuda/core/_linker.pyx index 4cadaf0d8dc..172b176c844 100644 --- a/cuda_core/cuda/core/_linker.pyx +++ b/cuda_core/cuda/core/_linker.pyx @@ -53,7 +53,6 @@ _keep_driver_in_stub: "cuda.bindings.driver.CUlinkState" _keep_nvjitlink_in_stub: "cuda.bindings.nvjitlink.nvJitLinkHandle" ctypedef const char* const_char_ptr -ctypedef void* void_ptr __all__ = ["Linker", "LinkerOptions"] @@ -156,10 +155,21 @@ cdef class Linker: def close(self) -> None: """Destroy this linker.""" + cdef vector[cydriver.CUjit_option] empty_keys + cdef vector[void*] empty_values if self._use_nvjitlink: self._nvjitlink_handle.reset() else: + if self._drv_log_bufs is not None: + if self._info_log is None: + self._info_log = self.get_info_log() + if self._error_log is None: + self._error_log = self.get_error_log() + # Destroy the CUlinkState before releasing storage referenced by it. self._culink_handle.reset() + self._drv_jit_keys.swap(empty_keys) + self._drv_jit_values.swap(empty_values) + self._drv_log_bufs = None @property def handle(self) -> LinkerHandleT: @@ -475,8 +485,8 @@ cdef inline int Linker_init(Linker self, tuple object_codes, object options) exc cdef cydriver.CUlinkState c_raw_culink cdef Py_ssize_t c_num_opts, i cdef vector[const_char_ptr] c_str_opts - cdef vector[cydriver.CUjit_option] c_jit_keys - cdef vector[void_ptr] c_jit_values + cdef cydriver.CUjit_option* c_drv_jit_keys_ptr + cdef void** c_drv_jit_values_ptr self._options = options = check_or_create_options(LinkerOptions, options, "Linker options") @@ -498,19 +508,24 @@ cdef inline int Linker_init(Linker self, tuple object_codes, object options) exc # the driver writes into via raw pointers during linking operations. self._drv_log_bufs = formatted_options c_num_opts = len(option_keys) - c_jit_keys.resize(c_num_opts) - c_jit_values.resize(c_num_opts) + self._drv_jit_keys.resize(c_num_opts) + self._drv_jit_values.resize(c_num_opts) for i in range(c_num_opts): - c_jit_keys[i] = option_keys[i] + self._drv_jit_keys[i] = option_keys[i] val = formatted_options[i] if isinstance(val, bytearray): - c_jit_values[i] = PyByteArray_AS_STRING(val) + self._drv_jit_values[i] = PyByteArray_AS_STRING(val) else: - c_jit_values[i] = int(val) + self._drv_jit_values[i] = int(val) + c_drv_jit_keys_ptr = self._drv_jit_keys.data() + c_drv_jit_values_ptr = self._drv_jit_values.data() try: with nogil: HANDLE_RETURN(cydriver.cuLinkCreate( - c_num_opts, c_jit_keys.data(), c_jit_values.data(), &c_raw_culink)) + c_num_opts, + c_drv_jit_keys_ptr, + c_drv_jit_values_ptr, + &c_raw_culink)) except CUDAError as e: Linker_annotate_error_log(self, e) raise @@ -624,11 +639,10 @@ cdef inline object Linker_link(Linker self, str target_type): raise code = (c_cubin_out)[:c_output_size] - # Linking is complete; cache the decoded log strings and release - # the driver's raw bytearray buffers (no longer written to). + # Linking is complete; cache the decoded logs. cuLinkDestroy may still + # dereference the raw log-buffer pointers, so retain them until close(). self._info_log = self.get_info_log() self._error_log = self.get_error_log() - self._drv_log_bufs = None return ObjectCode._init(bytes(code), target_type, name=self._options.name) diff --git a/cuda_core/tests/test_linker.py b/cuda_core/tests/test_linker.py index 942ccb10640..41cee4d9f31 100644 --- a/cuda_core/tests/test_linker.py +++ b/cuda_core/tests/test_linker.py @@ -8,6 +8,7 @@ from cuda.core import Device, Linker, LinkerOptions, Program, ProgramOptions, _linker from cuda.core._module import ObjectCode +from cuda.core._program import _can_load_generated_ptx from cuda.core._utils.cuda_utils import CUDAError ARCH = "sm_" + "".join(f"{i}" for i in Device().compute_capability) @@ -234,6 +235,47 @@ def test_linker_handle(compile_ptx_functions): assert int(handle) != 0 +@pytest.mark.agent_authored(model="gpt-5") +@pytest.mark.skipif(not is_culink_backend, reason="driver backend regression test") +def test_driver_linker_lifetime_no_heap_corruption(compile_ptx_functions): + if not _can_load_generated_ptx(): + pytest.skip("PTX version too new for current driver") + + linker = Linker(*compile_ptx_functions, options=LinkerOptions(arch=ARCH)) + linker.link("cubin") + linker.close() + del linker + + obj_a = Program(kernel_a, "c++", ProgramOptions(relocatable_device_code=True)).compile("ptx") + obj_b = Program(device_function_b, "c++", ProgramOptions(relocatable_device_code=True)).compile("ptx") + obj_c = Program(device_function_c, "c++", ProgramOptions(relocatable_device_code=True)).compile("ptx") + linker = Linker(obj_a, obj_b, obj_c, options=LinkerOptions(arch=ARCH)) + linker.link("cubin") + linker.close() + + +@pytest.mark.agent_authored(model="gpt-5") +@pytest.mark.skipif(not is_culink_backend, reason="driver backend regression test") +def test_driver_linker_preserves_error_log_after_close(init_cuda): + if not _can_load_generated_ptx(): + pytest.skip("PTX version too new for current driver") + + bad_kernel = """ +extern __device__ int Z(); +__global__ void A() { int r = Z(); } +""" + bad_obj = Program(bad_kernel, "c++", ProgramOptions(relocatable_device_code=True)).compile("ptx") + linker = Linker(bad_obj, options=LinkerOptions(arch=ARCH)) + with pytest.raises(CUDAError): + linker.link("cubin") + + error_log = linker.get_error_log() + assert error_log + linker.close() + assert linker.get_error_log() == error_log + assert isinstance(linker.get_info_log(), str) + + @pytest.mark.skipif(is_culink_backend, reason="nvjitlink options only tested with nvjitlink backend") def test_linker_options_nvjitlink_options_as_str(): """_prepare_nvjitlink_options(as_bytes=False) returns plain strings.""" From 93c448e3e0234d0ba3c4fbc2eb8cf4c0699f9715 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Wed, 8 Jul 2026 16:59:00 -0700 Subject: [PATCH 4/6] Isolate nvJitLink-specific cache tests --- cuda_core/tests/test_linker.py | 13 +++---------- cuda_core/tests/test_program_cache.py | 27 +++++++++++++++++++++------ 2 files changed, 24 insertions(+), 16 deletions(-) diff --git a/cuda_core/tests/test_linker.py b/cuda_core/tests/test_linker.py index 41cee4d9f31..9d95b5fd9c3 100644 --- a/cuda_core/tests/test_linker.py +++ b/cuda_core/tests/test_linker.py @@ -197,19 +197,12 @@ def test_linker_options_as_bytes_nvjitlink(): assert "-maxrregcount=32" in options_str -def test_linker_options_as_bytes_invalid_backend(): +@pytest.mark.parametrize("backend", ("invalid", "driver")) +def test_linker_options_as_bytes_invalid_backend(backend): """Test LinkerOptions.as_bytes() with invalid backend""" options = LinkerOptions(arch="sm_80") with pytest.raises(ValueError, match="only supports 'nvjitlink' backend"): - options.as_bytes("invalid") - - -@pytest.mark.skipif(not is_culink_backend, reason="driver backend test") -def test_linker_options_as_bytes_driver_not_supported(): - """Test that as_bytes() is not supported for driver backend""" - options = LinkerOptions(arch="sm_80") - with pytest.raises(RuntimeError, match="as_bytes\\(\\) only supports 'nvjitlink' backend"): - options.as_bytes("driver") + options.as_bytes(backend) def test_linker_logs_cached_after_link(compile_ptx_functions): diff --git a/cuda_core/tests/test_program_cache.py b/cuda_core/tests/test_program_cache.py index 963ec1cc04b..eb9f856482b 100644 --- a/cuda_core/tests/test_program_cache.py +++ b/cuda_core/tests/test_program_cache.py @@ -123,6 +123,14 @@ def _make_key(**overrides): return make_program_cache_key(**{**base, **overrides}) +@pytest.fixture +def require_nvjitlink_backend(): + from cuda.core import _linker + + if _linker._decide_nvjitlink_or_driver(): + pytest.skip("test requires the nvJitLink backend") + + def test_make_program_cache_key_returns_bytes(): key = _make_key() assert isinstance(key, bytes) @@ -369,9 +377,6 @@ def test_make_program_cache_key_ignores_name_expressions_for_non_nvrtc(code_type {"link_time_optimization": None}, id="lto_false_eq_none", ), - # ``time`` is a presence gate: the linker emits ``-time`` for any - # non-None value, so True / "path" produce the same flag. - pytest.param({"time": True}, {"time": "timing.csv"}, id="time_true_eq_path"), # ``no_cache`` has an ``is True`` gate; False and None equivalent. pytest.param({"no_cache": False}, {"no_cache": None}, id="no_cache_false_eq_none"), ], @@ -390,6 +395,17 @@ def test_make_program_cache_key_ptx_linker_equivalent_options_hash_same(a, b, mo assert k_a == k_b +@pytest.mark.agent_authored(model="gpt-5") +def test_make_program_cache_key_ptx_nvjitlink_time_values_hash_same(require_nvjitlink_backend, monkeypatch): + """nvJitLink emits ``-time`` for any non-None value.""" + from cuda.core.utils import _program_cache + + monkeypatch.setattr(_program_cache._keys, "_linker_backend_and_version", lambda _use_driver: ("nvJitLink", "12030")) + k_true = _make_key(code=".version 7.0", code_type="ptx", options=_opts(time=True)) + k_path = _make_key(code=".version 7.0", code_type="ptx", options=_opts(time="timing.csv")) + assert k_true == k_path + + @pytest.mark.parametrize( "field, a, b", [ @@ -999,9 +1015,8 @@ def test_make_program_cache_key_rejects_side_effect_options_nvrtc(option_kw, ext pytest.param({"time": "whatever.csv"}, id="time_path"), ], ) -def test_make_program_cache_key_accepts_side_effect_options_for_ptx(option_kw): - """The side-effect guard is NVRTC-specific: PTX (linker) and NVVM must - not be blocked by options whose side effects only apply under NVRTC.""" +def test_make_program_cache_key_accepts_side_effect_options_for_ptx(option_kw, require_nvjitlink_backend): + """nvJitLink ``-time`` writes only to its info log, not the filesystem.""" _make_key(code=".version 7.0", code_type="ptx", options=_opts(**option_kw)) # no raise From 32ab374bfab8ac9536d9ab3b8300895dade8685a Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Wed, 8 Jul 2026 17:00:52 -0700 Subject: [PATCH 5/6] Use public linker backend query in tests --- cuda_core/tests/test_program_cache.py | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/cuda_core/tests/test_program_cache.py b/cuda_core/tests/test_program_cache.py index eb9f856482b..2b42afb979b 100644 --- a/cuda_core/tests/test_program_cache.py +++ b/cuda_core/tests/test_program_cache.py @@ -125,9 +125,10 @@ def _make_key(**overrides): @pytest.fixture def require_nvjitlink_backend(): - from cuda.core import _linker + from cuda.core import Linker + from cuda.core.typing import CompilerBackendType - if _linker._decide_nvjitlink_or_driver(): + if Linker.which_backend() == CompilerBackendType.DRIVER: pytest.skip("test requires the nvJitLink backend") From e896ce0d10062f7e466145a39522daa78fd23d33 Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Thu, 9 Jul 2026 14:35:13 -0700 Subject: [PATCH 6/6] Use direct backend skips in program cache tests --- cuda_core/tests/test_program_cache.py | 19 ++++++++----------- 1 file changed, 8 insertions(+), 11 deletions(-) diff --git a/cuda_core/tests/test_program_cache.py b/cuda_core/tests/test_program_cache.py index 2b42afb979b..3bb579e2fe8 100644 --- a/cuda_core/tests/test_program_cache.py +++ b/cuda_core/tests/test_program_cache.py @@ -7,6 +7,10 @@ import pytest +from cuda.core import _linker + +is_culink_backend = _linker._decide_nvjitlink_or_driver() + def test_program_cache_resource_is_abstract(): from cuda.core.utils import ProgramCacheResource @@ -123,15 +127,6 @@ def _make_key(**overrides): return make_program_cache_key(**{**base, **overrides}) -@pytest.fixture -def require_nvjitlink_backend(): - from cuda.core import Linker - from cuda.core.typing import CompilerBackendType - - if Linker.which_backend() == CompilerBackendType.DRIVER: - pytest.skip("test requires the nvJitLink backend") - - def test_make_program_cache_key_returns_bytes(): key = _make_key() assert isinstance(key, bytes) @@ -396,8 +391,9 @@ def test_make_program_cache_key_ptx_linker_equivalent_options_hash_same(a, b, mo assert k_a == k_b +@pytest.mark.skipif(is_culink_backend, reason="test requires nvJitLink backend") @pytest.mark.agent_authored(model="gpt-5") -def test_make_program_cache_key_ptx_nvjitlink_time_values_hash_same(require_nvjitlink_backend, monkeypatch): +def test_make_program_cache_key_ptx_nvjitlink_time_values_hash_same(monkeypatch): """nvJitLink emits ``-time`` for any non-None value.""" from cuda.core.utils import _program_cache @@ -1016,7 +1012,8 @@ def test_make_program_cache_key_rejects_side_effect_options_nvrtc(option_kw, ext pytest.param({"time": "whatever.csv"}, id="time_path"), ], ) -def test_make_program_cache_key_accepts_side_effect_options_for_ptx(option_kw, require_nvjitlink_backend): +@pytest.mark.skipif(is_culink_backend, reason="test requires nvJitLink backend") +def test_make_program_cache_key_accepts_side_effect_options_for_ptx(option_kw): """nvJitLink ``-time`` writes only to its info log, not the filesystem.""" _make_key(code=".version 7.0", code_type="ptx", options=_opts(**option_kw)) # no raise