Merge branch 'main' into reduce_cudart

Author: leofang

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/tests/conftest.py

@@ -65,11 +65,5 @@ def pop_all_contexts():
     return pop_all_contexts
 
 
-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.",
-)
-
-
 # TODO: make the fixture more sophisticated using path finder
 skipif_need_cuda_headers = pytest.mark.skipif(os.environ.get("CUDA_PATH") is None, reason="need CUDA header")

cuda_core/tests/test_cuda_utils.py

@@ -3,7 +3,6 @@
 # SPDX-License-Identifier: LicenseRef-NVIDIA-SOFTWARE-LICENSE
 
 import pytest
-from conftest import skipif_testing_with_compute_sanitizer
 
 from cuda.bindings import driver, runtime
 from cuda.core.experimental._utils import cuda_utils
@@ -41,8 +40,6 @@ def test_runtime_cuda_error_explanations_health():
         assert not extra_expl
 
 
-# this test causes an API error when the driver is too old to know about all of the error codes
-@skipif_testing_with_compute_sanitizer
 def test_check_driver_error():
     num_unexpected = 0
     for error in driver.CUresult:

cuda_core/tests/test_event.py

@@ -7,7 +7,7 @@
 
 import numpy as np
 import pytest
-from conftest import skipif_need_cuda_headers, skipif_testing_with_compute_sanitizer
+from conftest import skipif_need_cuda_headers
 
 import cuda.core.experimental
 from cuda.core.experimental import Device, EventOptions, LaunchConfig, Program, ProgramOptions, launch
@@ -71,7 +71,6 @@ def test_is_done(init_cuda):
     assert event.is_done in (True, False)
 
 
-@skipif_testing_with_compute_sanitizer
 def test_error_timing_disabled():
     device = Device()
     device.set_current()
@@ -94,7 +93,6 @@ def test_error_timing_disabled():
         event2 - event1
 
 
-@skipif_testing_with_compute_sanitizer
 def test_error_timing_recorded():
     device = Device()
     device.set_current()
@@ -114,7 +112,6 @@ def test_error_timing_recorded():
         event3 - event2
 
 
-@skipif_testing_with_compute_sanitizer
 @skipif_need_cuda_headers  # libcu++
 @pytest.mark.skipif(tuple(int(i) for i in np.__version__.split(".")[:2]) < (2, 1), reason="need numpy 2.1.0+")
 def test_error_timing_incomplete():

cuda_core/tests/test_linker.py

@@ -3,7 +3,6 @@
 # SPDX-License-Identifier: LicenseRef-NVIDIA-SOFTWARE-LICENSE
 
 import pytest
-from conftest import skipif_testing_with_compute_sanitizer
 
 from cuda.core.experimental import Device, Linker, LinkerOptions, Program, ProgramOptions, _linker
 from cuda.core.experimental._module import ObjectCode
@@ -145,8 +144,6 @@ def test_linker_link_invalid_target_type(compile_ptx_functions):
         linker.link("invalid_target")
 
 
-# this test causes an API error when using the culink API
-@skipif_testing_with_compute_sanitizer
 def test_linker_get_error_log(compile_ptx_functions):
     options = LinkerOptions(name="ABC", arch=ARCH)