From 60170774b8582f0a45bac0ea3f5040c0433b2d25 Mon Sep 17 00:00:00 2001 From: Michael Droettboom Date: Tue, 21 Jul 2026 12:37:56 -0400 Subject: [PATCH] New API for nvrtc --- cuda_bindings/AGENTS.md | 5 + cuda_bindings/build_hooks.py | 2 +- .../cuda/bindings/_example_helpers/common.py | 22 +- cuda_bindings/cuda/bindings/_v2/__init__.py | 0 cuda_bindings/cuda/bindings/_v2/nvrtc.pxd | 60 + cuda_bindings/cuda/bindings/_v2/nvrtc.pyx | 1049 +++++++++++++++++ .../docs/source/release/13.4.0-notes.rst | 8 + cuda_bindings/examples/extra/jit_program.py | 33 +- cuda_bindings/tests/legacy_api/README.md | 4 + .../legacy_api/test_legacy_kernelParams.py | 802 +++++++++++++ .../tests/legacy_api/test_legacy_nvfatbin.py | 326 +++++ .../tests/legacy_api/test_legacy_nvjitlink.py | 210 ++++ .../tests/legacy_api/test_legacy_nvrtc.py | 37 + cuda_bindings/tests/test_kernelParams.py | 32 +- cuda_bindings/tests/test_nvfatbin.py | 42 +- cuda_bindings/tests/test_nvjitlink.py | 22 +- cuda_bindings/tests/test_nvrtc.py | 38 +- 17 files changed, 2552 insertions(+), 140 deletions(-) create mode 100644 cuda_bindings/cuda/bindings/_v2/__init__.py create mode 100644 cuda_bindings/cuda/bindings/_v2/nvrtc.pxd create mode 100644 cuda_bindings/cuda/bindings/_v2/nvrtc.pyx create mode 100644 cuda_bindings/tests/legacy_api/README.md create mode 100644 cuda_bindings/tests/legacy_api/test_legacy_kernelParams.py create mode 100644 cuda_bindings/tests/legacy_api/test_legacy_nvfatbin.py create mode 100644 cuda_bindings/tests/legacy_api/test_legacy_nvjitlink.py create mode 100644 cuda_bindings/tests/legacy_api/test_legacy_nvrtc.py diff --git a/cuda_bindings/AGENTS.md b/cuda_bindings/AGENTS.md index 12267ad12d0..8c544f8872a 100644 --- a/cuda_bindings/AGENTS.md +++ b/cuda_bindings/AGENTS.md @@ -42,6 +42,11 @@ subpackage in the `cuda-python` monorepo. - **Examples**: example coverage is pytest-based under `examples/`. - **Benchmarks**: run with `pytest --benchmark-only benchmarks/` when needed. +The `legacy_tests` subdirectory tests the old pre-v2 APIs of `driver`, `runtime` +and `nvrtc`. These test files should not be added to, only updated when +necessary to fix test failures. The canonical set of tests are those outside of +the `legacy_tests` subdirectory. + ## Build and environment notes - `CUDA_HOME` or `CUDA_PATH` must point to a valid CUDA Toolkit for source diff --git a/cuda_bindings/build_hooks.py b/cuda_bindings/build_hooks.py index f2a73b07218..a50133f9777 100644 --- a/cuda_bindings/build_hooks.py +++ b/cuda_bindings/build_hooks.py @@ -196,7 +196,7 @@ def _cleanup_dst_files(): # Build extension list extensions = [] - cuda_bindings_files = glob.glob("cuda/bindings/*.pyx") + cuda_bindings_files = glob.glob("cuda/bindings/*.pyx") + glob.glob("cuda/bindings/_v2/*.pyx") if sys.platform == "win32": cuda_bindings_files = [f for f in cuda_bindings_files if "cufile" not in f] diff --git a/cuda_bindings/cuda/bindings/_example_helpers/common.py b/cuda_bindings/cuda/bindings/_example_helpers/common.py index 8970944a31a..6335d3e3e4a 100644 --- a/cuda_bindings/cuda/bindings/_example_helpers/common.py +++ b/cuda_bindings/cuda/bindings/_example_helpers/common.py @@ -9,8 +9,8 @@ from cuda import pathfinder from cuda.bindings import driver as cuda -from cuda.bindings import nvrtc from cuda.bindings import runtime as cudart +from cuda.bindings._v2 import nvrtc from .helper_cuda import check_cuda_errors @@ -44,7 +44,7 @@ def __init__(self, code, dev_id): requirement_not_met(f'pathfinder.find_nvidia_header_directory("{libname}") returned None') include_dirs.append(hdr_dir) - prog = check_cuda_errors(nvrtc.nvrtcCreateProgram(str.encode(code), b"sourceCode.cu", 0, None, None)) + prog = nvrtc.create_program(str.encode(code), b"sourceCode.cu") # Initialize CUDA check_cuda_errors(cudart.cudaFree(0)) @@ -55,7 +55,7 @@ def __init__(self, code, dev_id): minor = check_cuda_errors( cudart.cudaDeviceGetAttribute(cudart.cudaDeviceAttr.cudaDevAttrComputeCapabilityMinor, dev_id) ) - _, nvrtc_minor = check_cuda_errors(nvrtc.nvrtcVersion()) + _, nvrtc_minor = nvrtc.version() use_cubin = nvrtc_minor >= 1 prefix = "sm" if use_cubin else "compute" arch_arg = bytes(f"--gpu-architecture={prefix}_{major}{minor}", "ascii") @@ -70,11 +70,9 @@ def __init__(self, code, dev_id): opts.append(f"--include-path={inc_dir}".encode()) try: - check_cuda_errors(nvrtc.nvrtcCompileProgram(prog, len(opts), opts)) - except RuntimeError as err: - log_size = check_cuda_errors(nvrtc.nvrtcGetProgramLogSize(prog)) - log = b" " * log_size - check_cuda_errors(nvrtc.nvrtcGetProgramLog(prog, log)) + nvrtc.compile_program(prog, opts) + except nvrtc.NvrtcError as err: + log = nvrtc.get_program_log(prog) import sys print(log.decode(), file=sys.stderr) # noqa: T201 @@ -82,13 +80,9 @@ def __init__(self, code, dev_id): sys.exit(1) if use_cubin: - data_size = check_cuda_errors(nvrtc.nvrtcGetCUBINSize(prog)) - data = b" " * data_size - check_cuda_errors(nvrtc.nvrtcGetCUBIN(prog, data)) + data = nvrtc.get_cubin(prog) else: - data_size = check_cuda_errors(nvrtc.nvrtcGetPTXSize(prog)) - data = b" " * data_size - check_cuda_errors(nvrtc.nvrtcGetPTX(prog, data)) + data = nvrtc.get_ptx(prog) self.module = check_cuda_errors(cuda.cuModuleLoadData(np.char.array(data))) diff --git a/cuda_bindings/cuda/bindings/_v2/__init__.py b/cuda_bindings/cuda/bindings/_v2/__init__.py new file mode 100644 index 00000000000..e69de29bb2d diff --git a/cuda_bindings/cuda/bindings/_v2/nvrtc.pxd b/cuda_bindings/cuda/bindings/_v2/nvrtc.pxd new file mode 100644 index 00000000000..1a0d1c811de --- /dev/null +++ b/cuda_bindings/cuda/bindings/_v2/nvrtc.pxd @@ -0,0 +1,60 @@ +# SPDX-FileCopyrightText: Copyright (c) 2025-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# +# SPDX-License-Identifier: Apache-2.0 +# +# This code was automatically generated across versions from 12.9.0 to 13.3.0. Do not modify it directly. + +# CYTHON-BINDINGS-GENERATED-DO-NOT-MODIFY-THIS-FILE: format=1; content-sha256=223716a3d6961dfe7231f2c88e76a9ab4c0a8d888cc4bf3e889bfbcbf8580346 +from libc.stdint cimport intptr_t + +from ..cynvrtc cimport * + + +############################################################################### +# Types +############################################################################### + +ctypedef nvrtcProgram Program + + +############################################################################### +# Enum +############################################################################### + +ctypedef nvrtcResult _Result + + +############################################################################### +# Functions +############################################################################### + +cpdef intptr_t create_program(bytes src, name, headers=*, include_names=*) except? 0 +cpdef compile_program(intptr_t prog, options=*) +cpdef set_flow_callback(intptr_t prog, intptr_t callback, intptr_t payload) +cpdef bytes get_lowered_name(intptr_t prog, bytes name_expression) +cpdef install_bundled_headers(bytes install_path, unsigned int flags) +cpdef tuple get_bundled_headers_info() +cpdef remove_bundled_headers(bytes install_path) + +cpdef str get_error_string(int result) +cpdef tuple version() +cpdef int get_num_supported_archs() except? -1 +cpdef object get_supported_archs() +cpdef destroy_program(intptr_t prog) +cpdef size_t get_ptx_size(intptr_t prog) except? 0 +cpdef bytes get_ptx(intptr_t prog) +cpdef size_t get_cubin_size(intptr_t prog) except? 0 +cpdef bytes get_cubin(intptr_t prog) +cpdef size_t get_ltoir_size(intptr_t prog) except? 0 +cpdef bytes get_ltoir(intptr_t prog) +cpdef size_t get_optix_ir_size(intptr_t prog) except? 0 +cpdef bytes get_optix_ir(intptr_t prog) +cpdef size_t get_program_log_size(intptr_t prog) except? 0 +cpdef bytes get_program_log(intptr_t prog) +cpdef add_name_expression(intptr_t prog, name_expression) +cpdef size_t get_pch_heap_size() except? 0 +cpdef set_pch_heap_size(size_t size) +cpdef int get_pch_create_status(intptr_t prog) except? -1 +cpdef size_t get_pch_heap_size_required(intptr_t prog) except? 0 +cpdef size_t get_tile_ir_size(intptr_t prog) except? 0 +cpdef bytes get_tile_ir(intptr_t prog) diff --git a/cuda_bindings/cuda/bindings/_v2/nvrtc.pyx b/cuda_bindings/cuda/bindings/_v2/nvrtc.pyx new file mode 100644 index 00000000000..379350ffa6f --- /dev/null +++ b/cuda_bindings/cuda/bindings/_v2/nvrtc.pyx @@ -0,0 +1,1049 @@ +# SPDX-FileCopyrightText: Copyright (c) 2025-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# +# SPDX-License-Identifier: Apache-2.0 +# +# This code was automatically generated across versions from 12.9.0 to 13.3.0. Do not modify it directly. +# CYTHON-BINDINGS-GENERATED-DO-NOT-MODIFY-THIS-FILE: format=1; content-sha256=e5360fc057cdd7b8e28b4d502821443d190e589b200dce1a234494ad6e9abf93 + + +# <<<< PREAMBLE CONTENT >>>> + +cimport cpython as _cyb_cpython +cimport cpython.buffer as _cyb_cpython_buffer +from cython cimport view as _cyb_view +from libc.stdlib cimport ( + calloc as _cyb_calloc, + free as _cyb_free, + malloc as _cyb_malloc, +) +from libc.string cimport ( + memcmp as _cyb_memcmp, + memcpy as _cyb_memcpy, +) + +from cuda.bindings._internal._fast_enum import FastEnum as _cyb_FastEnum + +import numpy as _numpy + +cdef _cyb___getbuffer(object self, _cyb_cpython.Py_buffer *buffer, void *ptr, int size, bint readonly): + buffer.buf = ptr + buffer.format = 'b' + buffer.internal = NULL + buffer.itemsize = 1 + buffer.len = size + buffer.ndim = 1 + buffer.obj = self + buffer.readonly = readonly + buffer.shape = &buffer.len + buffer.strides = &buffer.itemsize + buffer.suboffsets = NULL + +cdef _cyb_from_buffer(buffer, size, lowpp_type): + cdef _cyb_cpython.Py_buffer view + if _cyb_cpython.PyObject_GetBuffer(buffer, &view, _cyb_cpython_buffer.PyBUF_SIMPLE) != 0: + raise TypeError("buffer argument does not support the buffer protocol") + try: + if view.itemsize != 1: + raise ValueError("buffer itemsize must be 1 byte") + if view.len != size: + raise ValueError(f"buffer length must be {size} bytes") + return lowpp_type.from_ptr(view.buf, not view.readonly, buffer) + finally: + _cyb_cpython.PyBuffer_Release(&view) + +cdef _cyb_from_data(data, dtype_name, expected_dtype, lowpp_type): + # _numpy.recarray is a subclass of _numpy.ndarray, so implicitly handled here. + if isinstance(data, lowpp_type): + return data + if not isinstance(data, _numpy.ndarray): + raise TypeError("data argument must be a NumPy ndarray") + if data.size != 1: + raise ValueError("data array must have a size of 1") + if data.dtype != expected_dtype: + raise ValueError(f"data array must be of dtype {dtype_name}") + return lowpp_type.from_ptr(data.ctypes.data, not data.flags.writeable, data) + + +# <<<< END OF PREAMBLE CONTENT >>>> + +cimport cython # NOQA +from libcpp.vector cimport vector + +from cuda.bindings._internal._fast_enum import FastEnum as _FastEnum + + +############################################################################### +# Enum +############################################################################### + +class Result(_cyb_FastEnum): + """ + The enumerated type `nvrtcResult` defines API call result codes. NVRTC + API functions return `nvrtcResult` to indicate the call result. + + See `nvrtcResult`. + """ + SUCCESS = NVRTC_SUCCESS + ERROR_OUT_OF_MEMORY = NVRTC_ERROR_OUT_OF_MEMORY + ERROR_PROGRAM_CREATION_FAILURE = NVRTC_ERROR_PROGRAM_CREATION_FAILURE + ERROR_INVALID_INPUT = NVRTC_ERROR_INVALID_INPUT + ERROR_INVALID_PROGRAM = NVRTC_ERROR_INVALID_PROGRAM + ERROR_INVALID_OPTION = NVRTC_ERROR_INVALID_OPTION + ERROR_COMPILATION = NVRTC_ERROR_COMPILATION + ERROR_BUILTIN_OPERATION_FAILURE = NVRTC_ERROR_BUILTIN_OPERATION_FAILURE + ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION = NVRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION + ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION = NVRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION + ERROR_NAME_EXPRESSION_NOT_VALID = NVRTC_ERROR_NAME_EXPRESSION_NOT_VALID + ERROR_INTERNAL_ERROR = NVRTC_ERROR_INTERNAL_ERROR + ERROR_TIME_FILE_WRITE_FAILED = NVRTC_ERROR_TIME_FILE_WRITE_FAILED + ERROR_NO_PCH_CREATE_ATTEMPTED = NVRTC_ERROR_NO_PCH_CREATE_ATTEMPTED + ERROR_PCH_CREATE_HEAP_EXHAUSTED = NVRTC_ERROR_PCH_CREATE_HEAP_EXHAUSTED + ERROR_PCH_CREATE = NVRTC_ERROR_PCH_CREATE + ERROR_CANCELLED = NVRTC_ERROR_CANCELLED + ERROR_TIME_TRACE_FILE_WRITE_FAILED = NVRTC_ERROR_TIME_TRACE_FILE_WRITE_FAILED + ERROR_BUSY = NVRTC_ERROR_BUSY + + +class InstallHeadersFlag(_FastEnum): + """Flags for :func:`install_bundled_headers`.""" + SKIP_IF_EXISTS = ( + 0x0, + "Skip installation if version marker exists and version matches. " + "This is the default behavior when flags=0." + ) + FORCE_OVERWRITE = ( + 0x1, + "Clear existing directory contents before installation. " + "Guarantees consistency by removing any existing files first." + ) + NO_WAIT = ( + 0x2, + "Return NVRTC_ERROR_BUSY immediately if installation is in progress " + "by another process, instead of waiting for the lock. " + "Can be combined with FORCE_OVERWRITE using bitwise OR. " + "Do not wait for installation to complete." + ) + + +############################################################################### +# Error handling +############################################################################### + + +class NvrtcError(Exception): + def __init__(self, status): + self.status = status + s = get_error_string(status) + super(NvrtcError, self).__init__(s) + + def __reduce__(self): + return (type(self), (self.status,)) + +class OutOfMemoryError(NvrtcError): + pass +class ProgramCreationFailureError(NvrtcError): + pass +class InvalidInputError(NvrtcError): + pass +class InvalidProgramError(NvrtcError): + pass +class InvalidOptionError(NvrtcError): + pass +class CompilationError(NvrtcError): + pass +class BuiltinOperationFailureError(NvrtcError): + pass +class NoNameExpressionsAfterCompilationError(NvrtcError): + pass +class NoLoweredNamesBeforeCompilationError(NvrtcError): + pass +class NameExpressionNotValidError(NvrtcError): + pass +class InternalErrorError(NvrtcError): + pass +class TimeFileWriteFailedError(NvrtcError): + pass +class NoPchCreateAttemptedError(NvrtcError): + pass +class PchCreateHeapExhaustedError(NvrtcError): + pass +class PchCreateError(NvrtcError): + pass +class CancelledError(NvrtcError): + pass +class TimeTraceFileWriteFailedError(NvrtcError): + pass +class BusyError(NvrtcError): + pass +cdef object _nvrtc_error_factory(int status): + cdef object pystatus = status + if status == 1: + return OutOfMemoryError(pystatus) + elif status == 2: + return ProgramCreationFailureError(pystatus) + elif status == 3: + return InvalidInputError(pystatus) + elif status == 4: + return InvalidProgramError(pystatus) + elif status == 5: + return InvalidOptionError(pystatus) + elif status == 6: + return CompilationError(pystatus) + elif status == 7: + return BuiltinOperationFailureError(pystatus) + elif status == 8: + return NoNameExpressionsAfterCompilationError(pystatus) + elif status == 9: + return NoLoweredNamesBeforeCompilationError(pystatus) + elif status == 10: + return NameExpressionNotValidError(pystatus) + elif status == 11: + return InternalErrorError(pystatus) + elif status == 12: + return TimeFileWriteFailedError(pystatus) + elif status == 13: + return NoPchCreateAttemptedError(pystatus) + elif status == 14: + return PchCreateHeapExhaustedError(pystatus) + elif status == 15: + return PchCreateError(pystatus) + elif status == 16: + return CancelledError(pystatus) + elif status == 17: + return TimeTraceFileWriteFailedError(pystatus) + elif status == 18: + return BusyError(pystatus) + return NvrtcError(status) + +InternalError = InternalErrorError + + +@cython.profile(False) +cdef int check_status(int status) except 1 nogil: + if status != 0: + with gil: + raise _nvrtc_error_factory(status) + return status + + +############################################################################### +# POD definitions +############################################################################### + +cdef _get_bundled_headers_info_dtype_offsets(): + cdef nvrtcBundledHeadersInfo pod + return _numpy.dtype({ + 'names': ['available', 'compressed_size', 'uncompressed_size', 'cuda_version_major', 'cuda_version_minor', 'num_files'], + 'formats': [_numpy.int32, _numpy.uint64, _numpy.uint64, _numpy.int32, _numpy.int32, _numpy.uint32], + 'offsets': [ + (&(pod.available)) - (&pod), + (&(pod.compressedSize)) - (&pod), + (&(pod.uncompressedSize)) - (&pod), + (&(pod.cudaVersionMajor)) - (&pod), + (&(pod.cudaVersionMinor)) - (&pod), + (&(pod.numFiles)) - (&pod), + ], + 'itemsize': sizeof(nvrtcBundledHeadersInfo), + }) + +bundled_headers_info_dtype = _get_bundled_headers_info_dtype_offsets() + +cdef class BundledHeadersInfo: + """Empty-initialize an instance of `nvrtcBundledHeadersInfo`. + + + .. seealso:: `nvrtcBundledHeadersInfo` + """ + cdef: + nvrtcBundledHeadersInfo *_ptr + object _owner + bint _owned + bint _readonly + + def __init__(self): + self._ptr = _cyb_calloc(1, sizeof(nvrtcBundledHeadersInfo)) + if self._ptr == NULL: + raise MemoryError("Error allocating BundledHeadersInfo") + self._owner = None + self._owned = True + self._readonly = False + + def __dealloc__(self): + cdef nvrtcBundledHeadersInfo *ptr + if self._owned and self._ptr != NULL: + ptr = self._ptr + self._ptr = NULL + _cyb_free(ptr) + + def __repr__(self): + return f"<{__name__}.BundledHeadersInfo object at {hex(id(self))}>" + + @property + def ptr(self): + """Get the pointer address to the data as Python :class:`int`.""" + return (self._ptr) + + cdef intptr_t _get_ptr(self): + return (self._ptr) + + def __int__(self): + return (self._ptr) + + def __eq__(self, other): + cdef BundledHeadersInfo other_ + if not isinstance(other, BundledHeadersInfo): + return False + other_ = other + return (_cyb_memcmp((self._ptr), (other_._ptr), sizeof(nvrtcBundledHeadersInfo)) == 0) + + def __getbuffer__(self, _cyb_cpython.Py_buffer *buffer, int flags): + _cyb___getbuffer(self, buffer, self._ptr, sizeof(nvrtcBundledHeadersInfo), self._readonly) + + def __releasebuffer__(self, Py_buffer *buffer): + pass + + def __setitem__(self, key, val): + if key == 0 and isinstance(val, _numpy.ndarray): + self._ptr = _cyb_malloc(sizeof(nvrtcBundledHeadersInfo)) + if self._ptr == NULL: + raise MemoryError("Error allocating BundledHeadersInfo") + _cyb_memcpy(self._ptr, val.ctypes.data, sizeof(nvrtcBundledHeadersInfo)) + self._owner = None + self._owned = True + self._readonly = not val.flags.writeable + else: + setattr(self, key, val) + + @property + def available(self): + """int: Non-zero if bundled headers are available""" + return self._ptr[0].available + + @available.setter + def available(self, val): + if self._readonly: + raise ValueError("This BundledHeadersInfo instance is read-only") + self._ptr[0].available = val + + @property + def compressed_size(self): + """int: Size of compressed archive in bytes""" + return self._ptr[0].compressedSize + + @compressed_size.setter + def compressed_size(self, val): + if self._readonly: + raise ValueError("This BundledHeadersInfo instance is read-only") + self._ptr[0].compressedSize = val + + @property + def uncompressed_size(self): + """int: Estimated size when extracted in bytes""" + return self._ptr[0].uncompressedSize + + @uncompressed_size.setter + def uncompressed_size(self, val): + if self._readonly: + raise ValueError("This BundledHeadersInfo instance is read-only") + self._ptr[0].uncompressedSize = val + + @property + def cuda_version_major(self): + """int: CUDA major version of bundled headers""" + return self._ptr[0].cudaVersionMajor + + @cuda_version_major.setter + def cuda_version_major(self, val): + if self._readonly: + raise ValueError("This BundledHeadersInfo instance is read-only") + self._ptr[0].cudaVersionMajor = val + + @property + def cuda_version_minor(self): + """int: CUDA minor version of bundled headers""" + return self._ptr[0].cudaVersionMinor + + @cuda_version_minor.setter + def cuda_version_minor(self, val): + if self._readonly: + raise ValueError("This BundledHeadersInfo instance is read-only") + self._ptr[0].cudaVersionMinor = val + + @property + def num_files(self): + """int: Number of header files in the bundle""" + return self._ptr[0].numFiles + + @num_files.setter + def num_files(self, val): + if self._readonly: + raise ValueError("This BundledHeadersInfo instance is read-only") + self._ptr[0].numFiles = val + + @staticmethod + def from_buffer(buffer): + """Create an BundledHeadersInfo instance with the memory from the given buffer.""" + return _cyb_from_buffer(buffer, sizeof(nvrtcBundledHeadersInfo), BundledHeadersInfo) + + @staticmethod + def from_data(data): + """Create an BundledHeadersInfo instance wrapping the given NumPy array. + + Args: + data (_numpy.ndarray): a single-element array of dtype `bundled_headers_info_dtype` holding the data. + """ + return _cyb_from_data(data, "bundled_headers_info_dtype", bundled_headers_info_dtype, BundledHeadersInfo) + + @staticmethod + def from_ptr(intptr_t ptr, bint readonly=False, object owner=None): + """Create an BundledHeadersInfo instance wrapping the given pointer. + + Args: + ptr (intptr_t): pointer address as Python :class:`int` to the data. + owner (object): The Python object that owns the pointer. If not provided, data will be copied. + readonly (bool): whether the data is read-only (to the user). default is `False`. + """ + if ptr == 0: + raise ValueError("ptr must not be null (0)") + cdef BundledHeadersInfo obj = BundledHeadersInfo.__new__(BundledHeadersInfo) + if owner is None: + obj._ptr = _cyb_malloc(sizeof(nvrtcBundledHeadersInfo)) + if obj._ptr == NULL: + raise MemoryError("Error allocating BundledHeadersInfo") + _cyb_memcpy((obj._ptr), ptr, sizeof(nvrtcBundledHeadersInfo)) + obj._owner = None + obj._owned = True + else: + obj._ptr = ptr + obj._owner = owner + obj._owned = False + obj._readonly = readonly + return obj + + +############################################################################### +# Wrapper functions +############################################################################### + +cpdef intptr_t create_program(bytes src, name, headers=None, include_names=None) except? 0: + """nvrtcCreateProgram creates an instance of nvrtcProgram with the given input parameters. + + Args: + src (bytes): CUDA program source. + name (bytes | None): CUDA program name. ``None`` or ``""`` causes ``"default_program"`` + to be used. + headers (list[bytes] | None): Sources of the headers. ``None`` is treated as an empty + list (no headers). + include_names (list[bytes] | None): Name of each header by which it can be included in + the CUDA program source. Must have the same length as *headers*. + + Returns: + intptr_t: Opaque handle to the created program. Pass to :func:`destroy_program` when done. + + .. seealso:: `nvrtcCreateProgram` + """ + if headers is None: + headers = [] + if include_names is None: + include_names = [] + if len(headers) != len(include_names): + raise ValueError( + f"headers and include_names must have the same length " + f"({len(headers)} != {len(include_names)})" + ) + cdef int num_headers = len(headers) + cdef Program prog + cdef const char* c_src = src + cdef bytes _name = b"" if name is None else name + cdef const char* c_name = _name + cdef vector[const char*] cy_headers = headers + cdef vector[const char*] cy_include_names = include_names + cdef const char** hdr_data = NULL + cdef const char** inc_data = NULL + if num_headers: + hdr_data = cy_headers.data() + inc_data = cy_include_names.data() + with nogil: + __status__ = nvrtcCreateProgram(&prog, c_src, c_name, num_headers, hdr_data, inc_data) + check_status(__status__) + return prog + + +cpdef compile_program(intptr_t prog, options=None): + """nvrtcCompileProgram compiles the given program. + + It supports compile options listed in Supported Compile Options. + + Args: + prog (intptr_t): CUDA Runtime Compilation program. + options (list[bytes] | None): Compiler options as a list of byte strings. + May be ``None`` or an empty list for no options. + + .. seealso:: `nvrtcCompileProgram` + """ + if options is None: + options = [] + cdef int num_options = len(options) + cdef vector[const char*] cy_options = options + cdef const char** opt_data = NULL + if num_options: + opt_data = cy_options.data() + with nogil: + __status__ = nvrtcCompileProgram(prog, num_options, opt_data) + check_status(__status__) + + +cpdef set_flow_callback(intptr_t prog, intptr_t callback, intptr_t payload): + """nvrtcSetFlowCallback registers a callback that the compiler invokes during :func:`compile_program`. + + The callback signature must be ``int callback(void *param1, void *param2)``. + The compiler passes *payload* as *param1* and ``NULL`` as *param2* (reserved). + Return 1 to cancel compilation, 0 to continue; the callback must return + consistently, be thread-safe, and must not call any NVRTC/libnvvm/PTX APIs. + + Pass *callback* as a raw C function-pointer integer (e.g. via ``ctypes.cast``). + Pass 0 for *callback* to clear a previously registered callback. + + Args: + prog (intptr_t): CUDA Runtime Compilation program. + callback (intptr_t): C function pointer ``int (*)(void*, void*)`` cast to an integer, + or 0 to clear. + payload (intptr_t): Opaque pointer passed to the callback as its first argument. + + .. seealso:: `nvrtcSetFlowCallback` + """ + with nogil: + __status__ = nvrtcSetFlowCallback(prog, callback, payload) + check_status(__status__) + + +cpdef bytes get_lowered_name(intptr_t prog, bytes name_expression): + """nvrtcGetLoweredName extracts the lowered (mangled) name for a ``__global__`` function or ``__device__``/``__constant__`` variable. + + The memory containing the name is released when the program is destroyed by + :func:`destroy_program`. The identical *name_expression* must have been previously + provided to :func:`add_name_expression`. + + Args: + prog (intptr_t): CUDA Runtime Compilation program. + name_expression (bytes): Constant expression denoting the address of a + ``__global__`` function or ``__device__``/``__constant__`` variable. + + Returns: + bytes: C string containing the lowered (mangled) name, or ``None``. + + .. seealso:: `nvrtcGetLoweredName` + """ + cdef const char* c_name_expression = name_expression + cdef const char* lowered_name = NULL + with nogil: + __status__ = nvrtcGetLoweredName(prog, c_name_expression, &lowered_name) + check_status(__status__) + return lowered_name if lowered_name != NULL else None + + +cpdef install_bundled_headers(bytes install_path, unsigned int flags): + """nvrtcInstallBundledHeaders extracts CUDA headers bundled with NVRTC to a specified directory. + + NVRTC bundles a set of CUDA Toolkit headers and CCCL within libnvrtc-builtins. + After extraction, compile kernels by passing ``-I`` and + ``-I/cccl`` to :func:`compile_program`. A version marker file + (``.nvrtc_headers_version``) is created to track the installed version. + The function is thread-safe and process-safe; concurrent calls are serialized + using file locking. + + Args: + install_path (bytes): Path where headers should be extracted (UTF-8 encoded). + The directory is created if it does not exist. + flags (unsigned int): Bitwise OR of :class:`InstallHeadersFlag` values (or 0 + for the default ``SKIP_IF_EXISTS`` behaviour). + + Returns: + bytes | None: Detailed error message on failure, or ``None`` on success. + + .. seealso:: `nvrtcInstallBundledHeaders` + """ + cdef const char* c_install_path = install_path + cdef const char* error_log = NULL + with nogil: + __status__ = nvrtcInstallBundledHeaders(c_install_path, flags, &error_log) + check_status(__status__) + return error_log if error_log != NULL else None + + +cpdef tuple get_bundled_headers_info(): + """nvrtcGetBundledHeadersInfo queries information about the bundled headers without extracting them. + + Allows users to determine if bundled headers are available and get size estimates + before calling :func:`install_bundled_headers`. + + Returns: + tuple[BundledHeadersInfo, bytes | None]: Header information struct and an optional + detailed error message (``None`` on success). + + .. seealso:: `nvrtcGetBundledHeadersInfo` + """ + cdef BundledHeadersInfo info = BundledHeadersInfo() + cdef nvrtcBundledHeadersInfo* c_info = (info._get_ptr()) + cdef const char* error_log = NULL + with nogil: + __status__ = nvrtcGetBundledHeadersInfo(c_info, &error_log) + check_status(__status__) + return info, (error_log if error_log != NULL else None) + + +cpdef remove_bundled_headers(bytes install_path): + """nvrtcRemoveBundledHeaders removes previously installed bundled headers. + + Recursively removes all files and subdirectories within the installation + directory to help manage disk space. + + .. note:: This removes ALL contents of the specified directory, not just files + installed by NVRTC. Use with caution. + + Args: + install_path (bytes): Path where headers were previously installed; must be + the same path used with :func:`install_bundled_headers`. + + Returns: + bytes | None: Detailed error message on failure, or ``None`` on success. + + .. seealso:: `nvrtcRemoveBundledHeaders` + """ + cdef const char* c_install_path = install_path + cdef const char* error_log = NULL + with nogil: + __status__ = nvrtcRemoveBundledHeaders(c_install_path, &error_log) + check_status(__status__) + return error_log if error_log != NULL else None + + +cpdef str get_error_string(int result): + """nvrtcGetErrorString is a helper function that returns a string describing the given ``nvrtcResult`` code, e.g., NVRTC_SUCCESS to ``"NVRTC_SUCCESS"``. For unrecognized enumeration values, it returns ``"NVRTC_ERROR unknown"``. + + Args: + result (Result): CUDA Runtime Compilation API result code. + + .. seealso:: `nvrtcGetErrorString` + """ + cdef const char *_output_cstr_ + cdef bytes _output_ + with nogil: + _output_cstr_ = nvrtcGetErrorString(<_Result>result) + _output_ = _output_cstr_ + return _output_.decode() + + +cpdef tuple version(): + """nvrtcVersion sets the output parameters ``major`` and ``minor`` with the CUDA Runtime Compilation version number. + + Returns: + A 2-tuple containing: + + - int: CUDA Runtime Compilation major version number. + - int: CUDA Runtime Compilation minor version number. + + .. seealso:: `nvrtcVersion` + """ + cdef int major + cdef int minor + with nogil: + __status__ = nvrtcVersion(&major, &minor) + check_status(__status__) + return (major, minor) + + +cpdef int get_num_supported_archs() except? -1: + """nvrtcGetNumSupportedArchs sets the output parameter ``num_archs`` with the number of architectures supported by NVRTC. This can then be used to pass an array to ``nvrtcGetSupportedArchs`` to get the supported architectures. + + Returns: + int: number of supported architectures. + + .. seealso:: `nvrtcGetNumSupportedArchs` + """ + cdef int num_archs + with nogil: + __status__ = nvrtcGetNumSupportedArchs(&num_archs) + check_status(__status__) + return num_archs + + +cpdef object get_supported_archs(): + """nvrtcGetSupportedArchs populates the array passed via the output parameter ``supported_archs`` with the architectures supported by NVRTC. The array is sorted in the ascending order. The size of the array to be passed can be determined using ``nvrtcGetNumSupportedArchs``. + + Returns: + int: sorted array of supported architectures. + + .. seealso:: `nvrtcGetSupportedArchs` + """ + cdef int numArchs + with nogil: + __status__ = nvrtcGetNumSupportedArchs(&numArchs) + check_status(__status__) + if numArchs == 0: + return _cyb_view.array(shape=(1,), itemsize=sizeof(int), format="i", mode="c")[:0] + cdef _cyb_view.array supported_archs = _cyb_view.array(shape=(numArchs,), itemsize=sizeof(int), format="i", mode="c") + cdef int *supported_archs_ptr = (supported_archs.data) + with nogil: + __status__ = nvrtcGetSupportedArchs(supported_archs_ptr) + check_status(__status__) + return supported_archs + + +cpdef destroy_program(intptr_t prog): + """nvrtcDestroyProgram destroys the given program. + + Args: + prog (intptr_t): CUDA Runtime Compilation program. + + .. seealso:: `nvrtcDestroyProgram` + """ + cdef Program _prog_ = prog + with nogil: + __status__ = nvrtcDestroyProgram(&_prog_) + check_status(__status__) + + +cpdef size_t get_ptx_size(intptr_t prog) except? 0: + """nvrtcGetPTXSize sets the value of ``ptx_size_ret`` with the size of the PTX generated by the previous compilation of ``prog`` (including the trailing ``NULL``). + + Args: + prog (intptr_t): CUDA Runtime Compilation program. + + Returns: + size_t: Size of the generated PTX (including the trailing + ``NULL``). + + .. seealso:: `nvrtcGetPTXSize` + """ + cdef size_t ptx_size_ret + with nogil: + __status__ = nvrtcGetPTXSize(prog, &ptx_size_ret) + check_status(__status__) + return ptx_size_ret + + +cpdef bytes get_ptx(intptr_t prog): + """nvrtcGetPTX stores the PTX generated by the previous compilation of ``prog`` in the memory pointed by ``ptx``. + + Args: + prog (intptr_t): CUDA Runtime Compilation program. + + Returns: + char: Compiled result. + + .. seealso:: `nvrtcGetPTX` + """ + cdef size_t ptxSizeRet + with nogil: + __status__ = nvrtcGetPTXSize(prog, &ptxSizeRet) + check_status(__status__) + if ptxSizeRet == 0: + return b"" + cdef bytes _ptx_ = bytes(ptxSizeRet) + cdef char* ptx = _ptx_ + with nogil: + __status__ = nvrtcGetPTX(prog, ptx) + check_status(__status__) + return _ptx_ + + +cpdef size_t get_cubin_size(intptr_t prog) except? 0: + """nvrtcGetCUBINSize sets the value of ``cubin_size_ret`` with the size of the cubin generated by the previous compilation of ``prog``. The value of cubin_size_ret is set to 0 if the value specified to ``-arch`` is a virtual architecture instead of an actual architecture. + + Args: + prog (intptr_t): CUDA Runtime Compilation program. + + Returns: + size_t: Size of the generated cubin. + + .. seealso:: `nvrtcGetCUBINSize` + """ + cdef size_t cubin_size_ret + with nogil: + __status__ = nvrtcGetCUBINSize(prog, &cubin_size_ret) + check_status(__status__) + return cubin_size_ret + + +cpdef bytes get_cubin(intptr_t prog): + """nvrtcGetCUBIN stores the cubin generated by the previous compilation of ``prog`` in the memory pointed by ``cubin``. No cubin is available if the value specified to ``-arch`` is a virtual architecture instead of an actual architecture. The cubin does not contain code for the Tile functions (``__tile__`` / ``__tile_global__``) or variables (``__tile__``); use :func:`get_tile_ir` to extract the cuda_tile IR generated for Tile code. + + Args: + prog (intptr_t): CUDA Runtime Compilation program. + + Returns: + char: Compiled and assembled result. + + .. seealso:: `nvrtcGetCUBIN` + """ + cdef size_t cubinSizeRet + with nogil: + __status__ = nvrtcGetCUBINSize(prog, &cubinSizeRet) + check_status(__status__) + if cubinSizeRet == 0: + return b"" + cdef bytes _cubin_ = bytes(cubinSizeRet) + cdef char* cubin = _cubin_ + with nogil: + __status__ = nvrtcGetCUBIN(prog, cubin) + check_status(__status__) + return _cubin_ + + +cpdef size_t get_ltoir_size(intptr_t prog) except? 0: + """nvrtcGetLTOIRSize sets the value of ``ltoir_size_ret`` with the size of the LTO IR generated by the previous compilation of ``prog``. The value of ltoir_size_ret is set to 0 if the program was not compiled with ``-dlto``. + + Args: + prog (intptr_t): CUDA Runtime Compilation program. + + Returns: + size_t: Size of the generated LTO IR. + + .. seealso:: `nvrtcGetLTOIRSize` + """ + cdef size_t ltoir_size_ret + with nogil: + __status__ = nvrtcGetLTOIRSize(prog, <oir_size_ret) + check_status(__status__) + return ltoir_size_ret + + +cpdef bytes get_ltoir(intptr_t prog): + """nvrtcGetltoir stores the LTO IR generated by the previous compilation of ``prog`` in the memory pointed by ``ltoir``. No LTO IR is available if the program was compiled without ``-dlto``. + + Args: + prog (intptr_t): CUDA Runtime Compilation program. + + Returns: + char: Compiled result. + + .. seealso:: `nvrtcGetLTOIR` + """ + cdef size_t LTOIRSizeRet + with nogil: + __status__ = nvrtcGetLTOIRSize(prog, <OIRSizeRet) + check_status(__status__) + if LTOIRSizeRet == 0: + return b"" + cdef bytes _ltoir_ = bytes(LTOIRSizeRet) + cdef char* ltoir = _ltoir_ + with nogil: + __status__ = nvrtcGetLTOIR(prog, ltoir) + check_status(__status__) + return _ltoir_ + + +cpdef size_t get_optix_ir_size(intptr_t prog) except? 0: + """nvrtcGetOptiXIRSize sets the value of ``optixir_size_ret`` with the size of the OptiX IR generated by the previous compilation of ``prog``. The value of nvrtcGetOptiXIRSize is set to 0 if the program was compiled with options incompatible with OptiX IR generation. + + Args: + prog (intptr_t): CUDA Runtime Compilation program. + + Returns: + size_t: Size of the generated LTO IR. + + .. seealso:: `nvrtcGetOptiXIRSize` + """ + cdef size_t optixir_size_ret + with nogil: + __status__ = nvrtcGetOptiXIRSize(prog, &optixir_size_ret) + check_status(__status__) + return optixir_size_ret + + +cpdef bytes get_optix_ir(intptr_t prog): + """nvrtcGetOptiXIR stores the OptiX IR generated by the previous compilation of ``prog`` in the memory pointed by ``optixir``. No OptiX IR is available if the program was compiled with options incompatible with OptiX IR generation. + + Args: + prog (intptr_t): CUDA Runtime Compilation program. + + Returns: + char: Optix IR Compiled result. + + .. seealso:: `nvrtcGetOptiXIR` + """ + cdef size_t optixirSizeRet + with nogil: + __status__ = nvrtcGetOptiXIRSize(prog, &optixirSizeRet) + check_status(__status__) + if optixirSizeRet == 0: + return b"" + cdef bytes _optixir_ = bytes(optixirSizeRet) + cdef char* optixir = _optixir_ + with nogil: + __status__ = nvrtcGetOptiXIR(prog, optixir) + check_status(__status__) + return _optixir_ + + +cpdef size_t get_program_log_size(intptr_t prog) except? 0: + """nvrtcGetProgramLogSize sets ``log_size_ret`` with the size of the log generated by the previous compilation of ``prog`` (including the trailing ``NULL``). + + Args: + prog (intptr_t): CUDA Runtime Compilation program. + + Returns: + size_t: Size of the compilation log (including the trailing + ``NULL``). + + .. seealso:: `nvrtcGetProgramLogSize` + """ + cdef size_t log_size_ret + with nogil: + __status__ = nvrtcGetProgramLogSize(prog, &log_size_ret) + check_status(__status__) + return log_size_ret + + +cpdef bytes get_program_log(intptr_t prog): + """nvrtcGetProgramLog stores the log generated by the previous compilation of ``prog`` in the memory pointed by ``log``. + + Args: + prog (intptr_t): CUDA Runtime Compilation program. + + Returns: + char: Compilation log. + + .. seealso:: `nvrtcGetProgramLog` + """ + cdef size_t logSizeRet + with nogil: + __status__ = nvrtcGetProgramLogSize(prog, &logSizeRet) + check_status(__status__) + if logSizeRet == 0: + return b"" + cdef bytes _log_ = bytes(logSizeRet) + cdef char* log = _log_ + with nogil: + __status__ = nvrtcGetProgramLog(prog, log) + check_status(__status__) + return _log_ + + +cpdef add_name_expression(intptr_t prog, name_expression): + """nvrtcAddNameExpression notes the given name expression denoting the address of a global function or device/__constant__ variable. + + Args: + prog (intptr_t): CUDA Runtime Compilation program. + name_expression (str): constant expression denoting the + address of a global function or device/__constant__ + variable. + + .. seealso:: `nvrtcAddNameExpression` + """ + if not isinstance(name_expression, str): + raise TypeError("name_expression must be a Python str") + cdef bytes _temp_name_expression_ = (name_expression).encode() + cdef char* _name_expression_ = _temp_name_expression_ + with nogil: + __status__ = nvrtcAddNameExpression(prog, _name_expression_) + check_status(__status__) + + +cpdef size_t get_pch_heap_size() except? 0: + """retrieve the current size of the PCH Heap. + + Returns: + size_t: pointer to location where the size of the PCH Heap + will be stored. + + .. seealso:: `nvrtcGetPCHHeapSize` + """ + cdef size_t ret + with nogil: + __status__ = nvrtcGetPCHHeapSize(&ret) + check_status(__status__) + return ret + + +cpdef set_pch_heap_size(size_t size): + """set the size of the PCH Heap. + + Args: + size (size_t): requested size of the PCH Heap, in bytes. + + .. seealso:: `nvrtcSetPCHHeapSize` + """ + with nogil: + __status__ = nvrtcSetPCHHeapSize(size) + check_status(__status__) + + +cpdef int get_pch_create_status(intptr_t prog) except? -1: + """returns the PCH creation status. + + Args: + prog (intptr_t): CUDA Runtime Compilation program. + + .. seealso:: `nvrtcGetPCHCreateStatus` + """ + cdef int ret + with nogil: + ret = nvrtcGetPCHCreateStatus(prog) + return ret + + +cpdef size_t get_pch_heap_size_required(intptr_t prog) except? 0: + """retrieve the required size of the PCH heap required to compile the given program. + + Args: + prog (intptr_t): CUDA Runtime Compilation program. + + Returns: + size_t: pointer to location where the required size of the PCH + Heap will be stored. + + .. seealso:: `nvrtcGetPCHHeapSizeRequired` + """ + cdef size_t size + with nogil: + __status__ = nvrtcGetPCHHeapSizeRequired(prog, &size) + check_status(__status__) + return size + + +cpdef size_t get_tile_ir_size(intptr_t prog) except? 0: + """nvrtcGetTileIRSize sets the value of ``tile_ir_size_ret`` with the size of the cuda_tile IR generated by the previous compilation of ``prog``. + + Args: + prog (intptr_t): CUDA Runtime Compilation program. + + Returns: + size_t: Size of the generated cuda_tile IR. + + .. seealso:: `nvrtcGetTileIRSize` + """ + cdef size_t tile_ir_size_ret + with nogil: + __status__ = nvrtcGetTileIRSize(prog, &tile_ir_size_ret) + check_status(__status__) + return tile_ir_size_ret + + +cpdef bytes get_tile_ir(intptr_t prog): + """nvrtcGettile_ir stores the cuda_tile IR generated by the previous compilation of ``prog`` in the memory pointed by ``tile_ir``. + + Args: + prog (intptr_t): CUDA Runtime Compilation program. + + Returns: + char: Generated cuda_tile IR. + + .. seealso:: `nvrtcGetTileIR` + """ + cdef size_t TileIRSizeRet + with nogil: + __status__ = nvrtcGetTileIRSize(prog, &TileIRSizeRet) + check_status(__status__) + if TileIRSizeRet == 0: + return b"" + cdef bytes _tile_ir_ = bytes(TileIRSizeRet) + cdef char* tile_ir = _tile_ir_ + with nogil: + __status__ = nvrtcGetTileIR(prog, tile_ir) + check_status(__status__) + return _tile_ir_ + + +del _cyb_FastEnum diff --git a/cuda_bindings/docs/source/release/13.4.0-notes.rst b/cuda_bindings/docs/source/release/13.4.0-notes.rst index f269eca2a3d..603b4accd8c 100644 --- a/cuda_bindings/docs/source/release/13.4.0-notes.rst +++ b/cuda_bindings/docs/source/release/13.4.0-notes.rst @@ -13,6 +13,14 @@ Deprecation Notices removed in a future version. Python 3.10 reaches end of life in October 2026 per the `CPython support cycle `_. +Prerelease feature +------------------ + +A new version of the ``nvrtc`` API is available as ``cuda.bindings._v2.nvrtc``. The +primary improvements are: (1) raising exceptions rather than returning error +codes, (2) uses PEP8-compliant naming, and (3) is more performant. This API is +still experimental and subject to change. + Known issues ------------ diff --git a/cuda_bindings/examples/extra/jit_program.py b/cuda_bindings/examples/extra/jit_program.py index 7f55b2243ed..7a5cc1495fc 100644 --- a/cuda_bindings/examples/extra/jit_program.py +++ b/cuda_bindings/examples/extra/jit_program.py @@ -17,15 +17,15 @@ import numpy as np from cuda.bindings import driver as cuda -from cuda.bindings import nvrtc +from cuda.bindings._v2 import nvrtc def assert_drv(err): if isinstance(err, cuda.CUresult): if err != cuda.CUresult.CUDA_SUCCESS: raise RuntimeError(f"Cuda Error: {err}") - elif isinstance(err, nvrtc.nvrtcResult): - if err != nvrtc.nvrtcResult.NVRTC_SUCCESS: + elif isinstance(err, nvrtc.Result): + if err != nvrtc.Result.SUCCESS: raise RuntimeError(f"Nvrtc Error: {err}") else: raise RuntimeError(f"Unknown error type: {err}") @@ -57,8 +57,7 @@ def main(): assert_drv(err) # Create program - err, prog = nvrtc.nvrtcCreateProgram(str.encode(saxpy), b"saxpy.cu", 0, None, None) - assert_drv(err) + prog = nvrtc.create_program(str.encode(saxpy), b"saxpy.cu") # Get target architecture err, major = cuda.cuDeviceGetAttribute( @@ -69,38 +68,24 @@ def main(): cuda.CUdevice_attribute.CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, cu_device ) assert_drv(err) - err, nvrtc_major, nvrtc_minor = nvrtc.nvrtcVersion() - assert_drv(err) + nvrtc_major, nvrtc_minor = nvrtc.version() use_cubin = nvrtc_minor >= 1 prefix = "sm" if use_cubin else "compute" arch_arg = bytes(f"--gpu-architecture={prefix}_{major}{minor}", "ascii") # Compile program opts = [b"--fmad=false", arch_arg] - (err,) = nvrtc.nvrtcCompileProgram(prog, len(opts), opts) - assert_drv(err) + nvrtc.compile_program(prog, opts) # Get log from compilation - err, log_size = nvrtc.nvrtcGetProgramLogSize(prog) - assert_drv(err) - log = b" " * log_size - (err,) = nvrtc.nvrtcGetProgramLog(prog, log) - assert_drv(err) + log = nvrtc.get_program_log(prog) print(log.decode()) # Get data from compilation if use_cubin: - err, data_size = nvrtc.nvrtcGetCUBINSize(prog) - assert_drv(err) - data = b" " * data_size - (err,) = nvrtc.nvrtcGetCUBIN(prog, data) - assert_drv(err) + data = nvrtc.get_cubin(prog) else: - err, data_size = nvrtc.nvrtcGetPTXSize(prog) - assert_drv(err) - data = b" " * data_size - (err,) = nvrtc.nvrtcGetPTX(prog, data) - assert_drv(err) + data = nvrtc.get_ptx(prog) # Load data as module data and retrieve function data = np.char.array(data) diff --git a/cuda_bindings/tests/legacy_api/README.md b/cuda_bindings/tests/legacy_api/README.md new file mode 100644 index 00000000000..df02d4791dc --- /dev/null +++ b/cuda_bindings/tests/legacy_api/README.md @@ -0,0 +1,4 @@ +This directory contains tests for the legacy (non-v2-prefixed) APIs for +`driver`, `runtime` and `nvrtc`. These were created from snapshots of the +original version, and then the originals were updated to use the new API. Do +not add to the tests here, and do not update them unless necessary to fix a bug. diff --git a/cuda_bindings/tests/legacy_api/test_legacy_kernelParams.py b/cuda_bindings/tests/legacy_api/test_legacy_kernelParams.py new file mode 100644 index 00000000000..555d6a7284c --- /dev/null +++ b/cuda_bindings/tests/legacy_api/test_legacy_kernelParams.py @@ -0,0 +1,802 @@ +# SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +import ctypes + +import numpy as np +import pytest + +import cuda.bindings.driver as cuda +import cuda.bindings.nvrtc as nvrtc +import cuda.bindings.runtime as cudart + + +def ASSERT_DRV(err): + if isinstance(err, cuda.CUresult): + if err != cuda.CUresult.CUDA_SUCCESS: + raise RuntimeError(f"Cuda Error: {err}") + elif isinstance(err, cudart.cudaError_t): + if err != cudart.cudaError_t.cudaSuccess: + raise RuntimeError(f"Cudart Error: {err}") + elif isinstance(err, nvrtc.nvrtcResult): + if err != nvrtc.nvrtcResult.NVRTC_SUCCESS: + raise RuntimeError(f"Nvrtc Error: {err}") + else: + raise RuntimeError(f"Unknown error type: {err}") + + +def common_nvrtc(allKernelStrings, dev): + err, major = cuda.cuDeviceGetAttribute(cuda.CUdevice_attribute.CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, dev) + ASSERT_DRV(err) + err, minor = cuda.cuDeviceGetAttribute(cuda.CUdevice_attribute.CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, dev) + ASSERT_DRV(err) + err, _, nvrtc_minor = nvrtc.nvrtcVersion() + ASSERT_DRV(err) + use_cubin = nvrtc_minor >= 1 + prefix = "sm" if use_cubin else "compute" + arch_arg = bytes(f"--gpu-architecture={prefix}_{major}{minor}", "ascii") + + err, prog = nvrtc.nvrtcCreateProgram(str.encode(allKernelStrings), b"allKernelStrings.cu", 0, None, None) + ASSERT_DRV(err) + opts = (b"--fmad=false", arch_arg) + (err,) = nvrtc.nvrtcCompileProgram(prog, len(opts), opts) + + err_log, logSize = nvrtc.nvrtcGetProgramLogSize(prog) + ASSERT_DRV(err_log) + log = b" " * logSize + (err_log,) = nvrtc.nvrtcGetProgramLog(prog, log) + ASSERT_DRV(err_log) + result = log.decode() + if len(result) > 1: + print(result) + ASSERT_DRV(err) + + if use_cubin: + err, dataSize = nvrtc.nvrtcGetCUBINSize(prog) + ASSERT_DRV(err) + data = b" " * dataSize + (err,) = nvrtc.nvrtcGetCUBIN(prog, data) + ASSERT_DRV(err) + else: + err, dataSize = nvrtc.nvrtcGetPTXSize(prog) + ASSERT_DRV(err) + data = b" " * dataSize + (err,) = nvrtc.nvrtcGetPTX(prog, data) + ASSERT_DRV(err) + + err, module = cuda.cuModuleLoadData(np.char.array(data)) + ASSERT_DRV(err) + + return module + + +def test_kernelParams_empty(device): + kernelString = """\ + static __device__ bool isDone; + extern "C" __global__ + void empty_kernel() + { + isDone = true; + if (isDone) return; + } + """ + + module = common_nvrtc(kernelString, device) + + # cudaStructs kernel + err, kernel = cuda.cuModuleGetFunction(module, b"empty_kernel") + ASSERT_DRV(err) + + err, stream = cuda.cuStreamCreate(0) + ASSERT_DRV(err) + + (err,) = cuda.cuLaunchKernel( + kernel, + 1, + 1, + 1, # grid dim + 1, + 1, + 1, # block dim + 0, + stream, # shared mem and stream + ((), ()), + 0, + ) # arguments + ASSERT_DRV(err) + (err,) = cuda.cuLaunchKernel( + kernel, + 1, + 1, + 1, # grid dim + 1, + 1, + 1, # block dim + 0, + stream, # shared mem and stream + None, + 0, + ) # arguments + ASSERT_DRV(err) + + # Retrieve global and validate + isDone_host = ctypes.c_bool() + err, isDonePtr_device, isDonePtr_device_size = cuda.cuModuleGetGlobal(module, b"isDone") + ASSERT_DRV(err) + assert isDonePtr_device_size == ctypes.sizeof(ctypes.c_bool) + (err,) = cuda.cuMemcpyDtoHAsync(isDone_host, isDonePtr_device, ctypes.sizeof(ctypes.c_bool), stream) + ASSERT_DRV(err) + (err,) = cuda.cuStreamSynchronize(stream) + ASSERT_DRV(err) + assert isDone_host.value is True + + (err,) = cuda.cuStreamDestroy(stream) + ASSERT_DRV(err) + (err,) = cuda.cuModuleUnload(module) + ASSERT_DRV(err) + + +@pytest.mark.parametrize("use_ctypes_as_values", [False, True], ids=["no-ctypes", "ctypes"]) +def test_kernelParams(use_ctypes_as_values, device): + if use_ctypes_as_values: + assertValues_host = ( + ctypes.c_bool(True), + ctypes.c_char(b"Z"), + ctypes.c_wchar("Ā"), + ctypes.c_byte(-127), + ctypes.c_ubyte(255), + ctypes.c_short(1), + ctypes.c_ushort(1), + ctypes.c_int(2), + ctypes.c_uint(2), + ctypes.c_long(3), + ctypes.c_ulong(3), + ctypes.c_longlong(4), + ctypes.c_ulonglong(4), + ctypes.c_size_t(5), + ctypes.c_float(123.456), + ctypes.c_float(123.456), + ctypes.c_void_p(0xDEADBEEF), + ) + else: + assertValues_host = ( + True, + b"Z", + "Ā", + -127, + 255, + 90, + 72, + 85, + 82, + 66, + 65, + 86, + 90, + 33, + 123.456, + 123.456, + 0xDEADBEEF, + ) + assertTypes_host = ( + ctypes.c_bool, + ctypes.c_char, + ctypes.c_wchar, + ctypes.c_byte, + ctypes.c_ubyte, + ctypes.c_short, + ctypes.c_ushort, + ctypes.c_int, + ctypes.c_uint, + ctypes.c_long, + ctypes.c_ulong, + ctypes.c_longlong, + ctypes.c_ulonglong, + ctypes.c_size_t, + ctypes.c_float, + ctypes.c_double, + ctypes.c_void_p, + ) + + basicKernelString = """\ + extern "C" __global__ + void basic(bool b, + char c, wchar_t wc, + signed char byte, unsigned char ubyte, + short s, unsigned short us, + int i, unsigned int ui, + long l, unsigned long ul, + long long ll, unsigned long long ull, + size_t size, + float f, double d, + void *p, + bool *pb, + char *pc, wchar_t *pwc, + signed char *pbyte, unsigned char *pubyte, + short *ps, unsigned short *pus, + int *pi, unsigned int *pui, + long *pl, unsigned long *pul, + long long *pll, unsigned long long *pull, + size_t *psize, + float *pf, double *pd) + { + assert(b == {}); + assert(c == {}); + assert(wc == {}); + assert(byte == {}); + assert(ubyte == {}); + assert(s == {}); + assert(us == {}); + assert(i == {}); + assert(ui == {}); + assert(l == {}); + assert(ul == {}); + assert(ll == {}); + assert(ull == {}); + assert(size == {}); + assert(f == {}); + assert(d == {}); + assert(p == (void*){}); + *pb = b; + *pc = c; + *pwc = wc; + *pbyte = byte; + *pubyte = ubyte; + *ps = s; + *pus = us; + *pi = i; + *pui = ui; + *pl = l; + *pul = ul; + *pll = ll; + *pull = ull; + *psize = size; + *pf = f; + *pd = d; + } + """ + idx = 0 + while "{}" in basicKernelString: + val = assertValues_host[idx].value if use_ctypes_as_values else assertValues_host[idx] + if assertTypes_host[idx] == ctypes.c_float: + basicKernelString = basicKernelString.replace("{}", str(float(val)) + "f", 1) + elif assertTypes_host[idx] == ctypes.c_double: + basicKernelString = basicKernelString.replace("{}", str(float(val)), 1) + elif assertTypes_host[idx] == ctypes.c_char: + basicKernelString = basicKernelString.replace("{}", str(val)[1:], 1) + elif assertTypes_host[idx] == ctypes.c_wchar: + basicKernelString = basicKernelString.replace("{}", str(ord(val)), 1) + else: + basicKernelString = basicKernelString.replace("{}", str(int(val)), 1) + idx += 1 + + module = common_nvrtc(basicKernelString, device) + + err, kernel = cuda.cuModuleGetFunction(module, b"basic") + ASSERT_DRV(err) + + err, stream = cuda.cuStreamCreate(0) + ASSERT_DRV(err) + + # Prepare kernel + err, pb = cuda.cuMemAlloc(ctypes.sizeof(ctypes.c_bool)) + ASSERT_DRV(err) + err, pc = cuda.cuMemAlloc(ctypes.sizeof(ctypes.c_char)) + ASSERT_DRV(err) + err, pwc = cuda.cuMemAlloc(ctypes.sizeof(ctypes.c_wchar)) + ASSERT_DRV(err) + err, pbyte = cuda.cuMemAlloc(ctypes.sizeof(ctypes.c_byte)) + ASSERT_DRV(err) + err, pubyte = cuda.cuMemAlloc(ctypes.sizeof(ctypes.c_ubyte)) + ASSERT_DRV(err) + err, ps = cuda.cuMemAlloc(ctypes.sizeof(ctypes.c_short)) + ASSERT_DRV(err) + err, pus = cuda.cuMemAlloc(ctypes.sizeof(ctypes.c_ushort)) + ASSERT_DRV(err) + err, pi = cuda.cuMemAlloc(ctypes.sizeof(ctypes.c_int)) + ASSERT_DRV(err) + err, pui = cuda.cuMemAlloc(ctypes.sizeof(ctypes.c_uint)) + ASSERT_DRV(err) + err, pl = cuda.cuMemAlloc(ctypes.sizeof(ctypes.c_long)) + ASSERT_DRV(err) + err, pul = cuda.cuMemAlloc(ctypes.sizeof(ctypes.c_ulong)) + ASSERT_DRV(err) + err, pll = cuda.cuMemAlloc(ctypes.sizeof(ctypes.c_longlong)) + ASSERT_DRV(err) + err, pull = cuda.cuMemAlloc(ctypes.sizeof(ctypes.c_ulonglong)) + ASSERT_DRV(err) + err, psize = cuda.cuMemAlloc(ctypes.sizeof(ctypes.c_size_t)) + ASSERT_DRV(err) + err, pf = cuda.cuMemAlloc(ctypes.sizeof(ctypes.c_float)) + ASSERT_DRV(err) + err, pd = cuda.cuMemAlloc(ctypes.sizeof(ctypes.c_double)) + ASSERT_DRV(err) + + assertValues_device = (pb, pc, pwc, pbyte, pubyte, ps, pus, pi, pui, pl, pul, pll, pull, psize, pf, pd) + assertTypes_device = ( + None, + None, + None, + None, + None, + None, + None, + None, + None, + None, + None, + None, + None, + None, + None, + None, + ) + + basicKernelValues = assertValues_host + assertValues_device + basicKernelTypes = assertTypes_host + assertTypes_device + (err,) = cuda.cuLaunchKernel( + kernel, + 1, + 1, + 1, # grid dim + 1, + 1, + 1, # block dim + 0, + stream, # shared mem and stream + (basicKernelValues, basicKernelTypes), + 0, + ) # arguments + ASSERT_DRV(err) + + # Retrieve each dptr + host_params = tuple([valueType() for valueType in assertTypes_host[:-1]]) + for i in range(len(host_params)): + (err,) = cuda.cuMemcpyDtoHAsync( + host_params[i], assertValues_device[i], ctypes.sizeof(assertTypes_host[i]), stream + ) + ASSERT_DRV(err) + + # Validate retrieved values + (err,) = cuda.cuStreamSynchronize(stream) + ASSERT_DRV(err) + for i in range(len(host_params)): + val = basicKernelValues[i].value if use_ctypes_as_values else basicKernelValues[i] + if basicKernelTypes[i] == ctypes.c_float: + if use_ctypes_as_values: + assert val == host_params[i].value + else: + assert val == (int(host_params[i].value * 1000) / 1000) + else: + assert val == host_params[i].value + + (err,) = cuda.cuMemFree(pb) + ASSERT_DRV(err) + (err,) = cuda.cuMemFree(pc) + ASSERT_DRV(err) + (err,) = cuda.cuMemFree(pwc) + ASSERT_DRV(err) + (err,) = cuda.cuMemFree(pbyte) + ASSERT_DRV(err) + (err,) = cuda.cuMemFree(pubyte) + ASSERT_DRV(err) + (err,) = cuda.cuMemFree(ps) + ASSERT_DRV(err) + (err,) = cuda.cuMemFree(pus) + ASSERT_DRV(err) + (err,) = cuda.cuMemFree(pi) + ASSERT_DRV(err) + (err,) = cuda.cuMemFree(pui) + ASSERT_DRV(err) + (err,) = cuda.cuMemFree(pl) + ASSERT_DRV(err) + (err,) = cuda.cuMemFree(pul) + ASSERT_DRV(err) + (err,) = cuda.cuMemFree(pll) + ASSERT_DRV(err) + (err,) = cuda.cuMemFree(pull) + ASSERT_DRV(err) + (err,) = cuda.cuMemFree(psize) + ASSERT_DRV(err) + (err,) = cuda.cuMemFree(pf) + ASSERT_DRV(err) + (err,) = cuda.cuMemFree(pd) + ASSERT_DRV(err) + (err,) = cuda.cuStreamDestroy(stream) + ASSERT_DRV(err) + (err,) = cuda.cuModuleUnload(module) + ASSERT_DRV(err) + + +def test_kernelParams_types_cuda(device): + err, uvaSupported = cuda.cuDeviceGetAttribute( + cuda.CUdevice_attribute.CU_DEVICE_ATTRIBUTE_UNIFIED_ADDRESSING, device + ) + ASSERT_DRV(err) + + err, perr = cudart.cudaMalloc(ctypes.sizeof(ctypes.c_int)) + ASSERT_DRV(err) + err, pSurface_host = cudart.cudaHostAlloc(cudart.sizeof(cudart.cudaSurfaceObject_t), cudart.cudaHostAllocMapped) + ASSERT_DRV(err) + err, pDim3_host = cudart.cudaHostAlloc(cudart.sizeof(cudart.dim3), cudart.cudaHostAllocMapped) + ASSERT_DRV(err) + + # Get device pointer if UVM is not enabled + if uvaSupported: + kernelValues = ( + cudart.cudaError_t.cudaErrorUnknown, + perr, # enums + cudart.cudaSurfaceObject_t(248), + cudart.cudaSurfaceObject_t(_ptr=pSurface_host), # typedef of primative + cudart.dim3(), + cudart.dim3(_ptr=pDim3_host), + ) # struct + else: + err, pSurface_device = cudart.cudaHostGetDevicePointer(pSurface_host, 0) + ASSERT_DRV(err) + err, pDim3_device = cudart.cudaHostGetDevicePointer(pDim3_host, 0) + ASSERT_DRV(err) + kernelValues = ( + cudart.cudaError_t.cudaErrorUnknown, + perr, # enums + cudart.cudaSurfaceObject_t(248), + cudart.cudaSurfaceObject_t(_ptr=pSurface_device), # typedef of primative + cudart.dim3(), + cudart.dim3(_ptr=pDim3_device), + ) # struct + kernelTypes = (None, ctypes.c_void_p, None, ctypes.c_void_p, None, ctypes.c_void_p) + kernelValues[4].x = 1 + kernelValues[4].y = 2 + kernelValues[4].z = 3 + + kernelString = """\ + extern "C" __global__ + void structsCuda(cudaError_t err, cudaError_t *perr, + cudaSurfaceObject_t surface, cudaSurfaceObject_t *pSurface, + dim3 dim, dim3* pdim) + { + *perr = err; + *pSurface = surface; + pdim->x = dim.x; + pdim->y = dim.y; + pdim->z = dim.z; + } + """ + + module = common_nvrtc(kernelString, device) + + # cudaStructs kernel + err, kernel = cuda.cuModuleGetFunction(module, b"structsCuda") + ASSERT_DRV(err) + + err, stream = cuda.cuStreamCreate(0) + ASSERT_DRV(err) + + (err,) = cuda.cuLaunchKernel( + kernel, + 1, + 1, + 1, # grid dim + 1, + 1, + 1, # block dim + 0, + stream, # shared mem and stream + (kernelValues, kernelTypes), + 0, + ) # arguments + ASSERT_DRV(err) + + # Retrieve each dptr + host_err = ctypes.c_int() + (err,) = cudart.cudaMemcpyAsync( + ctypes.addressof(host_err), + perr, + ctypes.sizeof(ctypes.c_int()), + cudart.cudaMemcpyKind.cudaMemcpyDeviceToHost, + stream, + ) + ASSERT_DRV(err) + + # Validate kernel values + (err,) = cuda.cuStreamSynchronize(stream) + ASSERT_DRV(err) + cuda_err = cudart.cudaError_t(host_err.value) + + if uvaSupported: + assert kernelValues[0] == cuda_err + assert int(kernelValues[2]) == int(kernelValues[3]) + assert kernelValues[4].x == kernelValues[5].x + assert kernelValues[4].y == kernelValues[5].y + assert kernelValues[4].z == kernelValues[5].z + else: + surface_host = cudart.cudaSurfaceObject_t(_ptr=pSurface_host) + dim3_host = cudart.dim3(_ptr=pDim3_host) + assert kernelValues[0] == cuda_err + assert int(kernelValues[2]) == int(surface_host) + assert kernelValues[4].x == dim3_host.x + assert kernelValues[4].y == dim3_host.y + assert kernelValues[4].z == dim3_host.z + + (err,) = cudart.cudaFree(perr) + ASSERT_DRV(err) + (err,) = cudart.cudaFreeHost(pSurface_host) + ASSERT_DRV(err) + (err,) = cudart.cudaFreeHost(pDim3_host) + ASSERT_DRV(err) + (err,) = cuda.cuStreamDestroy(stream) + ASSERT_DRV(err) + (err,) = cuda.cuModuleUnload(module) + ASSERT_DRV(err) + + +def test_kernelParams_struct_custom(device): + err, uvaSupported = cuda.cuDeviceGetAttribute( + cuda.CUdevice_attribute.CU_DEVICE_ATTRIBUTE_UNIFIED_ADDRESSING, device + ) + ASSERT_DRV(err) + + kernelString = """\ + struct testStruct { + int value; + }; + + extern "C" __global__ + void structCustom(struct testStruct src, struct testStruct *dst) + { + dst->value = src.value; + } + """ + + module = common_nvrtc(kernelString, device) + + err, kernel = cuda.cuModuleGetFunction(module, b"structCustom") + ASSERT_DRV(err) + + err, stream = cuda.cuStreamCreate(0) + ASSERT_DRV(err) + + # structCustom kernel + class testStruct(ctypes.Structure): + _fields_ = [("value", ctypes.c_int)] + + err, pStruct_host = cudart.cudaHostAlloc(ctypes.sizeof(testStruct), cudart.cudaHostAllocMapped) + ASSERT_DRV(err) + + # Get device pointer if UVM is not enabled + if uvaSupported: + kernelValues = (testStruct(5), pStruct_host) + else: + err, pStruct_device = cudart.cudaHostGetDevicePointer(pStruct_host, 0) + ASSERT_DRV(err) + kernelValues = (testStruct(5), pStruct_device) + kernelTypes = (None, ctypes.c_void_p) + + (err,) = cuda.cuLaunchKernel( + kernel, + 1, + 1, + 1, # grid dim + 1, + 1, + 1, # block dim + 0, + stream, # shared mem and stream + (kernelValues, kernelTypes), + 0, + ) # arguments + ASSERT_DRV(err) + + # Validate kernel values + (err,) = cuda.cuStreamSynchronize(stream) + ASSERT_DRV(err) + struct_shared = testStruct.from_address(pStruct_host) + assert kernelValues[0].value == struct_shared.value + + (err,) = cudart.cudaFreeHost(pStruct_host) + ASSERT_DRV(err) + (err,) = cuda.cuStreamDestroy(stream) + ASSERT_DRV(err) + (err,) = cuda.cuModuleUnload(module) + ASSERT_DRV(err) + + +@pytest.mark.parametrize("pass_by_address", [False, True], ids=["by-address", "not-by-address"]) +def test_kernelParams_buffer_protocol(pass_by_address, device): + err, uvaSupported = cuda.cuDeviceGetAttribute( + cuda.CUdevice_attribute.CU_DEVICE_ATTRIBUTE_UNIFIED_ADDRESSING, device + ) + ASSERT_DRV(err) + + kernelString = """\ + struct testStruct { + int value; + }; + extern "C" __global__ + void testkernel(int i, int *pi, + float f, float *pf, + struct testStruct s, struct testStruct *ps) + { + *pi = i; + *pf = f; + ps->value = s.value; + } + """ + + module = common_nvrtc(kernelString, device) + + err, kernel = cuda.cuModuleGetFunction(module, b"testkernel") + ASSERT_DRV(err) + + err, stream = cuda.cuStreamCreate(0) + ASSERT_DRV(err) + + # testkernel kernel + class testStruct(ctypes.Structure): + _fields_ = [("value", ctypes.c_int)] + + err, pInt_host = cudart.cudaHostAlloc(ctypes.sizeof(ctypes.c_int), cudart.cudaHostAllocMapped) + ASSERT_DRV(err) + err, pFloat_host = cudart.cudaHostAlloc(ctypes.sizeof(ctypes.c_float), cudart.cudaHostAllocMapped) + ASSERT_DRV(err) + err, pStruct_host = cudart.cudaHostAlloc(ctypes.sizeof(testStruct), cudart.cudaHostAllocMapped) + ASSERT_DRV(err) + + # Get device pointer if UVM is not enabled + if uvaSupported: + kernelValues = ( + ctypes.c_int(1), + ctypes.c_void_p(pInt_host), + ctypes.c_float(123.456), + ctypes.c_void_p(pFloat_host), + testStruct(5), + ctypes.c_void_p(pStruct_host), + ) + else: + err, pInt_device = cudart.cudaHostGetDevicePointer(pInt_host, 0) + ASSERT_DRV(err) + err, pFloat_device = cudart.cudaHostGetDevicePointer(pFloat_host, 0) + ASSERT_DRV(err) + err, pStruct_device = cudart.cudaHostGetDevicePointer(pStruct_host, 0) + ASSERT_DRV(err) + kernelValues = ( + ctypes.c_int(1), + ctypes.c_void_p(pInt_device), + ctypes.c_float(123.456), + ctypes.c_void_p(pFloat_device), + testStruct(5), + ctypes.c_void_p(pStruct_device), + ) + + packagedParams = (ctypes.c_void_p * len(kernelValues))() + for idx in range(len(packagedParams)): + packagedParams[idx] = ctypes.addressof(kernelValues[idx]) + (err,) = cuda.cuLaunchKernel( + kernel, + 1, + 1, + 1, # grid dim + 1, + 1, + 1, # block dim + 0, + stream, # shared mem and stream + ctypes.addressof(packagedParams) if pass_by_address else packagedParams, + 0, + ) # arguments + ASSERT_DRV(err) + + # Validate kernel values + (err,) = cuda.cuStreamSynchronize(stream) + ASSERT_DRV(err) + assert kernelValues[0].value == ctypes.c_int.from_address(pInt_host).value + assert kernelValues[2].value == ctypes.c_float.from_address(pFloat_host).value + assert kernelValues[4].value == testStruct.from_address(pStruct_host).value + + (err,) = cudart.cudaFreeHost(pStruct_host) + ASSERT_DRV(err) + (err,) = cuda.cuStreamDestroy(stream) + ASSERT_DRV(err) + (err,) = cuda.cuModuleUnload(module) + ASSERT_DRV(err) + + +def test_kernelParams_buffer_protocol_numpy(device): + err, uvaSupported = cuda.cuDeviceGetAttribute( + cuda.CUdevice_attribute.CU_DEVICE_ATTRIBUTE_UNIFIED_ADDRESSING, device + ) + ASSERT_DRV(err) + + kernelString = """\ + struct testStruct { + int value; + }; + extern "C" __global__ + void testkernel(int i, int *pi, + float f, float *pf, + struct testStruct s, struct testStruct *ps) + { + *pi = i; + *pf = f; + ps->value = s.value; + } + """ + + module = common_nvrtc(kernelString, device) + + err, kernel = cuda.cuModuleGetFunction(module, b"testkernel") + ASSERT_DRV(err) + + err, stream = cuda.cuStreamCreate(0) + ASSERT_DRV(err) + + # testkernel kernel + testStruct = np.dtype([("value", np.int32)]) + + err, pInt_host = cudart.cudaHostAlloc(np.dtype(np.int32).itemsize, cudart.cudaHostAllocMapped) + ASSERT_DRV(err) + err, pFloat_host = cudart.cudaHostAlloc(np.dtype(np.float32).itemsize, cudart.cudaHostAllocMapped) + ASSERT_DRV(err) + err, pStruct_host = cudart.cudaHostAlloc(testStruct.itemsize, cudart.cudaHostAllocMapped) + ASSERT_DRV(err) + + # Get device pointer if UVM is not enabled + if uvaSupported: + kernelValues = ( + np.array(1, dtype=np.uint32), + np.array([pInt_host], dtype=np.uint64), + np.array(123.456, dtype=np.float32), + np.array([pFloat_host], dtype=np.uint64), + np.array([5], testStruct), + np.array([pStruct_host], dtype=np.uint64), + ) + else: + err, pInt_device = cudart.cudaHostGetDevicePointer(pInt_host, 0) + ASSERT_DRV(err) + err, pFloat_device = cudart.cudaHostGetDevicePointer(pFloat_host, 0) + ASSERT_DRV(err) + err, pStruct_device = cudart.cudaHostGetDevicePointer(pStruct_host, 0) + ASSERT_DRV(err) + kernelValues = ( + np.array(1, dtype=np.int32), + np.array([pInt_device], dtype=np.uint64), + np.array(123.456, dtype=np.float32), + np.array([pFloat_device], dtype=np.uint64), + np.array([5], testStruct), + np.array([pStruct_device], dtype=np.uint64), + ) + + packagedParams = np.array([arg.ctypes.data for arg in kernelValues], dtype=np.uint64) + (err,) = cuda.cuLaunchKernel( + kernel, + 1, + 1, + 1, # grid dim + 1, + 1, + 1, # block dim + 0, + stream, # shared mem and stream + packagedParams, + 0, + ) # arguments + ASSERT_DRV(err) + + # Validate kernel values + (err,) = cuda.cuStreamSynchronize(stream) + ASSERT_DRV(err) + + class numpy_address_wrapper: + def __init__(self, address, typestr): + self.__array_interface__ = {"data": (address, False), "typestr": typestr, "shape": (1,)} + + assert kernelValues[0] == np.array(numpy_address_wrapper(pInt_host, "; + .reg .b64 %rd<5>; + + + ld.param.u64 %rd1, [_Z6kernelPi_param_0]; + cvta.to.global.u64 %rd2, %rd1; + mov.u32 %r1, %tid.x; + mov.u32 %r2, %ctaid.x; + mov.u32 %r3, %ntid.x; + mad.lo.s32 %r4, %r2, %r3, %r1; + mul.wide.s32 %rd3, %r4, 4; + add.s64 %rd4, %rd2, %rd3; + ld.global.u32 %r5, [%rd4]; + add.s32 %r6, %r5, 1; + st.global.u32 [%rd4], %r6; + ret; + +}} +""" + +CODE = """ +int __device__ inc(int x) { + return x + 1; +} +""" + +# Base64 encoded TileIR generated by the toolshed/dump_cutile_b64.py script. +TILEIR_b64 = ( + "f1RpbGVJUgANAQAAgo0BCAECBgYBCwEECgCBAUQHBgQIEAAABgUMAQABBgUIEAALQwEICgEMAAYE" + "CBAAAwYFDAEABAYFCBAAD0MBCA4BEAAGBAgQAAYGBQwBAAcGBQgQABNDAQgSARQAMAUFBUIJDT4C" + "CgcEABkBFglCCRE+AgoHBAAcARYJAgoAABodQgkVZgEHBAAfIAEWCVwAAIQICADLy8vLy8vLg5oC" + "CMvLy8sBy8vLAAAAABfLy8vLy8vLBAAAAAAAAAAEAAAAAAAAAAQAAAAAAAAABAAAAAAAAAAEAAAA" + "AAAAAAQAAAAAAAAABAAAAAAAAAAEAAAAAAAAAAQAAAAAAAAABAAAAAAAAAAEAAAAAAAAAAQAAAAA" + "AAAABAAAAAAAAAAEAAAAAAAAAAUAAAAAAAAABgAAAAAAAAAGAAAAAAAAAAcAAAAAAAAABwAAAAAA" + "AAAIAAAAAAAAAAkAAAAAAAAACQAAAAAAAAAEAAAAAAAAAAnLy8sAAAAAAwAAAAUAAAAMAAAAEQAA" + "ABYAAAAbAAAAIAAAACUAAAACAAEBAQUBFwICAhcEAwMYBAQDAxkTBAMDGhEEAwMbEQQDAx0WBAMD" + "HgiFdATLy8sLy8vLAAAAAAEAAAACAAAAAwAAAAUAAAAIAAAACwAAABcAAAAYAAAALAAAADkAAAAA" + "AwcMAg0DAA0BABAJBAUFBAUFBAUFABEOAgEAAAAAAAAAgAEBAAAAAAAAAA8BEAAAAAgBAAAAAAAN" + "AgEQAAAAAAAAAIGIAQQFy8vLAAAAABIAAAAsAAAAPQAAAGoAAABkdW1wX2N1dGlsZV9iNjQucHkv" + "bG9jYWxob21lL2xvY2FsLXdhbmdtL3RveXZlY3Rvcl9hZGRfa2VybmVsL2xvY2FsaG9tZS9sb2Nh" + "bC13YW5nbS90b3kvZHVtcF9jdXRpbGVfYjY0LnB5c21fMTIwAA==" +) + + +def get_version() -> tuple[int, int]: + return nvfatbin.version() + + +@pytest.fixture(params=ARCHITECTURES) +def arch(request): + return request.param + + +@pytest.fixture(params=PTX_VERSIONS) +def ptx_version(request): + return request.param + + +@pytest.fixture +def PTX(arch, ptx_version): + return PTX_TEMPLATE.format(PTX_VERSION=ptx_version, ARCH=arch) + + +@pytest.fixture +def nvcc_smoke(tmpdir) -> str: + # TODO: Use cuda-pathfinder to locate nvcc on system. + nvcc = shutil.which("nvcc") + if nvcc is None: + pytest.skip("nvcc not found on PATH") + + # Smoke test: make sure nvcc is actually usable (toolkit + host compiler are set up), + # not merely present on PATH. + src = tmpdir / "nvcc_smoke.cu" + out = tmpdir / "nvcc_smoke.o" + with open(src, "w") as f: + f.write("") + try: + subprocess.run( # noqa: S603 + [nvcc, "-c", str(src), "-o", str(out)], + check=True, + capture_output=True, + ) + except subprocess.CalledProcessError as e: + stdout = (e.stdout or b"").decode(errors="replace") + stderr = (e.stderr or b"").decode(errors="replace") + pytest.skip( + "nvcc found on PATH but failed to compile a trivial input.\n" + f"command: {[nvcc, '-c', str(src), '-o', str(out)]!r}\n" + f"exit_code: {e.returncode}\n" + f"stdout:\n{stdout}\n" + f"stderr:\n{stderr}\n" + ) + + return nvcc + + +def _build_cubin(arch): + def CHECK_NVRTC(err): + if err != nvrtc.nvrtcResult.NVRTC_SUCCESS: + raise RuntimeError(repr(err)) + + err, program_handle = nvrtc.nvrtcCreateProgram(CODE.encode(), b"", 0, [], []) + CHECK_NVRTC(err) + err = nvrtc.nvrtcCompileProgram(program_handle, 1, [f"-arch={arch}".encode()])[0] + CHECK_NVRTC(err) + err, size = nvrtc.nvrtcGetCUBINSize(program_handle) + CHECK_NVRTC(err) + cubin = b" " * size + (err,) = nvrtc.nvrtcGetCUBIN(program_handle, cubin) + CHECK_NVRTC(err) + (err,) = nvrtc.nvrtcDestroyProgram(program_handle) + CHECK_NVRTC(err) + return cubin + + +@pytest.fixture +def CUBIN(arch): + return _build_cubin(arch) + + +# create a valid LTOIR input for testing +@pytest.fixture +def LTOIR(arch): + arch = arch.replace("sm", "compute") + + def CHECK_NVRTC(err): + if err != nvrtc.nvrtcResult.NVRTC_SUCCESS: + raise RuntimeError(repr(err)) + + empty_cplusplus_kernel = "__global__ void A() {}" + err, program_handle = nvrtc.nvrtcCreateProgram(empty_cplusplus_kernel.encode(), b"", 0, [], []) + CHECK_NVRTC(err) + err = nvrtc.nvrtcCompileProgram(program_handle, 1, [b"-dlto", f"-arch={arch}".encode()])[0] + CHECK_NVRTC(err) + err, size = nvrtc.nvrtcGetLTOIRSize(program_handle) + CHECK_NVRTC(err) + empty_kernel_ltoir = b" " * size + (err,) = nvrtc.nvrtcGetLTOIR(program_handle, empty_kernel_ltoir) + CHECK_NVRTC(err) + (err,) = nvrtc.nvrtcDestroyProgram(program_handle) + CHECK_NVRTC(err) + return empty_kernel_ltoir + + +@pytest.fixture +def OBJECT(arch, tmpdir, nvcc_smoke): + empty_cplusplus_kernel = "__global__ void A() {}" + with open(tmpdir / "object.cu", "w") as f: + f.write(empty_cplusplus_kernel) + + nvcc = nvcc_smoke + + # This is a test fixture that intentionally invokes a trusted tool (`nvcc`) to + # compile a temporary CUDA translation unit. + cmd = [nvcc, "-c", "-arch", arch, "-o", str(tmpdir / "object.o"), str(tmpdir / "object.cu")] + try: + subprocess.run( # noqa: S603 + cmd, + check=True, + capture_output=True, + ) + except subprocess.CalledProcessError as e: + stdout = (e.stdout or b"").decode(errors="replace") + stderr = (e.stderr or b"").decode(errors="replace") + raise RuntimeError( + "nvcc smoke test passed, but nvcc failed while compiling the test object.\n" + f"command: {cmd!r}\n" + f"exit_code: {e.returncode}\n" + f"stdout:\n{stdout}\n" + f"stderr:\n{stderr}\n" + ) from e + with open(tmpdir / "object.o", "rb") as f: + object = f.read() + + return object + + +@pytest.fixture +def TILEIR(tmpdir): + try: + binary_data = base64.b64decode(TILEIR_b64) + except binascii.Error as e: + raise ValueError( + "Base64 encoded TileIR is corrupted. Please regenerate the TileIR" + "by executing the toolshed/dump_cutile_b64.py script." + ) from e + return binary_data + + +@pytest.mark.parametrize("error_enum", nvfatbin.Result) +def test_get_error_string(error_enum): + es = nvfatbin.get_error_string(error_enum) + + if error_enum is nvfatbin.Result.SUCCESS: + assert es == "" + else: + assert es != "" + + +def test_nvfatbin_get_version(): + major, minor = nvfatbin.version() + assert major is not None + assert minor is not None + + +def test_nvfatbin_empty_create_and_destroy(): + handle = nvfatbin.create([], 0) + assert handle is not None + nvfatbin.destroy(handle) + + +def test_nvfatbin_invalid_input_create(): + with pytest.raises(nvfatbin.nvFatbinError, match="ERROR_UNRECOGNIZED_OPTION"): + nvfatbin.create(["--unsupported_option"], 1) + + +def test_nvfatbin_get_empty(): + handle = nvfatbin.create([], 0) + size = nvfatbin.size(handle) + + buffer = bytearray(size) + nvfatbin.get(handle, buffer) + + nvfatbin.destroy(handle) + + +def test_nvfatbin_add_ptx(PTX, arch): + arch_numeric = arch.split("_")[1] + + handle = nvfatbin.create([], 0) + nvfatbin.add_ptx(handle, PTX.encode(), len(PTX), arch_numeric, "add", f"-arch={arch}") + + buffer = bytearray(nvfatbin.size(handle)) + + nvfatbin.get(handle, buffer) + nvfatbin.destroy(handle) + + +def test_nvfatbin_add_cubin_ELF_SIZE_MISMATCH(): + cubin = _build_cubin("sm_80") + handle = nvfatbin.create([], 0) + with pytest.raises(nvfatbin.nvFatbinError, match="ERROR_ELF_ARCH_MISMATCH"): + nvfatbin.add_cubin(handle, cubin, len(cubin), "75", "inc") + + nvfatbin.destroy(handle) + + +def test_nvfatbin_add_cubin(CUBIN, arch): + arch_numeric = arch.split("_")[1] + + handle = nvfatbin.create([], 0) + nvfatbin.add_cubin(handle, CUBIN, len(CUBIN), arch_numeric, "inc") + + buffer = bytearray(nvfatbin.size(handle)) + + nvfatbin.get(handle, buffer) + nvfatbin.destroy(handle) + + +def test_nvfatbin_add_cubin_ELF_ARCH_MISMATCH(): + cubin = _build_cubin("sm_80") + handle = nvfatbin.create([], 0) + with pytest.raises(nvfatbin.nvFatbinError, match="ERROR_ELF_ARCH_MISMATCH"): + nvfatbin.add_cubin(handle, cubin, len(cubin), "75", "inc") + + nvfatbin.destroy(handle) + + +def test_nvdfatbin_add_ltoir(LTOIR, arch): + arch_numeric = arch.split("_")[1] + + handle = nvfatbin.create([], 0) + nvfatbin.add_ltoir(handle, LTOIR, len(LTOIR), arch_numeric, "inc", "") + + buffer = bytearray(nvfatbin.size(handle)) + + nvfatbin.get(handle, buffer) + nvfatbin.destroy(handle) + + +def test_nvfatbin_add_reloc(OBJECT): + handle = nvfatbin.create([], 0) + nvfatbin.add_reloc(handle, OBJECT, len(OBJECT)) + + buffer = bytearray(nvfatbin.size(handle)) + + nvfatbin.get(handle, buffer) + nvfatbin.destroy(handle) + + +@pytest.mark.skipif(get_version() < (13, 1), reason="TileIR API is not supported in CUDA < 13.1") +def test_nvfatbin_add_tile_ir(TILEIR): + handle = nvfatbin.create([], 0) + nvfatbin.add_tile_ir(handle, TILEIR, len(TILEIR), "VectorAdd", "") + + buffer = bytearray(nvfatbin.size(handle)) + + nvfatbin.get(handle, buffer) + nvfatbin.destroy(handle) diff --git a/cuda_bindings/tests/legacy_api/test_legacy_nvjitlink.py b/cuda_bindings/tests/legacy_api/test_legacy_nvjitlink.py new file mode 100644 index 00000000000..91d6ace1ce6 --- /dev/null +++ b/cuda_bindings/tests/legacy_api/test_legacy_nvjitlink.py @@ -0,0 +1,210 @@ +# SPDX-FileCopyrightText: Copyright (c) 2024-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +from contextlib import contextmanager + +import pytest + +from cuda.bindings import nvjitlink, nvrtc + + +@contextmanager +def nvjitlink_session(num_options, options): + """Create an nvJitLink handle and always destroy it (including on test failure).""" + handle = nvjitlink.create(num_options, options) + try: + yield handle + finally: + if handle != 0: + nvjitlink.destroy(handle) + + +# Establish a handful of compatible architectures and PTX versions to test with +ARCHITECTURES = ["sm_75", "sm_80", "sm_90", "sm_100"] +PTX_VERSIONS = ["6.4", "7.0", "8.5", "8.8"] + + +PTX_HEADER = """\ +.version {VERSION} +.target {ARCH} +.address_size 64 +""" + +PTX_KERNEL = """ +.visible .entry _Z6kernelPi( + .param .u64 _Z6kernelPi_param_0 +) +{ + .reg .pred %p<2>; + .reg .b32 %r<3>; + .reg .b64 %rd<3>; + + ld.param.u64 %rd1, [_Z6kernelPi_param_0]; + cvta.to.global.u64 %rd2, %rd1; + mov.u32 %r1, %tid.x; + st.global.u32 [%rd2+0], %r1; + ret; +} +""" + + +def _build_arch_ptx_parametrized_callable(): + av = tuple(zip(ARCHITECTURES, PTX_VERSIONS)) + return pytest.mark.parametrize( + ("arch", "ptx_bytes"), + [(a, (PTX_HEADER.format(VERSION=v, ARCH=a) + PTX_KERNEL).encode("utf-8")) for a, v in av], + ids=[f"{a}_{v}" for a, v in av], + ) + + +ARCH_PTX_PARAMETRIZED_CALLABLE = _build_arch_ptx_parametrized_callable() + + +def arch_ptx_parametrized(func): + return ARCH_PTX_PARAMETRIZED_CALLABLE(func) + + +def check_nvjitlink_usable(): + from cuda.bindings._internal import nvjitlink as inner_nvjitlink + + return inner_nvjitlink._inspect_function_pointer("__nvJitLinkVersion") != 0 + + +def check_nvjitlink_get_linked_ltoir_usable(): + from cuda.bindings._internal import nvjitlink as inner_nvjitlink + + return inner_nvjitlink._inspect_function_pointer("__nvJitLinkGetLinkedLTOIR") != 0 + + +pytestmark = pytest.mark.skipif( + not check_nvjitlink_usable(), reason="nvJitLink not usable, maybe not installed or too old (<12.3)" +) + + +# create a valid LTOIR input for testing +@pytest.fixture +def get_dummy_ltoir(): + def CHECK_NVRTC(err): + if err != nvrtc.nvrtcResult.NVRTC_SUCCESS: + raise RuntimeError(repr(err)) + + empty_cplusplus_kernel = "__global__ void A() {}" + err, program_handle = nvrtc.nvrtcCreateProgram(empty_cplusplus_kernel.encode(), b"", 0, [], []) + CHECK_NVRTC(err) + err = nvrtc.nvrtcCompileProgram(program_handle, 1, [b"-dlto"])[0] + CHECK_NVRTC(err) + err, size = nvrtc.nvrtcGetLTOIRSize(program_handle) + CHECK_NVRTC(err) + empty_kernel_ltoir = b" " * size + (err,) = nvrtc.nvrtcGetLTOIR(program_handle, empty_kernel_ltoir) + CHECK_NVRTC(err) + (err,) = nvrtc.nvrtcDestroyProgram(program_handle) + CHECK_NVRTC(err) + return empty_kernel_ltoir + + +def test_unrecognized_option_error(): + with pytest.raises(nvjitlink.nvJitLinkError, match="ERROR_UNRECOGNIZED_OPTION"): + nvjitlink.create(1, ["-fictitious_option"]) + + +def test_invalid_arch_error(): + with pytest.raises(nvjitlink.nvJitLinkError, match="ERROR_UNRECOGNIZED_OPTION"): + nvjitlink.create(1, ["-arch=sm_XX"]) + + +@pytest.mark.parametrize("option", ARCHITECTURES) +def test_create_and_destroy(option): + with nvjitlink_session(1, [f"-arch={option}"]) as handle: + assert handle != 0 + + +def test_create_and_destroy_bytes_options(): + with nvjitlink_session(1, [b"-arch=sm_80"]) as handle: + assert handle != 0 + + +@pytest.mark.parametrize("option", ARCHITECTURES) +def test_complete_empty(option): + with nvjitlink_session(1, [f"-arch={option}"]) as handle: + nvjitlink.complete(handle) + + +@arch_ptx_parametrized +def test_add_data(arch, ptx_bytes): + with nvjitlink_session(1, [f"-arch={arch}"]) as handle: + nvjitlink.add_data(handle, nvjitlink.InputType.ANY, ptx_bytes, len(ptx_bytes), "test_data") + nvjitlink.complete(handle) + + +@arch_ptx_parametrized +def test_add_file(arch, ptx_bytes, tmp_path): + with nvjitlink_session(1, [f"-arch={arch}"]) as handle: + file_path = tmp_path / "test_file.cubin" + file_path.write_bytes(ptx_bytes) + nvjitlink.add_file(handle, nvjitlink.InputType.ANY, str(file_path)) + nvjitlink.complete(handle) + + +@pytest.mark.parametrize("arch", ARCHITECTURES) +def test_get_error_log(arch): + with nvjitlink_session(1, [f"-arch={arch}"]) as handle: + nvjitlink.complete(handle) + log_size = nvjitlink.get_error_log_size(handle) + log = bytearray(log_size) + nvjitlink.get_error_log(handle, log) + assert len(log) == log_size + + +@arch_ptx_parametrized +def test_get_info_log(arch, ptx_bytes): + with nvjitlink_session(1, [f"-arch={arch}"]) as handle: + nvjitlink.add_data(handle, nvjitlink.InputType.ANY, ptx_bytes, len(ptx_bytes), "test_data") + nvjitlink.complete(handle) + log_size = nvjitlink.get_info_log_size(handle) + log = bytearray(log_size) + nvjitlink.get_info_log(handle, log) + assert len(log) == log_size + + +@arch_ptx_parametrized +def test_get_linked_cubin(arch, ptx_bytes): + with nvjitlink_session(1, [f"-arch={arch}"]) as handle: + nvjitlink.add_data(handle, nvjitlink.InputType.ANY, ptx_bytes, len(ptx_bytes), "test_data") + nvjitlink.complete(handle) + cubin_size = nvjitlink.get_linked_cubin_size(handle) + cubin = bytearray(cubin_size) + nvjitlink.get_linked_cubin(handle, cubin) + assert len(cubin) == cubin_size + + +@pytest.mark.parametrize("arch", ARCHITECTURES) +def test_get_linked_ptx(arch, get_dummy_ltoir): + with nvjitlink_session(3, [f"-arch={arch}", "-lto", "-ptx"]) as handle: + nvjitlink.add_data(handle, nvjitlink.InputType.LTOIR, get_dummy_ltoir, len(get_dummy_ltoir), "test_data") + nvjitlink.complete(handle) + ptx_size = nvjitlink.get_linked_ptx_size(handle) + ptx = bytearray(ptx_size) + nvjitlink.get_linked_ptx(handle, ptx) + assert len(ptx) == ptx_size + + +@pytest.mark.parametrize("arch", ARCHITECTURES) +@pytest.mark.skipif( + not check_nvjitlink_get_linked_ltoir_usable(), + reason="nvJitLinkGetLinkedLTOIR not available in installed nvJitLink", +) +def test_get_linked_ltoir(arch, get_dummy_ltoir): + with nvjitlink_session(2, [f"-arch={arch}", "-lto"]) as handle: + nvjitlink.add_data(handle, nvjitlink.InputType.LTOIR, get_dummy_ltoir, len(get_dummy_ltoir), "test_data") + nvjitlink.complete(handle) + ltoir_size = nvjitlink.get_linked_ltoir_size(handle) + ltoir = bytearray(ltoir_size) + nvjitlink.get_linked_ltoir(handle, ltoir) + assert len(ltoir) == ltoir_size + + +def test_package_version(): + ver = nvjitlink.version() + assert len(ver) == 2 + assert ver >= (12, 0) diff --git a/cuda_bindings/tests/legacy_api/test_legacy_nvrtc.py b/cuda_bindings/tests/legacy_api/test_legacy_nvrtc.py new file mode 100644 index 00000000000..90ff6766d41 --- /dev/null +++ b/cuda_bindings/tests/legacy_api/test_legacy_nvrtc.py @@ -0,0 +1,37 @@ +# SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +import pytest + +from cuda.bindings import nvrtc + + +def ASSERT_DRV(err): + if isinstance(err, nvrtc.nvrtcResult): + if err != nvrtc.nvrtcResult.NVRTC_SUCCESS: + raise RuntimeError(f"Nvrtc Error: {err}") + else: + raise RuntimeError(f"Unknown error type: {err}") + + +def nvrtcVersionLessThan(major, minor): + err, major_version, minor_version = nvrtc.nvrtcVersion() + ASSERT_DRV(err) + return major_version < major or (major == major_version and minor_version < minor) + + +@pytest.mark.skipif(nvrtcVersionLessThan(11, 3), reason="When nvrtcGetSupportedArchs was introduced") +def test_nvrtcGetSupportedArchs(): + err, supportedArchs = nvrtc.nvrtcGetSupportedArchs() + ASSERT_DRV(err) + assert len(supportedArchs) != 0 + + +@pytest.mark.skipif(nvrtcVersionLessThan(12, 1), reason="Preempt Segmentation Fault (see #499)") +def test_nvrtcGetLoweredName_failure(): + err, name = nvrtc.nvrtcGetLoweredName(None, b"I'm an elevated name!") + assert err == nvrtc.nvrtcResult.NVRTC_ERROR_INVALID_PROGRAM + assert name is None + err, name = nvrtc.nvrtcGetLoweredName(0, b"I'm another elevated name!") + assert err == nvrtc.nvrtcResult.NVRTC_ERROR_INVALID_PROGRAM + assert name is None diff --git a/cuda_bindings/tests/test_kernelParams.py b/cuda_bindings/tests/test_kernelParams.py index 555d6a7284c..793bbdebd68 100644 --- a/cuda_bindings/tests/test_kernelParams.py +++ b/cuda_bindings/tests/test_kernelParams.py @@ -6,8 +6,8 @@ import numpy as np import pytest +import cuda.bindings._v2.nvrtc as nvrtc import cuda.bindings.driver as cuda -import cuda.bindings.nvrtc as nvrtc import cuda.bindings.runtime as cudart @@ -18,9 +18,6 @@ def ASSERT_DRV(err): elif isinstance(err, cudart.cudaError_t): if err != cudart.cudaError_t.cudaSuccess: raise RuntimeError(f"Cudart Error: {err}") - elif isinstance(err, nvrtc.nvrtcResult): - if err != nvrtc.nvrtcResult.NVRTC_SUCCESS: - raise RuntimeError(f"Nvrtc Error: {err}") else: raise RuntimeError(f"Unknown error type: {err}") @@ -30,39 +27,24 @@ def common_nvrtc(allKernelStrings, dev): ASSERT_DRV(err) err, minor = cuda.cuDeviceGetAttribute(cuda.CUdevice_attribute.CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, dev) ASSERT_DRV(err) - err, _, nvrtc_minor = nvrtc.nvrtcVersion() - ASSERT_DRV(err) + _, nvrtc_minor = nvrtc.version() use_cubin = nvrtc_minor >= 1 prefix = "sm" if use_cubin else "compute" arch_arg = bytes(f"--gpu-architecture={prefix}_{major}{minor}", "ascii") - err, prog = nvrtc.nvrtcCreateProgram(str.encode(allKernelStrings), b"allKernelStrings.cu", 0, None, None) - ASSERT_DRV(err) + prog = nvrtc.create_program(str.encode(allKernelStrings), b"allKernelStrings.cu") opts = (b"--fmad=false", arch_arg) - (err,) = nvrtc.nvrtcCompileProgram(prog, len(opts), opts) + nvrtc.compile_program(prog, opts) - err_log, logSize = nvrtc.nvrtcGetProgramLogSize(prog) - ASSERT_DRV(err_log) - log = b" " * logSize - (err_log,) = nvrtc.nvrtcGetProgramLog(prog, log) - ASSERT_DRV(err_log) + log = nvrtc.get_program_log(prog) result = log.decode() if len(result) > 1: print(result) - ASSERT_DRV(err) if use_cubin: - err, dataSize = nvrtc.nvrtcGetCUBINSize(prog) - ASSERT_DRV(err) - data = b" " * dataSize - (err,) = nvrtc.nvrtcGetCUBIN(prog, data) - ASSERT_DRV(err) + data = nvrtc.get_cubin(prog) else: - err, dataSize = nvrtc.nvrtcGetPTXSize(prog) - ASSERT_DRV(err) - data = b" " * dataSize - (err,) = nvrtc.nvrtcGetPTX(prog, data) - ASSERT_DRV(err) + data = nvrtc.get_ptx(prog) err, module = cuda.cuModuleLoadData(np.char.array(data)) ASSERT_DRV(err) diff --git a/cuda_bindings/tests/test_nvfatbin.py b/cuda_bindings/tests/test_nvfatbin.py index 32b981d93da..1bb8d4fcc53 100644 --- a/cuda_bindings/tests/test_nvfatbin.py +++ b/cuda_bindings/tests/test_nvfatbin.py @@ -9,7 +9,8 @@ import pytest -from cuda.bindings import nvfatbin, nvrtc +from cuda.bindings import nvfatbin +from cuda.bindings._v2 import nvrtc ARCHITECTURES = ["sm_75", "sm_80", "sm_90", "sm_100"] PTX_VERSIONS = ["6.4", "7.0", "8.5", "8.8"] @@ -122,22 +123,9 @@ def nvcc_smoke(tmpdir) -> str: def _build_cubin(arch): - def CHECK_NVRTC(err): - if err != nvrtc.nvrtcResult.NVRTC_SUCCESS: - raise RuntimeError(repr(err)) - - err, program_handle = nvrtc.nvrtcCreateProgram(CODE.encode(), b"", 0, [], []) - CHECK_NVRTC(err) - err = nvrtc.nvrtcCompileProgram(program_handle, 1, [f"-arch={arch}".encode()])[0] - CHECK_NVRTC(err) - err, size = nvrtc.nvrtcGetCUBINSize(program_handle) - CHECK_NVRTC(err) - cubin = b" " * size - (err,) = nvrtc.nvrtcGetCUBIN(program_handle, cubin) - CHECK_NVRTC(err) - (err,) = nvrtc.nvrtcDestroyProgram(program_handle) - CHECK_NVRTC(err) - return cubin + program_handle = nvrtc.create_program(CODE.encode(), b"") + nvrtc.compile_program(program_handle, [f"-arch={arch}".encode()]) + return nvrtc.get_cubin(program_handle) @pytest.fixture @@ -149,24 +137,10 @@ def CUBIN(arch): @pytest.fixture def LTOIR(arch): arch = arch.replace("sm", "compute") - - def CHECK_NVRTC(err): - if err != nvrtc.nvrtcResult.NVRTC_SUCCESS: - raise RuntimeError(repr(err)) - empty_cplusplus_kernel = "__global__ void A() {}" - err, program_handle = nvrtc.nvrtcCreateProgram(empty_cplusplus_kernel.encode(), b"", 0, [], []) - CHECK_NVRTC(err) - err = nvrtc.nvrtcCompileProgram(program_handle, 1, [b"-dlto", f"-arch={arch}".encode()])[0] - CHECK_NVRTC(err) - err, size = nvrtc.nvrtcGetLTOIRSize(program_handle) - CHECK_NVRTC(err) - empty_kernel_ltoir = b" " * size - (err,) = nvrtc.nvrtcGetLTOIR(program_handle, empty_kernel_ltoir) - CHECK_NVRTC(err) - (err,) = nvrtc.nvrtcDestroyProgram(program_handle) - CHECK_NVRTC(err) - return empty_kernel_ltoir + program_handle = nvrtc.create_program(empty_cplusplus_kernel.encode(), b"") + nvrtc.compile_program(program_handle, [b"-dlto", f"-arch={arch}".encode()]) + return nvrtc.get_ltoir(program_handle) @pytest.fixture diff --git a/cuda_bindings/tests/test_nvjitlink.py b/cuda_bindings/tests/test_nvjitlink.py index 91d6ace1ce6..3fa6837a9a0 100644 --- a/cuda_bindings/tests/test_nvjitlink.py +++ b/cuda_bindings/tests/test_nvjitlink.py @@ -5,7 +5,8 @@ import pytest -from cuda.bindings import nvjitlink, nvrtc +from cuda.bindings import nvjitlink +from cuda.bindings._v2 import nvrtc @contextmanager @@ -84,23 +85,10 @@ def check_nvjitlink_get_linked_ltoir_usable(): # create a valid LTOIR input for testing @pytest.fixture def get_dummy_ltoir(): - def CHECK_NVRTC(err): - if err != nvrtc.nvrtcResult.NVRTC_SUCCESS: - raise RuntimeError(repr(err)) - empty_cplusplus_kernel = "__global__ void A() {}" - err, program_handle = nvrtc.nvrtcCreateProgram(empty_cplusplus_kernel.encode(), b"", 0, [], []) - CHECK_NVRTC(err) - err = nvrtc.nvrtcCompileProgram(program_handle, 1, [b"-dlto"])[0] - CHECK_NVRTC(err) - err, size = nvrtc.nvrtcGetLTOIRSize(program_handle) - CHECK_NVRTC(err) - empty_kernel_ltoir = b" " * size - (err,) = nvrtc.nvrtcGetLTOIR(program_handle, empty_kernel_ltoir) - CHECK_NVRTC(err) - (err,) = nvrtc.nvrtcDestroyProgram(program_handle) - CHECK_NVRTC(err) - return empty_kernel_ltoir + program_handle = nvrtc.create_program(empty_cplusplus_kernel.encode(), b"") + nvrtc.compile_program(program_handle, [b"-dlto"]) + return nvrtc.get_ltoir(program_handle) def test_unrecognized_option_error(): diff --git a/cuda_bindings/tests/test_nvrtc.py b/cuda_bindings/tests/test_nvrtc.py index 90ff6766d41..26eea71a100 100644 --- a/cuda_bindings/tests/test_nvrtc.py +++ b/cuda_bindings/tests/test_nvrtc.py @@ -3,35 +3,23 @@ import pytest -from cuda.bindings import nvrtc +from cuda.bindings._v2 import nvrtc -def ASSERT_DRV(err): - if isinstance(err, nvrtc.nvrtcResult): - if err != nvrtc.nvrtcResult.NVRTC_SUCCESS: - raise RuntimeError(f"Nvrtc Error: {err}") - else: - raise RuntimeError(f"Unknown error type: {err}") - - -def nvrtcVersionLessThan(major, minor): - err, major_version, minor_version = nvrtc.nvrtcVersion() - ASSERT_DRV(err) +def nvrtc_version_less_than(major, minor): + major_version, minor_version = nvrtc.version() return major_version < major or (major == major_version and minor_version < minor) -@pytest.mark.skipif(nvrtcVersionLessThan(11, 3), reason="When nvrtcGetSupportedArchs was introduced") -def test_nvrtcGetSupportedArchs(): - err, supportedArchs = nvrtc.nvrtcGetSupportedArchs() - ASSERT_DRV(err) - assert len(supportedArchs) != 0 +@pytest.mark.skipif(nvrtc_version_less_than(11, 3), reason="When nvrtcGetSupportedArchs was introduced") +def test_get_supported_archs(): + supported_archs = nvrtc.get_supported_archs() + assert len(supported_archs) != 0 -@pytest.mark.skipif(nvrtcVersionLessThan(12, 1), reason="Preempt Segmentation Fault (see #499)") -def test_nvrtcGetLoweredName_failure(): - err, name = nvrtc.nvrtcGetLoweredName(None, b"I'm an elevated name!") - assert err == nvrtc.nvrtcResult.NVRTC_ERROR_INVALID_PROGRAM - assert name is None - err, name = nvrtc.nvrtcGetLoweredName(0, b"I'm another elevated name!") - assert err == nvrtc.nvrtcResult.NVRTC_ERROR_INVALID_PROGRAM - assert name is None +@pytest.mark.skipif(nvrtc_version_less_than(12, 1), reason="Preempt Segmentation Fault (see #499)") +def test_get_lowered_name_failure(): + with pytest.raises(nvrtc.InvalidProgramError): + nvrtc.get_lowered_name(0, b"I'm an elevated name!") + with pytest.raises(nvrtc.InvalidProgramError): + nvrtc.get_lowered_name(0, b"I'm another elevated name!")