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