Merge branch 'main' into sg_cufile_bindings

Author: leofang

.github/workflows/build-docs.yml

@@ -56,7 +56,7 @@ jobs:
       # is resolved
 
       - name: Set up miniforge
-        uses: conda-incubator/setup-miniconda@505e6394dae86d6a5c7fbb6e3fb8938e3e863830  # v3.1.1
+        uses: conda-incubator/setup-miniconda@835234971496cad1653abb28a638a281cf32541f  # v3.2.0
         with:
           activate-environment: cuda-python-docs
           environment-file: ./cuda_python/docs/environment-docs.yml

.github/workflows/build-wheel.yml

@@ -71,7 +71,7 @@ jobs:
           env
       
       - name: Build cuda.core wheel
-        uses: pypa/cibuildwheel@faf86a6ed7efa889faf6996aa23820831055001a  # v2.23.3
+        uses: pypa/cibuildwheel@5f22145df44122af0f5a201f93cf0207171beca7  # v3.0.0
         env:
           CIBW_BUILD: ${{ env.CIBW_BUILD }}
           CIBW_ARCHS_LINUX: "native"
@@ -112,7 +112,7 @@ jobs:
           cuda-version: ${{ inputs.cuda-version }}
 
       - name: Build cuda.bindings wheel
-        uses: pypa/cibuildwheel@faf86a6ed7efa889faf6996aa23820831055001a  # v2.23.3
+        uses: pypa/cibuildwheel@5f22145df44122af0f5a201f93cf0207171beca7  # v3.0.0
         env:
           CIBW_BUILD: ${{ env.CIBW_BUILD }}
           CIBW_ARCHS_LINUX: "native"
@@ -185,7 +185,8 @@ jobs:
         id: setup-python2
         uses: actions/setup-python@a26af69be951a213d495a4c3e4e4022e16d87065  # v5.6.0
         with:
-          python-version: ${{ matrix.python-version }}
+          # workaround for actions/runner-images#12377 (the cached 3.13.4 is buggy on Windows)
+          python-version: ${{ matrix.python-version == '3.13' && '3.13.5' || matrix.python-version }}
 
       - name: Set up Python include paths
         run: |

.github/workflows/codeql.yml

@@ -31,13 +31,13 @@ jobs:
       uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683  # v4.2.2
 
     - name: Initialize CodeQL
-      uses: github/codeql-action/init@ff0a06e83cb2de871e5a09832bc6a81e7276941f  # v3.28.18
+      uses: github/codeql-action/init@ce28f5bb42b7a9f2c824e633a3f6ee835bab6858  # v3.29.0
       with:
         languages: ${{ matrix.language }}
         build-mode: ${{ matrix.build-mode }}
         queries: security-extended
 
     - name: Perform CodeQL Analysis
-      uses: github/codeql-action/analyze@ff0a06e83cb2de871e5a09832bc6a81e7276941f  # v3.28.18
+      uses: github/codeql-action/analyze@ce28f5bb42b7a9f2c824e633a3f6ee835bab6858  # v3.29.0
       with:
         category: "/language:${{matrix.language}}"

.github/workflows/release.yml

@@ -124,9 +124,13 @@ jobs:
           mkdir dist
           for p in ${{ inputs.component }}*
           do
