Merge branch 'main' into subtests2

Author: mdboom

.github/workflows/ci-pixi-source-test.yml

@@ -52,7 +52,11 @@ defaults:
     shell: bash --noprofile --norc -xeuo pipefail {0}
 
 env:
-  PIXI_VERSION: "v0.66.0" # keep in sync with the version developers run locally
+  # keep in sync with the version developers run locally. Must be >=0.71.0:
+  # older pixi re-ran the editable source build on every `pixi run`, recompiling
+  # all Cython extensions (#2138). The fix (content-addressed source-build cache,
+  # prefix-dev/pixi#6285 + #6123) also bumps the pixi.lock format to v7.
+  PIXI_VERSION: "v0.73.0"
 
 jobs:
   # ── PR guard: CPU-only build + import + placement smoke ──

README.md

@@ -5,7 +5,6 @@ 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 functionality
 * [cuda.bindings](https://nvidia.github.io/cuda-python/cuda-bindings/latest): Low-level Python bindings to CUDA C APIs
 * [cuda.pathfinder](https://nvidia.github.io/cuda-python/cuda-pathfinder/latest): Utilities for locating CUDA components installed in the user's Python environment
-* [cuda.coop](https://nvidia.github.io/cccl/unstable/python/coop.html): A Python module providing CCCL's reusable block-wide and warp-wide *device* primitives for use within Numba CUDA kernels
 * [cuda.compute](https://nvidia.github.io/cccl/unstable/python/compute/index.html): 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-mlir](https://nvidia.github.io/numba-cuda-mlir/): An evolution of Numba CUDA that improves upon its technical foundation and performance to provide the future of CUDA Python JIT compilation. It currently supports developing CUDA **SIMT** kernels in Python, providing Python bindings for accelerated device libraries, and serving as a compiler for user-defined functions in accelerated libraries.
 * [numba.cuda](https://nvidia.github.io/numba-cuda/): A Python DSL that exposes CUDA **SIMT** programming model and compiles a restricted subset of Python code into CUDA kernels and device functions

benchmarks/cuda_bindings/pixi.lock

@@ -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

benchmarks/cuda_core/pixi.lock

@@ -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]
 

cuda_bindings/AGENTS.md

@@ -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,25 +70,19 @@ 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
             print(err, file=sys.stderr)  # noqa: T201
             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)))
 

cuda_bindings/build_hooks.py

@@ -7,6 +7,28 @@
 #include <functional>
 #include <stdexcept>
 #include <string>
+#include <climits>
+#include <cstdint>
+
+// PyLong_AsInt entered the public/stable CPython API in 3.13. cuda.bindings
+// supports Python 3.10+, so provide a file-local backport for older builds.
+// This is a copy of the CPython implementation; it is `static` (unlike the
+// original) because this header is compiled into every extension module that
+// includes it, mirroring the other helpers below.
+#if PY_VERSION_HEX < 0x030D0000
+static int
+PyLong_AsInt(PyObject *obj)
+{
+    int overflow;
+    long result = PyLong_AsLongAndOverflow(obj, &overflow);
+    if (overflow || result > INT_MAX || result < INT_MIN) {
+        PyErr_SetString(PyExc_OverflowError,
+                        "Python int too large to convert to C int");
+        return -1;
+    }
+    return (int)result;
+}
+#endif
 
 static PyObject* ctypes_module = nullptr;
 
@@ -69,7 +91,13 @@ static void populate_feeders(PyTypeObject* target_t, PyTypeObject* source_t)
         {
             m_feeders[{target_t,source_t}] = [](void* ptr, PyObject* value) -> int
             {
-                *((int*)ptr) = (int)PyLong_AsLong(value);
+                // PyLong_AsInt range-checks against the 32-bit int slot and raises
+                // OverflowError itself, so an out-of-range value is rejected rather
+                // than silently truncated.
+                int v = PyLong_AsInt(value);
+                if (v == -1 && PyErr_Occurred())
+                    return -1;
+                *((int*)ptr) = v;
                 return sizeof(int);
             };
             return;
@@ -89,7 +117,23 @@ static void populate_feeders(PyTypeObject* target_t, PyTypeObject* source_t)
         {
             m_feeders[{target_t,source_t}] = [](void* ptr, PyObject* value) -> int
             {
-                *((int8_t*)ptr) = (int8_t)PyLong_AsLong(value);
+                // c_byte is an 8-bit slot with no dedicated CPython converter, so
+                // range-check explicitly against INT8_MIN/INT8_MAX. AsLongAndOverflow's
+                // `overflow` only flags values outside `long` (64-bit on LP64), so a
+                // value in that range would be silently truncated by (int8_t)v without
+                // the explicit bounds check. When overflow!=0, v is the -1 sentinel
+                // (not the real value), so that case must be caught before trusting v.
+                int overflow = 0;
+                long v = PyLong_AsLongAndOverflow(value, &overflow);
+                if (overflow == 0 && v == -1 && PyErr_Occurred())
+                    return -1;  // non-overflow conversion error; exception already set
+                if (overflow != 0 || v < INT8_MIN || v > INT8_MAX)
+                {
+                    PyErr_SetString(PyExc_OverflowError,
+                        "Python int is out of range for a c_byte (8-bit) kernel argument");
+                    return -1;
+                }
+                *((int8_t*)ptr) = (int8_t)v;
                 return sizeof(int8_t);
             };
             return;

cuda_bindings/cuda/bindings/_example_helpers/common.py

@@ -4,4 +4,4 @@
 # Include "param_packer.h" so its contents get compiled into every
 # Cython extension module that depends on param_packer.pxd.
 cdef extern from "param_packer.h":
-    int feed(void* ptr, object o, object ct)
+    int feed(void* ptr, object o, object ct) except? -1

cuda_bindings/cuda/bindings/_lib/param_packer.h

@@ -162,6 +162,7 @@ cdef class _HelperCUcoredumpSettings:
     cdef cydriver.CUcoredumpSettings_enum _attrib
     cdef bint _is_getter
     cdef size_t _size
+    cdef object _references  # keeps caller bytes alive so _charstar stays valid
 
     # Return values
     cdef bint _bool