diff --git a/cuda_core/cuda/core/_linker.pxd b/cuda_core/cuda/core/_linker.pxd index 1b7d39fd1d4..0d4c226baa8 100644 --- a/cuda_core/cuda/core/_linker.pxd +++ b/cuda_core/cuda/core/_linker.pxd @@ -18,6 +18,7 @@ cdef class Linker: vector[cydriver.CUjit_option] _drv_jit_keys vector[void*] _drv_jit_values bint _use_nvjitlink + bint _has_ptx_or_cubin_input 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() diff --git a/cuda_core/cuda/core/_linker.pyi b/cuda_core/cuda/core/_linker.pyi index 4fed2399f7a..95700199d54 100644 --- a/cuda_core/cuda/core/_linker.pyi +++ b/cuda_core/cuda/core/_linker.pyi @@ -21,6 +21,8 @@ const_char_ptr: TypeAlias = bytes __all__ = ['Linker', 'LinkerOptions'] LinkerHandleT = Union['cuda.bindings.nvjitlink.nvJitLinkHandle', 'cuda.bindings.driver.CUlinkState'] _driver = None +_nvjitlink = None +_nvjitlink_version = None _inited = False _use_nvjitlink_backend = None _nvjitlink_input_types = None @@ -50,7 +52,9 @@ class Linker: Parameters ---------- target_type : ObjectCodeFormatType | str - The type of the target output. Must be either "cubin" or "ptx". + The type of the target output. Must be "cubin", "ptx", or + "ltoir". Linked LTOIR output requires + ``link_time_optimization=True`` and nvJitLink 13.3 or newer. Returns ------- @@ -61,6 +65,16 @@ class Linker: Ensure that input object codes were compiled with appropriate flags for linking (e.g., relocatable device code enabled). + + A CUBIN produced with ``incremental=True`` can be passed directly + to another :class:`Linker`, but it can still contain unresolved + device references and should be finalized before execution. + + ``"ltoir"`` output contains only the LTOIR carried by the inputs. + Direct PTX and CUBIN inputs are rejected because they carry no + LTOIR. FATBIN, host object, and library inputs are accepted but + may carry no LTOIR, in which case they contribute nothing to the + output. """ def get_error_log(self) -> str: """Get the error log generated by the linker. @@ -137,6 +151,11 @@ class LinkerOptions: link_time_optimization : bool, optional Perform link time optimization. Default: False. + incremental : bool, optional + Perform an incremental link. The result can be passed + directly to a later :class:`Linker`. Requires nvJitLink 13.2 or newer + and is not supported by the driver linker backend. + Default: False. ptx : bool, optional Emit PTX after linking instead of CUBIN; only supported with ``link_time_optimization=True``. Default: False. @@ -216,6 +235,7 @@ class LinkerOptions: split_compile_extended: int | None = None no_cache: bool | None = None numba_debug: bool | None = None + incremental: bool | None = None def __post_init__(self) -> None: ... def _prepare_nvjitlink_options(self, as_bytes: bool=False) -> list[bytes] | list[str]: ... @@ -241,6 +261,10 @@ class LinkerOptions: If nvJitLink backend is not available. """ +def _require_nvjitlink_version(minimum_version: tuple[int, int], feature: str) -> None: + """Check that the cached nvJitLink runtime meets a feature's requirement.""" +def _linked_ltoir_output_module(): + """Return bindings that can retrieve linked LTOIR without a Cython dependency.""" def _nvjitlink_has_version_symbol(nvjitlink) -> bool: ... def _decide_nvjitlink_or_driver() -> bool: """Return True if falling back to the cuLink* driver APIs.""" diff --git a/cuda_core/cuda/core/_linker.pyx b/cuda_core/cuda/core/_linker.pyx index f104e2d158b..ccfa35bfc16 100644 --- a/cuda_core/cuda/core/_linker.pyx +++ b/cuda_core/cuda/core/_linker.pyx @@ -18,6 +18,7 @@ from cuda.bindings cimport cynvjitlink from ._resource_handles cimport ( as_cu, + as_intptr, as_py, create_culink_handle, create_nvjitlink_handle, @@ -101,7 +102,9 @@ cdef class Linker: Parameters ---------- target_type : ObjectCodeFormatType | str - The type of the target output. Must be either "cubin" or "ptx". + The type of the target output. Must be "cubin", "ptx", or + "ltoir". Linked LTOIR output requires + ``link_time_optimization=True`` and nvJitLink 13.3 or newer. Returns ------- @@ -112,6 +115,16 @@ cdef class Linker: Ensure that input object codes were compiled with appropriate flags for linking (e.g., relocatable device code enabled). + + A CUBIN produced with ``incremental=True`` can be passed directly + to another :class:`Linker`, but it can still contain unresolved + device references and should be finalized before execution. + + ``"ltoir"`` output contains only the LTOIR carried by the inputs. + Direct PTX and CUBIN inputs are rejected because they carry no + LTOIR. FATBIN, host object, and library inputs are accepted but + may carry no LTOIR, in which case they contribute nothing to the + output. """ Linker_check_open(self) return Linker_link(self, str(target_type)) @@ -256,6 +269,11 @@ class LinkerOptions: link_time_optimization : bool, optional Perform link time optimization. Default: False. + incremental : bool, optional + Perform an incremental link. The result can be passed + directly to a later :class:`Linker`. Requires nvJitLink 13.2 or newer + and is not supported by the driver linker backend. + Default: False. ptx : bool, optional Emit PTX after linking instead of CUBIN; only supported with ``link_time_optimization=True``. Default: False. @@ -336,6 +354,7 @@ class LinkerOptions: split_compile_extended: int | None = None no_cache: bool | None = None numba_debug: bool | None = None + incremental: bool | None = None def __post_init__(self) -> None: _lazy_init() @@ -371,6 +390,8 @@ class LinkerOptions: options.append("-verbose") if self.link_time_optimization: options.append("-lto") + if self.incremental: + options.append("-r") if self.ptx: options.append("-ptx") if self.optimization_level is not None: @@ -450,6 +471,8 @@ class LinkerOptions: if self.link_time_optimization: formatted_options.append(1) option_keys.append(_driver.CUjit_option.CU_JIT_LTO) + if self.incremental: + raise ValueError("incremental option is not supported by the driver API") if self.ptx: raise ValueError("ptx option is not supported by the driver API") if self.optimization_level is not None: @@ -532,8 +555,13 @@ cdef inline int Linker_init(Linker self, tuple object_codes, object options) exc cdef void** c_drv_jit_values_ptr self._options = options = check_or_create_options(LinkerOptions, options, "Linker options") + self._has_ptx_or_cubin_input = False + if options.incremental and options.ptx: + raise ValueError("incremental and ptx output options cannot be used together") if _use_nvjitlink_backend: + if options.incremental: + _require_nvjitlink_version((13, 2), "incremental linking") self._use_nvjitlink = True options_bytes = options._prepare_nvjitlink_options(as_bytes=True) c_num_opts = len(options_bytes) @@ -639,17 +667,34 @@ cdef inline void Linker_add_code_object(Linker self, object object_code) except Linker_annotate_error_log(self, e) raise + if object_code.code_type in ("ptx", "cubin"): + self._has_ptx_or_cubin_input = True + cdef inline object Linker_link(Linker self, str target_type): """Complete linking and return the result as ObjectCode.""" - if target_type not in ("cubin", "ptx"): + if target_type not in ("cubin", "ptx", "ltoir"): raise ValueError(f"Unsupported target type: {target_type}") + if self._options.incremental and target_type == "ptx": + raise ValueError("PTX output is not supported for incremental linking") + if target_type == "ltoir": + if not self._use_nvjitlink: + raise ValueError("LTOIR output is not supported by the driver API") + if not self._options.link_time_optimization: + raise ValueError("LTOIR output requires link_time_optimization=True") + if self._has_ptx_or_cubin_input: + raise ValueError( + 'LTOIR output is not supported with "ptx" or "cubin" inputs; ' + "they carry no LTOIR and would be omitted from the output" + ) + nvjitlink_module = _linked_ltoir_output_module() cdef cynvjitlink.nvJitLinkHandle c_nvjitlink_h cdef cydriver.CUlinkState c_culink_state cdef size_t c_output_size = 0 cdef char* c_code_ptr cdef void* c_cubin_out = NULL + cdef intptr_t c_handle if self._use_nvjitlink: c_nvjitlink_h = as_cu(self._nvjitlink_handle) @@ -663,7 +708,7 @@ cdef inline object Linker_link(Linker self, str target_type): with nogil: HANDLE_RETURN_NVJITLINK(c_nvjitlink_h, cynvjitlink.nvJitLinkGetLinkedCubin(c_nvjitlink_h, c_code_ptr)) - else: + elif target_type == "ptx": HANDLE_RETURN_NVJITLINK(c_nvjitlink_h, cynvjitlink.nvJitLinkGetLinkedPtxSize(c_nvjitlink_h, &c_output_size)) code = bytearray(c_output_size) @@ -671,6 +716,11 @@ cdef inline object Linker_link(Linker self, str target_type): with nogil: HANDLE_RETURN_NVJITLINK(c_nvjitlink_h, cynvjitlink.nvJitLinkGetLinkedPtx(c_nvjitlink_h, c_code_ptr)) + else: + c_handle = as_intptr(self._nvjitlink_handle) + output_size = nvjitlink_module.get_linked_ltoir_size(c_handle) + code = bytearray(output_size) + nvjitlink_module.get_linked_ltoir(c_handle, code) else: c_culink_state = as_cu(self._culink_handle) try: @@ -702,6 +752,8 @@ cdef inline void Linker_annotate_error_log(Linker self, object e): # TODO: revisit this treatment for py313t builds _driver = None # populated if nvJitLink cannot be used +_nvjitlink = None # populated if nvJitLink can be used +_nvjitlink_version = None _inited = False _use_nvjitlink_backend = None # set by _decide_nvjitlink_or_driver() @@ -710,6 +762,31 @@ _nvjitlink_input_types = None _driver_input_types = None +def _require_nvjitlink_version(minimum_version: tuple[int, int], feature: str) -> None: + """Check that the cached nvJitLink runtime meets a feature's requirement.""" + if _nvjitlink_version < minimum_version: + required = ".".join(str(component) for component in minimum_version) + detected = ".".join(str(component) for component in _nvjitlink_version) + raise RuntimeError(f"{feature} requires nvJitLink {required} or newer; found {detected}") + + +# TODO(#2783): Replace this Python-level dispatch with direct cimports once +# the cuda-bindings runtime floor includes the linked-LTOIR getters. +def _linked_ltoir_output_module(): + """Return bindings that can retrieve linked LTOIR without a Cython dependency.""" + _require_nvjitlink_version((13, 3), "LTOIR output") + missing = [ + name + for name in ("get_linked_ltoir_size", "get_linked_ltoir") + if not hasattr(_nvjitlink, name) + ] + if missing: + raise RuntimeError( + "LTOIR output requires cuda-bindings with " + " and ".join(missing) + ) + return _nvjitlink + + def _nvjitlink_has_version_symbol(nvjitlink) -> bool: # This condition is equivalent to testing for version >= 12.3 return bool(nvjitlink._inspect_function_pointer("__nvJitLinkVersion")) @@ -718,10 +795,13 @@ def _nvjitlink_has_version_symbol(nvjitlink) -> bool: # Note: this function is reused in the tests def _decide_nvjitlink_or_driver() -> bool: """Return True if falling back to the cuLink* driver APIs.""" - global _driver, _use_nvjitlink_backend + global _driver, _nvjitlink, _nvjitlink_version, _use_nvjitlink_backend if _use_nvjitlink_backend is not None: return not _use_nvjitlink_backend + _nvjitlink = None + _nvjitlink_version = None + warn_txt_common = ( "the driver APIs will be used instead, which do not support" " minor version compatibility or linking LTO IRs." @@ -742,6 +822,9 @@ def _decide_nvjitlink_or_driver() -> bool: ) else: if has_version_symbol: + detected_version = nvjitlink_module.version() + _nvjitlink = nvjitlink_module + _nvjitlink_version = detected_version _use_nvjitlink_backend = True return False # Use nvjitlink warn_txt = ( diff --git a/cuda_core/cuda/core/_module.pyi b/cuda_core/cuda/core/_module.pyi index e73c47f471f..be128213e17 100644 --- a/cuda_core/cuda/core/_module.pyi +++ b/cuda_core/cuda/core/_module.pyi @@ -355,13 +355,13 @@ class ObjectCode: """ @staticmethod def from_object(module: bytes | str, *, name: str='', symbol_mapping: dict[str, str] | None=None) -> ObjectCode: - """Create an :class:`ObjectCode` instance from an existing object code. + """Create an :class:`ObjectCode` instance from a host object containing device code. Parameters ---------- module : bytes | str - Either a bytes object containing the in-memory object code to load, or - a file path string pointing to the on-disk object code to load. + Either a bytes object containing the in-memory host object to load, or + a file path string pointing to the on-disk host object to load. name : str | None A human-readable identifier representing this code object. symbol_mapping : dict | None diff --git a/cuda_core/cuda/core/_module.pyx b/cuda_core/cuda/core/_module.pyx index a350f14887f..b530a4225b7 100644 --- a/cuda_core/cuda/core/_module.pyx +++ b/cuda_core/cuda/core/_module.pyx @@ -741,13 +741,13 @@ cdef class ObjectCode: @staticmethod def from_object(module: bytes | str, *, name: str = "", symbol_mapping: dict[str, str] | None = None) -> ObjectCode: - """Create an :class:`ObjectCode` instance from an existing object code. + """Create an :class:`ObjectCode` instance from a host object containing device code. Parameters ---------- module : bytes | str - Either a bytes object containing the in-memory object code to load, or - a file path string pointing to the on-disk object code to load. + Either a bytes object containing the in-memory host object to load, or + a file path string pointing to the on-disk host object to load. name : str | None A human-readable identifier representing this code object. symbol_mapping : dict | None diff --git a/cuda_core/cuda/core/typing.py b/cuda_core/cuda/core/typing.py index 1bf9bb7c0d2..6471608c033 100644 --- a/cuda_core/cuda/core/typing.py +++ b/cuda_core/cuda/core/typing.py @@ -84,7 +84,7 @@ class ObjectCodeFormatType(StrEnum): * ``CUBIN`` — device-native CUDA binary. * ``LTOIR`` — LTO (link-time optimization) IR for later linking. * ``FATBIN`` — fat binary bundling multiple device images. - * ``OBJECT`` — relocatable device object. + * ``OBJECT`` — host object containing device code. * ``LIBRARY`` — device code library. """ diff --git a/cuda_core/docs/source/release/1.3.0-notes.rst b/cuda_core/docs/source/release/1.3.0-notes.rst index 6280f127285..dc5d59de966 100644 --- a/cuda_core/docs/source/release/1.3.0-notes.rst +++ b/cuda_core/docs/source/release/1.3.0-notes.rst @@ -9,6 +9,15 @@ New features ------------ +- Added incremental linking to :class:`Linker` through + ``LinkerOptions(incremental=True)`` on nvJitLink 13.2 or newer. A partial + native result is returned as a CUBIN and can be passed directly to another + linker. :meth:`Linker.link` also accepts ``"ltoir"`` on nvJitLink 13.3 or + newer so incremental LTO chains can retain IR until their final link. Direct + PTX and CUBIN inputs are rejected for this target because they carry no LTOIR + and nvJitLink would otherwise silently omit them. + (`#2369 `__) + - Added the read-only :attr:`ManagedBuffer.last_prefetch_location` property, which reports the destination requested by the most recent explicit prefetch across the buffer. It returns a :class:`Device` or :class:`Host`, diff --git a/cuda_core/tests/test_linker.py b/cuda_core/tests/test_linker.py index c80071fa405..4e57c843b3f 100644 --- a/cuda_core/tests/test_linker.py +++ b/cuda_core/tests/test_linker.py @@ -5,9 +5,20 @@ import inspect import warnings +import numpy as np import pytest -from cuda.core import Device, Linker, LinkerOptions, Program, ProgramOptions, _linker +from cuda.core import ( + Device, + LaunchConfig, + LegacyPinnedMemoryResource, + Linker, + LinkerOptions, + Program, + ProgramOptions, + _linker, + launch, +) from cuda.core._module import ObjectCode from cuda.core._program import _can_load_generated_ptx from cuda.core._utils.cuda_utils import CUDAError @@ -27,7 +38,11 @@ from cuda.bindings import nvjitlink nvJitLinkError = nvjitlink.nvJitLinkError + nvjitlink_version = nvjitlink.version() + has_linked_ltoir_bindings = all(hasattr(nvjitlink, name) for name in ("get_linked_ltoir_size", "get_linked_ltoir")) else: + nvjitlink_version = (0, 0) + has_linked_ltoir_bindings = False class nvJitLinkError(Exception): pass @@ -86,8 +101,7 @@ def compile_ltoir_functions(init_cuda): LinkerOptions(arch=ARCH, variables_used=["var1", "var2"]), LinkerOptions(arch=ARCH, variables_used=("var1", "var2")), ] - version = nvjitlink.version() - if version >= (12, 5): + if nvjitlink_version >= (12, 5): options.append(LinkerOptions(arch=ARCH, no_cache=True)) @@ -198,6 +212,14 @@ def test_linker_options_as_bytes_nvjitlink(): assert "-maxrregcount=32" in options_str +@pytest.mark.agent_authored(model="gpt-5.6") +@pytest.mark.skipif(is_culink_backend, reason="as_bytes() only supported for nvjitlink backend") +@pytest.mark.parametrize("value,expected_count", [(None, 0), (False, 0), (True, 1)]) +def test_linker_options_incremental_as_bytes(value, expected_count): + options = LinkerOptions(arch="sm_80", incremental=value) + assert options.as_bytes().count(b"-r") == expected_count + + @pytest.mark.parametrize("backend", ("invalid", "driver")) def test_linker_options_as_bytes_invalid_backend(backend): """Test LinkerOptions.as_bytes() with invalid backend""" @@ -320,11 +342,37 @@ def fake_decide(): assert result == "nvJitLink" assert called, "_decide_nvjitlink_or_driver was not called" + @pytest.mark.agent_authored(model="gpt-5.6") + def test_which_backend_caches_nvjitlink_module_and_version(self, monkeypatch): + class NvJitLink: + version_calls = 0 + + @classmethod + def version(cls): + cls.version_calls += 1 + return (13, 4) + + monkeypatch.setattr(_linker, "_use_nvjitlink_backend", None) + monkeypatch.setattr(_linker, "_nvjitlink", None) + monkeypatch.setattr(_linker, "_nvjitlink_version", None) + monkeypatch.setattr(_linker, "_optional_cuda_import", lambda _name: NvJitLink) + monkeypatch.setattr(_linker, "_nvjitlink_has_version_symbol", lambda _nvjitlink: True) + + assert Linker.which_backend() == "nvJitLink" + assert _linker._nvjitlink is NvJitLink + assert _linker._nvjitlink_version == (13, 4) + assert NvJitLink.version_calls == 1 + + assert Linker.which_backend() == "nvJitLink" + assert NvJitLink.version_calls == 1 + @pytest.mark.agent_authored(model="grok-4.5") def test_which_backend_falls_back_when_nvjitlink_too_old(self, monkeypatch): """Regression test for #2408: old nvJitLink must not crash which_backend().""" monkeypatch.setattr(_linker, "_use_nvjitlink_backend", None) monkeypatch.setattr(_linker, "_driver", None) + monkeypatch.setattr(_linker, "_nvjitlink", None) + monkeypatch.setattr(_linker, "_nvjitlink_version", None) def fake__optional_cuda_import(modname, probe_function=None): assert modname == "cuda.bindings.nvjitlink" @@ -346,6 +394,8 @@ def test_which_backend_falls_back_when_dylib_missing(self, monkeypatch): monkeypatch.setattr(_linker, "_use_nvjitlink_backend", None) monkeypatch.setattr(_linker, "_driver", None) + monkeypatch.setattr(_linker, "_nvjitlink", None) + monkeypatch.setattr(_linker, "_nvjitlink_version", None) def raise_missing(_nvjitlink): raise DynamicLibNotFoundError("missing") @@ -453,6 +503,13 @@ def test_prepare_driver_options_unsupported_raises(driver_binding, kwargs, match opts._prepare_driver_options() +@pytest.mark.agent_authored(model="gpt-5.6") +def test_prepare_driver_options_rejects_incremental(driver_binding): + options = LinkerOptions(incremental=True) + with pytest.raises(ValueError, match="incremental option is not supported by the driver API"): + options._prepare_driver_options() + + @pytest.mark.agent_authored(model="claude-opus-5") @pytest.mark.parametrize("value", [True, False]) def test_numba_debug_warns_and_is_ignored(value): @@ -500,3 +557,247 @@ def test_as_bytes_nvjitlink_unavailable(monkeypatch): opts = LinkerOptions(arch="sm_80") with pytest.raises(RuntimeError, match="nvJitLink backend is not available"): opts.as_bytes("nvjitlink") + + +@pytest.mark.agent_authored(model="gpt-5.6") +def test_require_nvjitlink_version_reports_required_and_detected_versions(monkeypatch): + monkeypatch.setattr(_linker, "_nvjitlink_version", (13, 1)) + + with pytest.raises(RuntimeError, match=r"requires nvJitLink 13\.2 or newer; found 13\.1"): + _linker._require_nvjitlink_version((13, 2), "incremental linking") + + +@pytest.mark.agent_authored(model="gpt-5.6") +def test_require_nvjitlink_version_accepts_boundary_version(monkeypatch): + monkeypatch.setattr(_linker, "_nvjitlink_version", (13, 2)) + + _linker._require_nvjitlink_version((13, 2), "incremental linking") + + +@pytest.mark.agent_authored(model="gpt-5.6") +def test_linked_ltoir_output_requires_new_enough_runtime(monkeypatch): + monkeypatch.setattr(_linker, "_nvjitlink_version", (13, 2)) + + with pytest.raises(RuntimeError, match=r"LTOIR output requires nvJitLink 13\.3 or newer; found 13\.2"): + _linker._linked_ltoir_output_module() + + +@pytest.mark.agent_authored(model="gpt-5.6") +def test_linked_ltoir_output_requires_new_enough_bindings(monkeypatch): + class NvJitLinkWithoutLinkedLtoir: + pass + + monkeypatch.setattr(_linker, "_nvjitlink", NvJitLinkWithoutLinkedLtoir) + monkeypatch.setattr(_linker, "_nvjitlink_version", (13, 3)) + + with pytest.raises(RuntimeError, match="cuda-bindings with get_linked_ltoir_size and get_linked_ltoir"): + _linker._linked_ltoir_output_module() + + +incremental_caller = r""" +extern "C" __device__ int incremental_helper(); +extern "C" __global__ void incremental_kernel(int* result) { + if (threadIdx.x == 0 && blockIdx.x == 0) { + *result = incremental_helper(); + } +} +""" + +incremental_helper = r""" +extern "C" __device__ int incremental_helper() { return 42; } +""" + + +def _compile_incremental_inputs(target_type): + if target_type == "ptx": + options = ProgramOptions(relocatable_device_code=True) + else: + options = ProgramOptions(link_time_optimization=True) + caller = Program(incremental_caller, "c++", options).compile(target_type) + helper = Program(incremental_helper, "c++", options).compile(target_type) + return caller, helper + + +def _launch_incrementally_linked_kernel(device, linked_code): + kernel = linked_code.get_kernel("incremental_kernel") + stream = device.create_stream() + try: + with LegacyPinnedMemoryResource().allocate(4) as host_buffer: + result = np.from_dlpack(host_buffer).view(np.int32) + try: + with device.memory_resource.allocate(4, stream=stream) as device_buffer: + result[:] = 0 + + launch(stream, LaunchConfig(grid=1, block=1), kernel, device_buffer) + device_buffer.copy_to(host_buffer, stream=stream) + stream.sync() + actual = int(result[0]) + finally: + # Drop the DLPack view before releasing its pinned allocation. + result = None + + assert actual == 42 + finally: + try: + stream.sync() + finally: + stream.close() + + +@pytest.mark.human_reviewed +@pytest.mark.skipif( + is_culink_backend or nvjitlink_version < (13, 2), + reason="incremental linking requires nvJitLink 13.2 or newer", +) +def test_incremental_cubin_round_trip(init_cuda): + caller, helper = _compile_incremental_inputs("ptx") + + partial = Linker(caller, options=LinkerOptions(arch=ARCH, incremental=True)).link("cubin") + assert partial.code_type == "cubin" + + resolved_partial = Linker( + partial, + helper, + options=LinkerOptions(arch=ARCH, incremental=True), + ).link("cubin") + assert resolved_partial.code_type == "cubin" + _launch_incrementally_linked_kernel(init_cuda, resolved_partial) + + final = Linker(resolved_partial, options=LinkerOptions(arch=ARCH)).link("cubin") + _launch_incrementally_linked_kernel(init_cuda, final) + + +@pytest.mark.agent_authored(model="gpt-5.6") +@pytest.mark.skipif( + is_culink_backend or nvjitlink_version < (13, 3) or not has_linked_ltoir_bindings, + reason="linked LTOIR output requires nvJitLink 13.3 or newer and matching cuda-bindings", +) +def test_incremental_ltoir_round_trip(init_cuda): + caller, helper = _compile_incremental_inputs("ltoir") + incremental_options = LinkerOptions( + arch=ARCH, + incremental=True, + link_time_optimization=True, + ) + + partial = Linker(caller, options=incremental_options).link("ltoir") + assert partial.code_type == "ltoir" + + resolved_partial = Linker(partial, helper, options=incremental_options).link("ltoir") + assert resolved_partial.code_type == "ltoir" + + final = Linker( + resolved_partial, + options=LinkerOptions(arch=ARCH, link_time_optimization=True), + ).link("cubin") + _launch_incrementally_linked_kernel(init_cuda, final) + + +@pytest.mark.agent_authored(model="gpt-5.6") +@pytest.mark.skipif( + is_culink_backend or nvjitlink_version < (13, 3) or not has_linked_ltoir_bindings, + reason="linked LTOIR output requires nvJitLink 13.3 or newer and matching cuda-bindings", +) +def test_complete_ltoir_round_trip_matches_direct_cubin(init_cuda): + caller, helper = _compile_incremental_inputs("ltoir") + options = LinkerOptions(arch=ARCH, link_time_optimization=True) + + linked_ltoir = Linker(caller, helper, options=options).link("ltoir") + direct_cubin = Linker(caller, helper, options=options).link("cubin") + round_trip_cubin = Linker(linked_ltoir, options=options).link("cubin") + + assert linked_ltoir.code_type == "ltoir" + assert round_trip_cubin.code == direct_cubin.code + _launch_incrementally_linked_kernel(init_cuda, round_trip_cubin) + + +@pytest.mark.agent_authored(model="gpt-5.6") +@pytest.mark.skipif( + is_culink_backend or nvjitlink_version < (13, 2), + reason="incremental linking requires nvJitLink 13.2 or newer", +) +def test_incremental_lto_cubin_round_trip(init_cuda): + caller, helper = _compile_incremental_inputs("ltoir") + partial = Linker( + caller, + options=LinkerOptions( + arch=ARCH, + incremental=True, + link_time_optimization=True, + ), + ).link("cubin") + + assert partial.code_type == "cubin" + assert partial.code.startswith(b"\x7fELF") + assert int.from_bytes(partial.code[16:18], "little") == 1 # ET_REL + + final = Linker( + partial, + helper, + options=LinkerOptions(arch=ARCH, link_time_optimization=True), + ).link("cubin") + _launch_incrementally_linked_kernel(init_cuda, final) + + +@pytest.mark.agent_authored(model="gpt-5.6") +@pytest.mark.skipif( + is_culink_backend or nvjitlink_version < (13, 3) or not has_linked_ltoir_bindings, + reason="linked LTOIR output requires nvJitLink 13.3 or newer and matching cuda-bindings", +) +@pytest.mark.parametrize("non_ltoir_type", ("ptx", "cubin")) +def test_ltoir_output_rejects_inputs_without_ltoir(init_cuda, non_ltoir_type): + caller, _ = _compile_incremental_inputs("ltoir") + other_kernel = 'extern "C" __global__ void other_kernel() {}' + non_ltoir_input = Program( + other_kernel, + "c++", + ProgramOptions(relocatable_device_code=True), + ).compile(non_ltoir_type) + linker = Linker( + caller, + non_ltoir_input, + options=LinkerOptions( + arch=ARCH, + incremental=True, + link_time_optimization=True, + ), + ) + + with pytest.raises(ValueError, match='LTOIR output is not supported with "ptx" or "cubin" inputs'): + linker.link("ltoir") + + +@pytest.mark.agent_authored(model="gpt-5.6") +@pytest.mark.skipif( + is_culink_backend or nvjitlink_version < (13, 2), + reason="incremental linking requires nvJitLink 13.2 or newer", +) +def test_incremental_link_rejects_ptx_output(compile_ptx_functions): + linker = Linker( + *compile_ptx_functions, + options=LinkerOptions(arch=ARCH, incremental=True), + ) + with pytest.raises(ValueError, match="PTX output is not supported for incremental linking"): + linker.link("ptx") + + +@pytest.mark.agent_authored(model="gpt-5.6") +def test_incremental_link_rejects_ptx_option(compile_ptx_functions): + with pytest.raises(ValueError, match="incremental and ptx output options cannot be used together"): + Linker( + *compile_ptx_functions, + options=LinkerOptions( + arch=ARCH, + incremental=True, + link_time_optimization=True, + ptx=True, + ), + ) + + +@pytest.mark.agent_authored(model="gpt-5.6") +@pytest.mark.skipif(is_culink_backend, reason="LTOIR output requires nvJitLink") +def test_ltoir_output_requires_lto(compile_ptx_functions): + linker = Linker(*compile_ptx_functions, options=LinkerOptions(arch=ARCH)) + with pytest.raises(ValueError, match="LTOIR output requires link_time_optimization=True"): + linker.link("ltoir") diff --git a/cuda_core/tests/test_optional_dependency_imports.py b/cuda_core/tests/test_optional_dependency_imports.py index b08b7d344d9..c5dff38116a 100644 --- a/cuda_core/tests/test_optional_dependency_imports.py +++ b/cuda_core/tests/test_optional_dependency_imports.py @@ -18,12 +18,16 @@ def restore_optional_import_state(): saved_nvvm_attempted = _program._nvvm_import_attempted saved_driver = _linker._driver saved_inited = _linker._inited + saved_nvjitlink = _linker._nvjitlink + saved_nvjitlink_version = _linker._nvjitlink_version saved_use_nvjitlink = _linker._use_nvjitlink_backend _program._nvvm_module = None _program._nvvm_import_attempted = False _linker._driver = None _linker._inited = False + _linker._nvjitlink = None + _linker._nvjitlink_version = None _linker._use_nvjitlink_backend = None yield @@ -32,6 +36,8 @@ def restore_optional_import_state(): _program._nvvm_import_attempted = saved_nvvm_attempted _linker._driver = saved_driver _linker._inited = saved_inited + _linker._nvjitlink = saved_nvjitlink + _linker._nvjitlink_version = saved_nvjitlink_version _linker._use_nvjitlink_backend = saved_use_nvjitlink @@ -168,12 +174,22 @@ def fake__optional_cuda_import(modname, probe_function=None): assert _linker._use_nvjitlink_backend is False -@pytest.mark.agent_authored(model="grok-4.5") +@pytest.mark.agent_authored(model="gpt-5.6") def test_decide_nvjitlink_or_driver_selects_nvjitlink_when_version_symbol_present(monkeypatch): + version_calls = 0 + + class FakeModule: + def version(self): + nonlocal version_calls + version_calls += 1 + return (13, 4) + + nvjitlink_module = FakeModule() + def fake__optional_cuda_import(modname, probe_function=None): assert modname == "cuda.bindings.nvjitlink" assert probe_function is None - return object() + return nvjitlink_module monkeypatch.setattr(_linker, "_optional_cuda_import", fake__optional_cuda_import) monkeypatch.setattr(_linker, "_nvjitlink_has_version_symbol", lambda _nvjitlink: True) @@ -182,21 +198,24 @@ def fake__optional_cuda_import(modname, probe_function=None): assert use_driver_backend is False assert _linker._use_nvjitlink_backend is True + assert _linker._nvjitlink is nvjitlink_module + assert _linker._nvjitlink_version == (13, 4) + assert version_calls == 1 -@pytest.mark.agent_authored(model="grok-4.5") -def test_decide_nvjitlink_or_driver_does_not_call_version(monkeypatch): - """Regression guard for #2408: must not call module.version().""" +@pytest.mark.agent_authored(model="gpt-5.6") +def test_decide_nvjitlink_or_driver_does_not_call_version_when_symbol_missing(monkeypatch): + """Regression guard for #2408: old nvJitLink must not call module.version().""" called = {"version": False, "inspect": False} class FakeModule: def version(self): called["version"] = True - raise AssertionError("module.version() must not be used for nvJitLink probing") + raise AssertionError("module.version() must not be called when its symbol is missing") def fake_has_version(_nvjitlink): called["inspect"] = True - return True + return False def fake__optional_cuda_import(modname, probe_function=None): assert modname == "cuda.bindings.nvjitlink" @@ -206,6 +225,9 @@ def fake__optional_cuda_import(modname, probe_function=None): monkeypatch.setattr(_linker, "_optional_cuda_import", fake__optional_cuda_import) monkeypatch.setattr(_linker, "_nvjitlink_has_version_symbol", fake_has_version) - assert _linker._decide_nvjitlink_or_driver() is False + with pytest.warns(RuntimeWarning, match="too old \\(<12.3\\)"): + assert _linker._decide_nvjitlink_or_driver() is True assert called["inspect"] is True assert called["version"] is False + assert _linker._nvjitlink is None + assert _linker._nvjitlink_version is None