+            # exclude cython test artifacts
+            if [[ "${p}" == *-tests ]]; then
+              continue
+            fi
             mv ${p}/*.whl dist/
           done
-          rmdir ${{ inputs.component }}*
+          rm -rf ${{ inputs.component }}*
 
       - name: Publish package distributions to PyPI
         if: ${{ inputs.wheel-dst == 'pypi' }}

README.md

@@ -4,11 +4,10 @@ CUDA Python is the home for accessing NVIDIA’s CUDA platform from Python. It c
 
 * [cuda.core](https://nvidia.github.io/cuda-python/cuda-core/latest): Pythonic access to CUDA Runtime and other core functionalities
 * [cuda.bindings](https://nvidia.github.io/cuda-python/cuda-bindings/latest): Low-level Python bindings to CUDA C APIs
-* [cuda.cooperative](https://nvidia.github.io/cccl/cuda_cooperative/): A Python package providing CCCL's reusable block-wide and warp-wide *device* primitives for use within Numba CUDA kernels
-* [cuda.parallel](https://nvidia.github.io/cccl/cuda_parallel/): A Python package for easy access to CCCL's highly efficient and customizable parallel algorithms, like `sort`, `scan`, `reduce`, `transform`, etc, that are callable on the *host*
+* [cuda.cccl.cooperative](https://nvidia.github.io/cccl/cuda_cooperative/): A Python module providing CCCL's reusable block-wide and warp-wide *device* primitives for use within Numba CUDA kernels
+* [cuda.cccl.parallel](https://nvidia.github.io/cccl/cuda_parallel/): A Python module for easy access to CCCL's highly efficient and customizable parallel algorithms, like `sort`, `scan`, `reduce`, `transform`, etc, that are callable on the *host*
 * [numba.cuda](https://nvidia.github.io/numba-cuda/): Numba's target for CUDA GPU programming by directly compiling a restricted subset of Python code into CUDA kernels and device functions following the CUDA execution model.
-
-For access to NVIDIA CPU & GPU Math Libraries, please refer to [nvmath-python](https://docs.nvidia.com/cuda/nvmath-python/latest).
+* [nvmath-python](https://docs.nvidia.com/cuda/nvmath-python/latest): Pythonic access to NVIDIA CPU & GPU Math Libraries, with both [*host*](https://docs.nvidia.com/cuda/nvmath-python/latest/overview.html#host-apis) and [*device* (nvmath.device)](https://docs.nvidia.com/cuda/nvmath-python/latest/overview.html#device-apis) APIs. It also provides low-level Python bindings to host C APIs ([nvmath.bindings](https://docs.nvidia.com/cuda/nvmath-python/latest/bindings/index.html)).
 
 CUDA Python is currently undergoing an overhaul to improve existing and bring up new components. All of the previously available functionalities from the `cuda-python` package will continue to be available, please refer to the [cuda.bindings](https://nvidia.github.io/cuda-python/cuda-bindings/latest) documentation for installation guide and further detail.
 

ci/tools/setup-sanitizer

@@ -12,11 +12,10 @@ set -euo pipefail
 if [[ "${SETUP_SANITIZER}" == 1 ]]; then
   COMPUTE_SANITIZER="${CUDA_HOME}/bin/compute-sanitizer"
   COMPUTE_SANITIZER_VERSION=$(${COMPUTE_SANITIZER} --version | grep -Eo "[0-9]{4}\.[0-9]\.[0-9]" | sed -e 's/\.//g')
-  SANITIZER_CMD="${COMPUTE_SANITIZER} --target-processes=all --launch-timeout=0 --tool=memcheck --error-exitcode=1"
+  SANITIZER_CMD="${COMPUTE_SANITIZER} --target-processes=all --launch-timeout=0 --tool=memcheck --error-exitcode=1 --report-api-errors=no"
   if [[ "$COMPUTE_SANITIZER_VERSION" -ge 202111 ]]; then
     SANITIZER_CMD="${SANITIZER_CMD} --padding=32"
   fi
-  echo "CUDA_PYTHON_TESTING_WITH_COMPUTE_SANITIZER=1" >> $GITHUB_ENV
 else
   SANITIZER_CMD=""
 fi

cuda_bindings/cuda/bindings/_internal/nvvm.pxd

@@ -11,6 +11,7 @@ from ..cynvvm cimport *
 # Wrapper functions
 ###############################################################################
 
+cdef const char* _nvvmGetErrorString(nvvmResult result) except?NULL nogil
 cdef nvvmResult _nvvmVersion(int* major, int* minor) except?_NVVMRESULT_INTERNAL_LOADING_ERROR nogil
 cdef nvvmResult _nvvmIRVersion(int* majorIR, int* minorIR, int* majorDbg, int* minorDbg) except?_NVVMRESULT_INTERNAL_LOADING_ERROR nogil
 cdef nvvmResult _nvvmCreateProgram(nvvmProgram* prog) except?_NVVMRESULT_INTERNAL_LOADING_ERROR nogil

cuda_bindings/cuda/bindings/_internal/nvvm_linux.pyx

@@ -36,6 +36,7 @@ cdef extern from "<dlfcn.h>" nogil:
 cdef bint __py_nvvm_init = False
 cdef void* __cuDriverGetVersion = NULL
 
+cdef void* __nvvmGetErrorString = NULL
 cdef void* __nvvmVersion = NULL
 cdef void* __nvvmIRVersion = NULL
 cdef void* __nvvmCreateProgram = NULL
@@ -82,6 +83,13 @@ cdef int _check_or_init_nvvm() except -1 nogil:
     handle = NULL
 
     # Load function
+    global __nvvmGetErrorString
+    __nvvmGetErrorString = dlsym(RTLD_DEFAULT, 'nvvmGetErrorString')
+    if __nvvmGetErrorString == NULL:
+        if handle == NULL:
+            handle = load_library(driver_ver)
+        __nvvmGetErrorString = dlsym(handle, 'nvvmGetErrorString')
+
     global __nvvmVersion
     __nvvmVersion = dlsym(RTLD_DEFAULT, 'nvvmVersion')
     if __nvvmVersion == NULL:
@@ -181,6 +189,9 @@ cpdef dict _inspect_function_pointers():
     _check_or_init_nvvm()
     cdef dict data = {}
 
+    global __nvvmGetErrorString
+    data["__nvvmGetErrorString"] = <intptr_t>__nvvmGetErrorString
+
     global __nvvmVersion
     data["__nvvmVersion"] = <intptr_t>__nvvmVersion
 
@@ -232,6 +243,16 @@ cpdef _inspect_function_pointer(str name):
 # Wrapper functions
 ###############################################################################
 
+cdef const char* _nvvmGetErrorString(nvvmResult result) except?NULL nogil:
+    global __nvvmGetErrorString
+    _check_or_init_nvvm()
+    if __nvvmGetErrorString == NULL:
+        with gil:
+            raise FunctionNotFoundError("function nvvmGetErrorString is not found")
+    return (<const char* (*)(nvvmResult) noexcept nogil>__nvvmGetErrorString)(
+        result)
+
+
 cdef nvvmResult _nvvmVersion(int* major, int* minor) except?_NVVMRESULT_INTERNAL_LOADING_ERROR nogil:
     global __nvvmVersion
     _check_or_init_nvvm()

cuda_bindings/cuda/bindings/_internal/nvvm_windows.pyx

@@ -23,6 +23,7 @@ LOAD_LIBRARY_SEARCH_DLL_LOAD_DIR = 0x00000100
 cdef bint __py_nvvm_init = False
 cdef void* __cuDriverGetVersion = NULL
 
+cdef void* __nvvmGetErrorString = NULL
 cdef void* __nvvmVersion = NULL
 cdef void* __nvvmIRVersion = NULL
 cdef void* __nvvmCreateProgram = NULL
@@ -62,6 +63,12 @@ cdef int _check_or_init_nvvm() except -1 nogil:
         handle = path_finder._load_nvidia_dynamic_library("nvvm").handle
 
         # Load function
+        global __nvvmGetErrorString
+        try:
+            __nvvmGetErrorString = <void*><intptr_t>win32api.GetProcAddress(handle, 'nvvmGetErrorString')
+        except:
+            pass
+
         global __nvvmVersion
         try:
             __nvvmVersion = <void*><intptr_t>win32api.GetProcAddress(handle, 'nvvmVersion')
@@ -149,6 +156,9 @@ cpdef dict _inspect_function_pointers():
     _check_or_init_nvvm()
     cdef dict data = {}
 
+    global __nvvmGetErrorString
+    data["__nvvmGetErrorString"] = <intptr_t>__nvvmGetErrorString
+
     global __nvvmVersion
     data["__nvvmVersion"] = <intptr_t>__nvvmVersion
 
@@ -200,6 +210,16 @@ cpdef _inspect_function_pointer(str name):
 # Wrapper functions
 ###############################################################################
 
+cdef const char* _nvvmGetErrorString(nvvmResult result) except?NULL nogil:
+    global __nvvmGetErrorString
+    _check_or_init_nvvm()
+    if __nvvmGetErrorString == NULL:
+        with gil:
+            raise FunctionNotFoundError("function nvvmGetErrorString is not found")
+    return (<const char* (*)(nvvmResult) noexcept nogil>__nvvmGetErrorString)(
+        result)
+
+
 cdef nvvmResult _nvvmVersion(int* major, int* minor) except?_NVVMRESULT_INTERNAL_LOADING_ERROR nogil:
     global __nvvmVersion
     _check_or_init_nvvm()

cuda_bindings/cuda/bindings/cynvvm.pxd

@@ -33,6 +33,7 @@ ctypedef void* nvvmProgram 'nvvmProgram'
 # Functions
 ###############################################################################
 
+cdef const char* nvvmGetErrorString(nvvmResult result) except?NULL nogil
 cdef nvvmResult nvvmVersion(int* major, int* minor) except?_NVVMRESULT_INTERNAL_LOADING_ERROR nogil
 cdef nvvmResult nvvmIRVersion(int* majorIR, int* minorIR, int* majorDbg, int* minorDbg) except?_NVVMRESULT_INTERNAL_LOADING_ERROR nogil
 cdef nvvmResult nvvmCreateProgram(nvvmProgram* prog) except?_NVVMRESULT_INTERNAL_LOADING_ERROR nogil

cuda_bindings/cuda/bindings/cynvvm.pyx

@@ -11,6 +11,10 @@ from ._internal cimport nvvm as _nvvm
 # Wrapper functions
 ###############################################################################
 
+cdef const char* nvvmGetErrorString(nvvmResult result) except?NULL nogil:
+    return _nvvm._nvvmGetErrorString(result)
+
+
 cdef nvvmResult nvvmVersion(int* major, int* minor) except?_NVVMRESULT_INTERNAL_LOADING_ERROR nogil:
     return _nvvm._nvvmVersion(major, minor)
 

cuda_bindings/cuda/bindings/nvvm.pxd

@@ -27,6 +27,7 @@ ctypedef nvvmResult _Result
 # Functions
 ###############################################################################
 
+cpdef str get_error_string(int result)
 cpdef tuple version()
 cpdef tuple ir_version()
 cpdef intptr_t create_program() except? 0

cuda_bindings/cuda/bindings/nvvm.pyx

@@ -73,6 +73,19 @@ cpdef destroy_program(intptr_t prog):
     check_status(status)
 
 
+cpdef str get_error_string(int result):
+    """Get the message string for the given ``nvvmResult`` code.
+
+    Args:
+        result (Result): NVVM API result code.
+
+    .. seealso:: `nvvmGetErrorString`
+    """
+    cdef bytes _output_
+    _output_ = nvvmGetErrorString(<_Result>result)
+    return _output_.decode()
+
+
 cpdef tuple version():
     """Get the NVVM version.
 

cuda_bindings/docs/source/environment_variables.md

@@ -11,8 +11,3 @@
 ## Runtime Environment Variables
 
 - `CUDA_PYTHON_CUDA_PER_THREAD_DEFAULT_STREAM` : When set to 1, the default stream is the per-thread default stream. When set to 0, the default stream is the legacy default stream. This defaults to 0, for the legacy default stream. See [Stream Synchronization Behavior](https://docs.nvidia.com/cuda/cuda-runtime-api/stream-sync-behavior.html) for an explanation of the legacy and per-thread default streams.
-
-
-## Test-Time Environment Variables
-
-- `CUDA_PYTHON_TESTING_WITH_COMPUTE_SANITIZER` : When set to 1, tests are skipped that would cause [compute-sanitizer](https://docs.nvidia.com/compute-sanitizer/ComputeSanitizer/index.html) to raise an error.

cuda_bindings/tests/conftest.py

@@ -1,15 +1,9 @@
 # Copyright 2025 NVIDIA Corporation.  All rights reserved.
 # SPDX-License-Identifier: LicenseRef-NVIDIA-SOFTWARE-LICENSE
 
-import os
 
 import pytest
 
-skipif_testing_with_compute_sanitizer = pytest.mark.skipif(
-    os.environ.get("CUDA_PYTHON_TESTING_WITH_COMPUTE_SANITIZER", "0") == "1",
-    reason="The compute-sanitizer is running, and this test causes an API error.",
-)
-
 
 def pytest_configure(config):
     config.custom_info = []

cuda_bindings/tests/test_cuda.py

@@ -7,7 +7,6 @@
 
 import numpy as np
 import pytest
-from conftest import skipif_testing_with_compute_sanitizer
 
 import cuda.cuda as cuda
 import cuda.cudart as cudart
@@ -80,7 +79,6 @@ def test_cuda_memcpy():
     assert err == cuda.CUresult.CUDA_SUCCESS
 
 
-@skipif_testing_with_compute_sanitizer
 def test_cuda_array():
     (err,) = cuda.cuInit(0)
     assert err == cuda.CUresult.CUDA_SUCCESS
@@ -234,7 +232,6 @@ def test_cuda_uuid_list_access():
     assert err == cuda.CUresult.CUDA_SUCCESS
 
 
-@skipif_testing_with_compute_sanitizer
 def test_cuda_cuModuleLoadDataEx():
     (err,) = cuda.cuInit(0)
     assert err == cuda.CUresult.CUDA_SUCCESS
@@ -622,7 +619,6 @@ def test_cuda_coredump_attr():
     assert err == cuda.CUresult.CUDA_SUCCESS
 
 
-@skipif_testing_with_compute_sanitizer
 def test_get_error_name_and_string():
     (err,) = cuda.cuInit(0)
     assert err == cuda.CUresult.CUDA_SUCCESS
@@ -952,7 +948,6 @@ def test_CUmemDecompressParams_st():
     assert int(desc.dstActBytes) == 0
 
 
-@skipif_testing_with_compute_sanitizer
 def test_all_CUresult_codes():
     max_code = int(max(cuda.CUresult))
     # Smoke test. CUDA_ERROR_UNKNOWN = 999, but intentionally using literal value.
@@ -985,21 +980,18 @@ def test_all_CUresult_codes():
     assert num_good >= 76  # CTK 11.0.3_450.51.06
 
 
-@skipif_testing_with_compute_sanitizer
 def test_cuKernelGetName_failure():
     err, name = cuda.cuKernelGetName(0)
     assert err == cuda.CUresult.CUDA_ERROR_INVALID_VALUE
     assert name is None
 
 
-@skipif_testing_with_compute_sanitizer
 def test_cuFuncGetName_failure():
     err, name = cuda.cuFuncGetName(0)
     assert err == cuda.CUresult.CUDA_ERROR_INVALID_VALUE
     assert name is None
 
 
-@skipif_testing_with_compute_sanitizer
 @pytest.mark.skipif(
     driverVersionLessThan(12080) or not supportsCudaAPI("cuCheckpointProcessGetState"),
     reason="When API was introduced",

cuda_bindings/tests/test_cudart.py

@@ -6,7 +6,6 @@
 
 import numpy as np
 import pytest
-from conftest import skipif_testing_with_compute_sanitizer
 
 import cuda.cuda as cuda
 import cuda.cudart as cudart
@@ -67,7 +66,6 @@ def test_cudart_memcpy():
     assertSuccess(err)
 
 
-@skipif_testing_with_compute_sanitizer
 def test_cudart_hostRegister():
     # Use hostRegister API to check for correct enum return values
     page_size = 80

cuda_bindings/tests/test_nvjitlink.py

@@ -66,12 +66,13 @@ def check_nvjitlink_usable():
 def get_dummy_ltoir():
     def CHECK_NVRTC(err):
         if err != nvrtc.nvrtcResult.NVRTC_SUCCESS:
-            raise RuntimeError(f"Nvrtc Error: {err}")
+            raise RuntimeError(repr(err))
 
     empty_cplusplus_kernel = "__global__ void A() {}"
     err, program_handle = nvrtc.nvrtcCreateProgram(empty_cplusplus_kernel.encode(), b"", 0, [], [])
     CHECK_NVRTC(err)
-    nvrtc.nvrtcCompileProgram(program_handle, 1, [b"-dlto"])
+    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

cuda_bindings/tests/test_nvvm.py

@@ -194,6 +194,20 @@ def get_program_log(prog):
     return buffer.decode(errors="backslashreplace")
 
 
+def test_get_error_string():
+    num_success = 0
+    num_errors = 0
+    for enum_obj in nvvm.Result:
+        es = nvvm.get_error_string(enum_obj)
+        if enum_obj is nvvm.Result.SUCCESS:
+            num_success += 1
+        else:
+            assert es.startswith("NVVM_ERROR")
+            num_errors += 1
+    assert num_success == 1
+    assert num_errors > 1  # smoke check is sufficient
+
+
 def test_nvvm_version():
     ver = nvvm.version()
     assert len(ver) == 2

cuda_core/cuda/core/_version.py

@@ -2,4 +2,4 @@
 #
 # SPDX-License-Identifier: Apache-2.0
 
-__version__ = "0.2.0"
+__version__ = "0.3.0"