@@ -11,14 +11,7 @@
diff --git a/cuda_bindings_12/docs/source/contribute.md b/cuda_bindings_12/docs/source/contribute.md
index 4620da882c7..e2f95fab6f2 100644
--- a/cuda_bindings_12/docs/source/contribute.md
+++ b/cuda_bindings_12/docs/source/contribute.md
@@ -12,6 +12,6 @@ Thank you for your interest in contributing to `cuda-bindings`! Based on the typ
- Before starting work on an existing issue, comment on the issue to
express your interest and wait to be assigned by a maintainer. This
helps avoid redundant effort.
- - Follow the repository [contribution guide](https://github.com/NVIDIA/cuda-python/blob/12.9.x/CONTRIBUTING.md),
+ - Follow the repository [contribution guide](https://github.com/NVIDIA/cuda-python/blob/main/CONTRIBUTING.md),
including signing off each commit under the Developer Certificate of
Origin (DCO) and cryptographically signing commits.
diff --git a/cuda_bindings_12/docs/source/install.md b/cuda_bindings_12/docs/source/install.md
index 7fac1517623..5284536dd1b 100644
--- a/cuda_bindings_12/docs/source/install.md
+++ b/cuda_bindings_12/docs/source/install.md
@@ -5,7 +5,7 @@
`cuda.bindings` supports the same platforms as CUDA. Runtime dependencies are:
* Linux (x86-64, arm64) and Windows (x86-64)
-* Python 3.9 - 3.13
+* Python 3.10 - 3.15
* Driver: Linux (450.80.02 or later) Windows (456.38 or later)
* Optionally, NVRTC, nvJitLink, and NVVM from CUDA Toolkit 12.x
@@ -19,12 +19,12 @@ Starting from v12.8.0, `cuda-python` becomes a meta package which currently depe
## Installing from PyPI
```console
-$ pip install -U cuda-python
+$ pip install -U "cuda-python~=12.9.0"
```
Install all optional dependencies with:
```{code-block} shell
-pip install -U cuda-python[all]
+pip install -U "cuda-python[all]~=12.9.0"
```
Where the optional dependencies are:
@@ -37,7 +37,7 @@ Where the optional dependencies are:
## Installing from Conda
```console
-$ conda install -c conda-forge cuda-python
+$ conda install -c conda-forge cuda-python=12.9
```
diff --git a/cuda_bindings_12/docs/source/motivation.md b/cuda_bindings_12/docs/source/motivation.md
index 5b8879f2bc5..de71bee4d08 100644
--- a/cuda_bindings_12/docs/source/motivation.md
+++ b/cuda_bindings_12/docs/source/motivation.md
@@ -4,7 +4,7 @@
NVIDIA’s CUDA Python provides [Cython](https://cython.org/) bindings and Python
wrappers for the driver and runtime API for existing toolkits and libraries to
simplify GPU-based accelerated processing. Python is one of the most popular
-programming languages for science, engineering, data analytics, and deep
+programming languages for science, engineering, data analytics, and deep
learning applications. The goal of CUDA Python is to unify
the Python ecosystem with a single set of interfaces that provide full coverage
of and access to the CUDA host APIs from Python.
@@ -19,18 +19,18 @@ science, and AI.
[Anaconda](https://www.anaconda.com/) that can compile Python code for execution
on CUDA-capable GPUs, provides Python developers with an easy entry into
GPU-accelerated computing and a path for using increasingly sophisticated CUDA
-code with a minimum of new syntax and jargon. Numba has its own CUDA driver API
-bindings that can now be replaced with CUDA Python. With CUDA Python and Numba,
+code with a minimum of new syntax and jargon. Numba has its own CUDA driver API
+bindings that can now be replaced with CUDA Python. With CUDA Python and Numba,
you get the best of both worlds: rapid iterative development with Python and the
speed of a compiled language targeting both CPUs and NVIDIA GPUs.
[CuPy](https://cupy.dev/) is a
[NumPy](https://numpy.org/)/[SciPy](https://www.scipy.org/) compatible Array
library, from [Preferred Networks](https://www.preferred.jp/en/), for
-GPU-accelerated computing with Python. CUDA Python simplifies the CuPy build
-and allows for a faster and smaller memory footprint when importing the CuPy
-Python module. In the future, when more CUDA Toolkit libraries are supported,
-CuPy will have a lighter maintenance overhead and have fewer wheels to
+GPU-accelerated computing with Python. CUDA Python simplifies the CuPy build
+and allows for a faster and smaller memory footprint when importing the CuPy
+Python module. In the future, when more CUDA Toolkit libraries are supported,
+CuPy will have a lighter maintenance overhead and have fewer wheels to
release. Users benefit from a faster CUDA runtime!
Our goal is to help unify the Python CUDA ecosystem with a single standard set
diff --git a/cuda_bindings_12/docs/source/overview.md b/cuda_bindings_12/docs/source/overview.md
index 5eb083d20d4..50278ce8afb 100644
--- a/cuda_bindings_12/docs/source/overview.md
+++ b/cuda_bindings_12/docs/source/overview.md
@@ -33,7 +33,7 @@ API](https://docs.nvidia.com/cuda/cuda-driver-api/index.html) and
[NVRTC](https://docs.nvidia.com/cuda/nvrtc/index.html) modules from the `cuda.bindings`
package. Next, we consider how to store host data and pass it to the device. Different
approaches can be used to accomplish this and are described in [Preparing kernel
-arguments](https://nvidia.github.io/cuda-python/cuda-bindings/latest/overview.html#preparing-kernel-arguments).
+arguments](#preparing-kernel-arguments).
In this example, we will use NumPy to store host data and pass it to the device, so let's
import this dependency as well.
@@ -53,7 +53,8 @@ def _cudaGetErrorEnum(error):
elif isinstance(error, nvrtc.nvrtcResult):
return nvrtc.nvrtcGetErrorString(error)[1]
else:
- raise RuntimeError('Unknown error type: {}'.format(error))
+ raise RuntimeError("Unknown error type: {}".format(error))
+
def checkCudaErrors(result):
if result[0].value:
@@ -103,9 +104,13 @@ checkCudaErrors(driver.cuInit(0))
cuDevice = checkCudaErrors(driver.cuDeviceGet(0))
# Derive target architecture for device 0
-major = checkCudaErrors(driver.cuDeviceGetAttribute(driver.CUdevice_attribute.CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, cuDevice))
-minor = checkCudaErrors(driver.cuDeviceGetAttribute(driver.CUdevice_attribute.CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, cuDevice))
-arch_arg = bytes(f'--gpu-architecture=compute_{major}{minor}', 'ascii')
+major = checkCudaErrors(
+ driver.cuDeviceGetAttribute(driver.CUdevice_attribute.CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, cuDevice)
+)
+minor = checkCudaErrors(
+ driver.cuDeviceGetAttribute(driver.CUdevice_attribute.CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, cuDevice)
+)
+arch_arg = bytes(f"--gpu-architecture=compute_{major}{minor}", "ascii")
# Create program
prog = checkCudaErrors(nvrtc.nvrtcCreateProgram(str.encode(saxpy), b"saxpy.cu", 0, [], []))
@@ -178,12 +183,8 @@ dOutclass = checkCudaErrors(driver.cuMemAlloc(bufferSize))
stream = checkCudaErrors(driver.cuStreamCreate(0))
-checkCudaErrors(driver.cuMemcpyHtoDAsync(
- dXclass, hX.ctypes.data, bufferSize, stream
-))
-checkCudaErrors(driver.cuMemcpyHtoDAsync(
- dYclass, hY.ctypes.data, bufferSize, stream
-))
+checkCudaErrors(driver.cuMemcpyHtoDAsync(dXclass, hX.ctypes.data, bufferSize, stream))
+checkCudaErrors(driver.cuMemcpyHtoDAsync(dYclass, hY.ctypes.data, bufferSize, stream))
```
With data prep and resources allocation finished, the kernel is ready to be
@@ -213,23 +214,23 @@ args = np.array([arg.ctypes.data for arg in args], dtype=np.uint64)
Now the kernel can be launched:
```python
-checkCudaErrors(driver.cuLaunchKernel(
- kernel,
- NUM_BLOCKS, # grid x dim
- 1, # grid y dim
- 1, # grid z dim
- NUM_THREADS, # block x dim
- 1, # block y dim
- 1, # block z dim
- 0, # dynamic shared memory
- stream, # stream
- args.ctypes.data, # kernel arguments
- 0, # extra (ignore)
-))
-
-checkCudaErrors(driver.cuMemcpyDtoHAsync(
- hOut.ctypes.data, dOutclass, bufferSize, stream
-))
+checkCudaErrors(
+ driver.cuLaunchKernel(
+ kernel,
+ NUM_BLOCKS, # grid x dim
+ 1, # grid y dim
+ 1, # grid z dim
+ NUM_THREADS, # block x dim
+ 1, # block y dim
+ 1, # block z dim
+ 0, # dynamic shared memory
+ stream, # stream
+ args.ctypes.data, # kernel arguments
+ 0, # extra (ignore)
+ )
+)
+
+checkCudaErrors(driver.cuMemcpyDtoHAsync(hOut.ctypes.data, dOutclass, bufferSize, stream))
checkCudaErrors(driver.cuStreamSynchronize(stream))
```
@@ -245,7 +246,7 @@ in the designated stream are finished.
# Assert values are same after running kernel
hZ = a * hX + hY
if not np.allclose(hOut, hZ):
- raise ValueError("Error outside tolerance for host-device vectors")
+ raise ValueError("Error outside tolerance for host-device vectors")
```
Perform verification of the data to ensure correctness and finish the code with
@@ -281,7 +282,7 @@ nsys profile -s none -t cuda --stats=true
* -
- C++
- - Python
+ - Python
* - Kernel execution
- 352µs
- 352µs
@@ -302,6 +303,7 @@ maximize performance ({numref}`Figure 1`).
Screenshot of Nsight Compute CLI output of `cuda.bindings` example.
```
+(preparing-kernel-arguments)=
## Preparing kernel arguments
The `cuLaunchKernel` API bindings retain low-level CUDA argument preparation requirements:
@@ -430,14 +432,21 @@ kernelParams = np.array([arg.ctypes.data for arg in kernelValues], dtype=np.intp
The launch API supports [Buffer Protocol](https://docs.python.org/3/c-api/buffer.html) objects, therefore we can pass the array object directly.
```python
-checkCudaErrors(cuda.cuLaunchKernel(
- kernel,
- 1, 1, 1, # grid dim
- 1, 1, 1, # block dim
- 0, stream, # shared mem and stream
- kernelParams=kernelParams,
- extra=0,
-))
+checkCudaErrors(
+ cuda.cuLaunchKernel(
+ kernel,
+ 1,
+ 1,
+ 1, # grid dim
+ 1,
+ 1,
+ 1, # block dim
+ 0,
+ stream, # shared mem and stream
+ kernelParams=kernelParams,
+ extra=0,
+ )
+)
```
### Using ctypes
@@ -461,6 +470,7 @@ For this example the result becomes:
class testStruct(ctypes.Structure):
_fields_ = [("value", ctypes.c_int)]
+
# Allocate device memory
pInt = checkCudaErrors(cudart.cudaMalloc(ctypes.sizeof(ctypes.c_int)))
pFloat = checkCudaErrors(cudart.cudaMalloc(ctypes.sizeof(ctypes.c_float)))
@@ -496,21 +506,28 @@ In all three cases, the API call will fetch the underlying pointer value and con
With the setup complete, the kernel can be launched:
```python
-checkCudaErrors(cuda.cuLaunchKernel(
- kernel,
- 1, 1, 1, # grid dim
- 1, 1, 1, # block dim
- 0, stream, # shared mem and stream
- kernelParams=(kernelValues, kernelTypes),
- extra=0,
-))
+checkCudaErrors(
+ cuda.cuLaunchKernel(
+ kernel,
+ 1,
+ 1,
+ 1, # grid dim
+ 1,
+ 1,
+ 1, # block dim
+ 0,
+ stream, # shared mem and stream
+ kernelParams=(kernelValues, kernelTypes),
+ extra=0,
+ )
+)
```
### CUDA objects
-Certain CUDA kernels use native CUDA types as their parameters such as `cudaTextureObject_t`. These types require special handling since they're neither a primitive ctype nor a custom user type. Since `cuda.bindings` exposes each of them as Python classes, they each implement `getPtr()` and `__int__()`. These two callables used to support the NumPy and ctypes approach. The difference between each call is further described under [Tips and Tricks](https://nvidia.github.io/cuda-python/cuda-bindings/latest/tips_and_tricks.html#).
+Certain CUDA kernels use native CUDA types as their parameters such as `cudaTextureObject_t`. These types require special handling since they're neither a primitive ctype nor a custom user type. Since `cuda.bindings` exposes each of them as Python classes, they each implement `getPtr()` and `__int__()`. These two callables used to support the NumPy and ctypes approach. The difference between each call is further described under [Tips and Tricks](tips_and_tricks.rst).
-For this example, lets use the `transformKernel` from [examples/0_Introduction/simpleCubemapTexture_test.py](https://github.com/NVIDIA/cuda-python/blob/main/cuda_bindings/examples/0_Introduction/simpleCubemapTexture_test.py):
+For this example, lets use the `transformKernel` from [examples/0_Introduction/simpleCubemapTexture_test.py](https://github.com/NVIDIA/cuda-python/blob/main/cuda_bindings_12/examples/0_Introduction/simpleCubemapTexture_test.py):
```python
simpleCubemapTexture = """\
@@ -521,6 +538,7 @@ __global__ void transformKernel(float *g_odata, int width, cudaTextureObject_t t
}
"""
+
def main():
...
d_data = checkCudaErrors(cudart.cudaMalloc(size))
diff --git a/cuda_bindings_12/docs/source/release.rst b/cuda_bindings_12/docs/source/release.rst
index 8a19737abef..f844bfc251d 100644
--- a/cuda_bindings_12/docs/source/release.rst
+++ b/cuda_bindings_12/docs/source/release.rst
@@ -6,31 +6,6 @@ Release Notes
.. toctree::
:maxdepth: 3
+ :glob:
- 12.X.Y
- 12.9.0
- 12.8.0
- 12.6.2
- 12.6.1
- 12.6.0
- 12.5.0
- 12.4.0
- 12.3.0
- 12.2.1
- 12.2.0
- 12.1.0
- 12.0.0
- 11.8.7
- 11.8.6
- 11.8.5
- 11.8.4
- 11.8.3
- 11.8.2
- 11.8.1
- 11.8.0
- 11.7.1
- 11.7.0
- 11.6.1
- 11.6.0
- 11.5.0
- 11.4.0
+ release/*[0-9]-notes
diff --git a/cuda_bindings_12/docs/source/release/11.7.1-notes.md b/cuda_bindings_12/docs/source/release/11.7.1-notes.md
index 2997c9da56c..981e00aa751 100644
--- a/cuda_bindings_12/docs/source/release/11.7.1-notes.md
+++ b/cuda_bindings_12/docs/source/release/11.7.1-notes.md
@@ -8,6 +8,7 @@ Released on June 29, 2022
## Limitations
+(cuda-bindings-11-7-1-source-builds)=
### Source builds
CUDA Python no longer re-declares CUDA types, instead it uses the types from CUDA C headers. As such source builds now need to access to latest CTK headers. In particular:
diff --git a/cuda_bindings_12/docs/source/release/11.8.0-notes.md b/cuda_bindings_12/docs/source/release/11.8.0-notes.md
index c5bf9f71c32..02996130d0d 100644
--- a/cuda_bindings_12/docs/source/release/11.8.0-notes.md
+++ b/cuda_bindings_12/docs/source/release/11.8.0-notes.md
@@ -10,7 +10,7 @@ Released on October 3, 2022
### Source Builds
-CUDA Python source builds now parse CUDA headers located in $CUDA_HOME directory, enabling/disabling types and APIs if defined. Therefore this removes the need for CTK headers to have all types defined. By allowing minor variations, previous [11.7.1 mobile platform workaround](https://nvidia.github.io/cuda-python/release/11.7.1-notes.html#source-builds) is no longer needed.
+CUDA Python source builds now parse CUDA headers located in $CUDA_HOME directory, enabling/disabling types and APIs if defined. Therefore this removes the need for CTK headers to have all types defined. By allowing minor variations, previous [11.7.1 mobile platform workaround](#cuda-bindings-11-7-1-source-builds) is no longer needed.
It's still required that source builds use the latest CTK headers (i.e. “$CUDA_HOME/include” has latest CTK headers).
diff --git a/cuda_bindings_12/docs/source/release/12.9.1-notes.rst b/cuda_bindings_12/docs/source/release/12.9.1-notes.rst
new file mode 100644
index 00000000000..bb536b8a135
--- /dev/null
+++ b/cuda_bindings_12/docs/source/release/12.9.1-notes.rst
@@ -0,0 +1,51 @@
+.. SPDX-FileCopyrightText: Copyright (c) 2025-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+.. SPDX-License-Identifier: Apache-2.0
+
+.. module:: cuda.bindings
+
+``cuda-bindings`` 12.9.1 Release notes
+======================================
+
+Released on Aug 6, 2025
+
+
+Highlights
+----------
+
+* A utility module :mod:`cuda.bindings.utils` is added
+
+ * Using ``int(cuda_obj)`` to retrieve the underlying address of a CUDA object is deprecated and
+ subject to future removal. Please switch to use :func:`~cuda.bindings.utils.get_cuda_native_handle`
+ instead.
+
+* The ``cuda.bindings.cufile`` Python module was added, wrapping the
+ `cuFile C APIs `_.
+ Supported on Linux only.
+
+ * Currently using this module requires NumPy to be present. Any recent NumPy 1.x or 2.x should work.
+
+* Python bindings in every module, including ``driver``, ``runtime``, and ``nvrtc``, now have the GIL
+ released before calling the underlying C APIs.
+
+
+Bug fixes
+---------
+
+* Fix a library loading bug that preferred shared libraries without a SOVERSION.
+
+
+Miscellaneous
+-------------
+
+* All Python bindings now have the GIL released when calling into the underlying C APIs.
+* Added PTX utilities including :func:`~utils.get_minimal_required_cuda_ver_from_ptx_ver` and :func:`~utils.get_ptx_ver`.
+* Common CUDA objects such as :class:`~runtime.cudaStream_t` now compare equal if the underlying address is the same.
+* Add a binding to ``nvvmGetErrorString()``.
+* Build the bindings with Cython profile hooks disabled.
+* The internal pathfinder module is now isolated to a standalone package ``cuda-pathfinder`` and made as a required dependency.
+
+
+Known issues
+------------
+
+* Updating from older versions (v12.6.2.post1 and below) via ``pip install -U cuda-python`` might not work. Please do a clean re-installation by uninstalling ``pip uninstall -y cuda-python`` followed by installing ``pip install cuda-python``.
diff --git a/cuda_bindings_12/docs/source/release/12.9.2-notes.rst b/cuda_bindings_12/docs/source/release/12.9.2-notes.rst
new file mode 100644
index 00000000000..11d3bf0af31
--- /dev/null
+++ b/cuda_bindings_12/docs/source/release/12.9.2-notes.rst
@@ -0,0 +1,21 @@
+.. SPDX-FileCopyrightText: Copyright (c) 2025-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+.. SPDX-License-Identifier: Apache-2.0
+
+.. module:: cuda.bindings
+
+``cuda-bindings`` 12.9.2 Release notes
+======================================
+
+Released on Aug 18, 2025
+
+
+Highlights
+----------
+
+* Make populating the internal symbol table thread-safe.
+
+
+Known issues
+------------
+
+* Updating from older versions (v12.6.2.post1 and below) via ``pip install -U cuda-python`` might not work. Please do a clean re-installation by uninstalling ``pip uninstall -y cuda-python`` followed by installing ``pip install cuda-python``.
diff --git a/cuda_bindings_12/docs/source/release/12.9.3-notes.rst b/cuda_bindings_12/docs/source/release/12.9.3-notes.rst
new file mode 100644
index 00000000000..66876a7f3eb
--- /dev/null
+++ b/cuda_bindings_12/docs/source/release/12.9.3-notes.rst
@@ -0,0 +1,29 @@
+.. SPDX-FileCopyrightText: Copyright (c) 2025-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+.. SPDX-License-Identifier: Apache-2.0
+
+.. module:: cuda.bindings
+
+``cuda-bindings`` 12.9.3 Release notes
+======================================
+
+Released on Oct 9, 2025
+
+
+Highlights
+----------
+
+* This is the last release that officially supports Python 3.9.
+* Python 3.14 is supported.
+* **Experimental** free-threaded builds for Python 3.13/3.14 are made available. Any bugs can be reported to `our GitHub repo `_. More details are available in our :ref:`support` docs.
+* Automatic CUDA library path detection based on ``CUDA_HOME``, eliminating the need to manually set ``LIBRARY_PATH`` environment variables for installation.
+* The Python overhead of calling functions in CUDA bindings in ``driver``, ``runtime`` and ``nvrtc`` has been reduced by approximately 30%.
+* On Windows, the ``pywin32`` dependency has been removed. The necessary Windows API functions are now accessed directly.
+* Updated the ``cuda.bindings.runtime`` module to statically link against the CUDA Runtime library from CUDA Toolkit 12.9.1.
+* ``cyruntime.getLocalRuntimeVersion`` now uses pathfinder to find the CUDA runtime.
+
+
+Known issues
+------------
+
+* Updating from older versions (v12.6.2.post1 and below) via ``pip install -U cuda-python`` might not work. Please do a clean re-installation by uninstalling ``pip uninstall -y cuda-python`` followed by installing ``pip install cuda-python``.
+* The graphics APIs in ``cuda.bindings.runtime`` are inadvertently disabled in 12.9.3. Users needing these APIs should update to 12.9.4.
diff --git a/cuda_bindings_12/docs/source/release/12.9.4-notes.rst b/cuda_bindings_12/docs/source/release/12.9.4-notes.rst
new file mode 100644
index 00000000000..79abcd35dba
--- /dev/null
+++ b/cuda_bindings_12/docs/source/release/12.9.4-notes.rst
@@ -0,0 +1,23 @@
+.. SPDX-FileCopyrightText: Copyright (c) 2025-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+.. SPDX-License-Identifier: Apache-2.0
+
+.. module:: cuda.bindings
+
+``cuda-bindings`` 12.9.4 Release notes
+======================================
+
+
+
+Highlights
+----------
+
+
+Bug fixes
+---------
+
+* The graphics APIs in ``cuda.bindings.runtime`` were inadvertently disabled in 12.9.3. This has been fixed.
+
+Known issues
+------------
+
+* Updating from older versions (v12.6.2.post1 and below) via ``pip install -U cuda-python`` might not work. Please do a clean re-installation by uninstalling ``pip uninstall -y cuda-python`` followed by installing ``pip install cuda-python``.
diff --git a/cuda_bindings_12/docs/source/release/12.9.5-notes.rst b/cuda_bindings_12/docs/source/release/12.9.5-notes.rst
new file mode 100644
index 00000000000..13c49394a29
--- /dev/null
+++ b/cuda_bindings_12/docs/source/release/12.9.5-notes.rst
@@ -0,0 +1,29 @@
+.. SPDX-FileCopyrightText: Copyright (c) 2025-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+.. SPDX-License-Identifier: Apache-2.0
+
+.. module:: cuda.bindings
+
+``cuda-bindings`` 12.9.5 Release notes
+======================================
+
+Highlights
+----------
+
+* Added ``__cuda_stream__`` protocol support to ``driver.CUStream`` class, enabling better interoperability with libraries that expect this protocol.
+* Python 3.9 support was dropped (end of life).
+
+Bug fixes
+---------
+
+* Fixed ``cuStreamBeginCaptureToGraph`` to allow the ``dependencyData`` argument to be optional, matching the underlying CUDA API behavior.
+
+Experimental
+------------
+
+* Experimental NVML bindings were added under ``cuda.bindings._nvml``.
+
+Known issues
+------------
+
+* Updating from older versions (v12.6.2.post1 and below) via ``pip install -U cuda-python`` might not work. Please do a clean re-installation by uninstalling ``pip uninstall -y cuda-python`` followed by installing ``pip install cuda-python``.
+* The graphics APIs in ``cuda.bindings.runtime`` are inadvertently disabled in 12.9.3. Users needing these APIs should update to 12.9.4 or higher.
diff --git a/cuda_bindings_12/docs/source/release/12.9.6-notes.rst b/cuda_bindings_12/docs/source/release/12.9.6-notes.rst
new file mode 100644
index 00000000000..1589796ed9a
--- /dev/null
+++ b/cuda_bindings_12/docs/source/release/12.9.6-notes.rst
@@ -0,0 +1,64 @@
+.. SPDX-FileCopyrightText: Copyright (c) 2025-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+.. SPDX-License-Identifier: Apache-2.0
+
+.. module:: cuda.bindings
+
+``cuda-bindings`` 12.9.6 Release notes
+======================================
+
+Highlights
+----------
+
+* ``cuda.bindings.nvml`` has graduated from experimental (``cuda.bindings._nvml``)
+ to a fully supported public module with extensive handwritten Pythonic API
+ coverage spanning ~170 functions across system queries, device discovery,
+ memory, power, clocks, utilization, thermals, NVLink, and device configuration.
+ (`PR #1524 `_,
+ `PR #1548 `_)
+* Add ``nvFatbin`` bindings.
+ (`PR #1467 `_)
+* Performance improvement: ``cuda.bindings`` now uses a faster ``enum``
+ implementation, rather than the standard library's ``enum.IntEnum``.
+ This leads to much faster import times, and slightly faster attribute access
+ times.
+ (`PR #1581 `_)
+* Multiple performance improvements cumulatively reducing Python-to-C call
+ overhead through faster ``void *`` conversion, faster result returning,
+ optimized enum-to-vector conversion, and stack-allocated small arrays.
+
+Bugfixes
+--------
+
+* Fixed an issue where the ``CU_POINTER_ATTRIBUTE_DEVICE_ORDINAL`` attribute was
+ retrieved as an unsigned int, rather than a signed int.
+ (`PR #1451 `_)
+* Fixed a use-after-free in ``_HelperInputVoidPtr`` properties when backed by
+ Python buffer objects.
+ (`PR #1629 `_)
+
+Miscellaneous
+-------------
+
+* Faster ``void *`` conversion using stack-allocated buffers instead of heap
+ allocation.
+ (`PR #1616 `_)
+* Faster returning of results from driver, runtime, and NVRTC bindings.
+ (`PR #1647 `_,
+ `PR #1656 `_)
+* Faster conversion of enum sequences to vectors by eliminating temporary
+ Python objects.
+ (`PR #1667 `_)
+* Stack-allocated small numeric arrays in driver bindings, reducing heap
+ allocation overhead.
+ (`PR #1545 `_)
+* NVML bindings now use ``cuda_pathfinder`` for library discovery, consistent
+ with other CUDA libraries.
+ (`PR #1661 `_)
+* ``CUDA_HOME`` is no longer required at metadata resolution time (e.g.
+ ``pip install --dry-run``, ``uv lock``); it is only needed at actual build time.
+ (`PR #1652 `_)
+
+Known issues
+------------
+
+* Updating from older versions (v12.6.2.post1 and below) via ``pip install -U cuda-python`` might not work. Please do a clean re-installation by uninstalling ``pip uninstall -y cuda-python`` followed by installing ``pip install cuda-python``.
diff --git a/cuda_bindings_12/docs/source/release/12.9.7-notes.rst b/cuda_bindings_12/docs/source/release/12.9.7-notes.rst
new file mode 100644
index 00000000000..842f1f3923e
--- /dev/null
+++ b/cuda_bindings_12/docs/source/release/12.9.7-notes.rst
@@ -0,0 +1,53 @@
+.. SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+.. SPDX-License-Identifier: Apache-2.0
+
+.. module:: cuda.bindings
+
+``cuda-bindings`` 12.9.7 Release notes
+======================================
+
+Bugfixes
+--------
+
+* Fixed ``CUDA_PYTHON_CUDA_PER_THREAD_DEFAULT_STREAM=0`` incorrectly enabling
+ per-thread default stream mode.
+ (`PR #2110 `_)
+
+* Fixed a use-after-free in ``cudaGraphGetEdges``, ``cudaGraphNodeGetDependencies``,
+ ``cudaGraphNodeGetDependentNodes``, ``cudaStreamGetCaptureInfo``, and their
+ driver-API counterparts (``cuGraphGetEdges``, ``cuGraphNodeGetDependencies``,
+ ``cuGraphNodeGetDependentNodes``, ``cuStreamGetCaptureInfo``). The returned
+ ``cudaGraphEdgeData``/``CUgraphEdgeData`` wrappers were backed by a scratch
+ buffer that was freed before the call returned, leaving every wrapper holding
+ a dangling pointer. The returned wrappers now own deep copies of the edge
+ data.
+ (`Issue #1804 `_,
+ `PR #2110 `_)
+
+* Fixed a double-free in the generated setters for list-valued struct members
+ (e.g. ``CUlaunchConfig.attrs``, ``CUDA_MEM_ALLOC_NODE_PARAMS.accessDescs``,
+ external-semaphore and batch-mem-op node parameter arrays, and their runtime
+ counterparts). Assigning an empty list freed the internal buffer but left
+ the cached pointer non-NULL, so a subsequent assignment or ``__dealloc__``
+ would call ``free()`` again on the dangling pointer.
+ (`PR #2115 `_)
+
+Miscellaneous
+-------------
+
+* NVRTC bindings now use pre-generated Cython files and no longer require
+ pyclibrary header parsing at build time.
+ (`PR #1957 `_)
+
+* Improved generated documentation and argument names, including the ``ind_ex``
+ argument naming bug.
+ (`PR #1928 `_,
+ `PR #2110 `_)
+
+* Source archives now include git archival metadata for setuptools-scm.
+ (`PR #1756 `_)
+
+Known issues
+------------
+
+* Updating from older versions (v12.6.2.post1 and below) via ``pip install -U cuda-python`` might not work. Please do a clean re-installation by uninstalling ``pip uninstall -y cuda-python`` followed by installing ``pip install cuda-python``.
diff --git a/cuda_bindings_12/docs/source/release/12.X.Y-notes.rst b/cuda_bindings_12/docs/source/release/12.X.Y-notes.rst
index f5e5d00a172..a6af0583fa6 100644
--- a/cuda_bindings_12/docs/source/release/12.X.Y-notes.rst
+++ b/cuda_bindings_12/docs/source/release/12.X.Y-notes.rst
@@ -1,6 +1,8 @@
.. SPDX-FileCopyrightText: Copyright (c) 2025-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
.. SPDX-License-Identifier: Apache-2.0
+:orphan:
+
.. module:: cuda.bindings
``cuda-bindings`` 12.X.Y Release notes
diff --git a/cuda_bindings_12/docs/source/support.rst b/cuda_bindings_12/docs/source/support.rst
index 05a60221a05..379bdb9d9ab 100644
--- a/cuda_bindings_12/docs/source/support.rst
+++ b/cuda_bindings_12/docs/source/support.rst
@@ -1,6 +1,8 @@
.. SPDX-FileCopyrightText: Copyright (c) 2025-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
.. SPDX-License-Identifier: Apache-2.0
+.. _support:
+
``cuda.bindings`` Support Policy
================================
@@ -10,14 +12,14 @@ The ``cuda.bindings`` module has the following support policy:
third number in the version string), however, is reserved to reflect Python-only changes and
is out of sync with the Toolkit patch version.
2. The module is actively maintained to support the latest CUDA major version and its prior major
- version. For example, as of writing the bindings for CUDA 11 & 12 are maintained. Any fix in the
- latest bindings would be backported to the prior major version.
+ version. Both active source trees live on the main branch so applicable fixes can update the
+ CUDA 12 and CUDA 13 packages in one change.
3. The module supports `CUDA minor version compatibility`_, meaning that ``cuda.bindings`` 12.x
supports any Toolkit 12.y. (Whether or not a binding API would actually correctly function
depends on the underlying driver and the Toolkit versions, as described in the compatibility
documentation.)
4. The module supports all Python versions following the `CPython EOL schedule`_. As of writing
- Python 3.9 - 3.13 are supported.
+ Python 3.10 - 3.15 are supported.
5. The module exposes a Cython layer from which types and functions could be ``cimport``'d. While
we strive to keep this layer stable, due to Cython limitations a new *minor* release of this
module could require Cython layer users to rebuild their projects and update their pinning to
diff --git a/cuda_bindings_12/docs/source/tips_and_tricks.rst b/cuda_bindings_12/docs/source/tips_and_tricks.rst
index 5de1be1bb64..d54006d78f9 100644
--- a/cuda_bindings_12/docs/source/tips_and_tricks.rst
+++ b/cuda_bindings_12/docs/source/tips_and_tricks.rst
@@ -27,9 +27,9 @@ All of the Python classes do not manage the lifetime of the underlying CUDA C ob
Getting and setting attributes of extension types
=================================================
-While the bindings outwardly present the attributes of extension types in a pythonic way, they can't always be interacted with in a Pythonic style. Often the getters/setters (__getitem__(), __setitem__()) are actually a translation step to convert values between Python and C. For example, in some cases, attempting to modify an attribute in place, will lead to unexpected behavior due to the design of the underlying implementation. For this reason, users should use the getters and setters directly when interacting with extension types.
+While the bindings outwardly present the attributes of extension types in a pythonic way, they can't always be interacted with in a Pythonic style. Often the getters/setters (__getitem__(), __setitem__()) are actually a translation step to convert values between Python and C. For example, in some cases, attempting to modify an attribute in place, will lead to unexpected behavior due to the design of the underlying implementation. For this reason, users should use the getters and setters directly when interacting with extension types.
-An example of this is the :class:`~cuda.bindings.driver.CULaunchConfig` type.
+An example of this is the :class:`~cuda.bindings.driver.CULaunchConfig` type.
.. code-block:: python
@@ -37,7 +37,7 @@ An example of this is the :class:`~cuda.bindings.driver.CULaunchConfig` type.
cfg.numAttrs += 1
attr = cuda.CUlaunchAttribute()
-
+
...
# This works. We are passing the new attribute to the setter
diff --git a/cuda_bindings_12/docs/versions.json b/cuda_bindings_12/docs/versions.json
deleted file mode 100644
index cc1299896a2..00000000000
--- a/cuda_bindings_12/docs/versions.json
+++ /dev/null
@@ -1,7 +0,0 @@
-{
- "latest" : "latest",
- "12.9.0" : "12.9.0",
- "12.8.0" : "12.8.0",
- "12.6.2" : "12.6.2",
- "12.6.1" : "12.6.1"
-}
diff --git a/cuda_bindings_12/examples/0_Introduction/clock_nvrtc_test.py b/cuda_bindings_12/examples/0_Introduction/clock_nvrtc_test.py
index dcf94438fdd..ab9f4f2d996 100644
--- a/cuda_bindings_12/examples/0_Introduction/clock_nvrtc_test.py
+++ b/cuda_bindings_12/examples/0_Introduction/clock_nvrtc_test.py
@@ -1,4 +1,4 @@
-# SPDX-FileCopyrightText: Copyright (c) 2021-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: Apache-2.0
import platform
diff --git a/cuda_bindings_12/examples/0_Introduction/simpleCubemapTexture_test.py b/cuda_bindings_12/examples/0_Introduction/simpleCubemapTexture_test.py
index 8cc409a4000..4ca0f29f5dd 100644
--- a/cuda_bindings_12/examples/0_Introduction/simpleCubemapTexture_test.py
+++ b/cuda_bindings_12/examples/0_Introduction/simpleCubemapTexture_test.py
@@ -1,4 +1,4 @@
-# SPDX-FileCopyrightText: Copyright (c) 2021-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: Apache-2.0
import ctypes
diff --git a/cuda_bindings_12/examples/0_Introduction/simpleP2P_test.py b/cuda_bindings_12/examples/0_Introduction/simpleP2P_test.py
index 3d5dc871481..9f6d053eba9 100644
--- a/cuda_bindings_12/examples/0_Introduction/simpleP2P_test.py
+++ b/cuda_bindings_12/examples/0_Introduction/simpleP2P_test.py
@@ -1,4 +1,4 @@
-# SPDX-FileCopyrightText: Copyright (c) 2021-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: Apache-2.0
import ctypes
diff --git a/cuda_bindings_12/examples/0_Introduction/simpleZeroCopy_test.py b/cuda_bindings_12/examples/0_Introduction/simpleZeroCopy_test.py
index 49a0edf4558..356a5ac56b1 100644
--- a/cuda_bindings_12/examples/0_Introduction/simpleZeroCopy_test.py
+++ b/cuda_bindings_12/examples/0_Introduction/simpleZeroCopy_test.py
@@ -1,4 +1,4 @@
-# SPDX-FileCopyrightText: Copyright (c) 2021-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: Apache-2.0
import ctypes
diff --git a/cuda_bindings_12/examples/0_Introduction/systemWideAtomics_test.py b/cuda_bindings_12/examples/0_Introduction/systemWideAtomics_test.py
index bc51dcdcc96..170170d5dbe 100644
--- a/cuda_bindings_12/examples/0_Introduction/systemWideAtomics_test.py
+++ b/cuda_bindings_12/examples/0_Introduction/systemWideAtomics_test.py
@@ -1,4 +1,4 @@
-# SPDX-FileCopyrightText: Copyright (c) 2021-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: Apache-2.0
import ctypes
diff --git a/cuda_bindings_12/examples/0_Introduction/vectorAddDrv_test.py b/cuda_bindings_12/examples/0_Introduction/vectorAddDrv_test.py
index dcb6cf014c9..0d0f9634e60 100644
--- a/cuda_bindings_12/examples/0_Introduction/vectorAddDrv_test.py
+++ b/cuda_bindings_12/examples/0_Introduction/vectorAddDrv_test.py
@@ -1,4 +1,4 @@
-# SPDX-FileCopyrightText: Copyright (c) 2021-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: Apache-2.0
import ctypes
diff --git a/cuda_bindings_12/examples/0_Introduction/vectorAddMMAP_test.py b/cuda_bindings_12/examples/0_Introduction/vectorAddMMAP_test.py
index 16eafbe3e74..24bd26ab860 100644
--- a/cuda_bindings_12/examples/0_Introduction/vectorAddMMAP_test.py
+++ b/cuda_bindings_12/examples/0_Introduction/vectorAddMMAP_test.py
@@ -1,4 +1,4 @@
-# SPDX-FileCopyrightText: Copyright (c) 2021-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: Apache-2.0
import ctypes
diff --git a/cuda_bindings_12/examples/2_Concepts_and_Techniques/streamOrderedAllocation_test.py b/cuda_bindings_12/examples/2_Concepts_and_Techniques/streamOrderedAllocation_test.py
index 56c7578db13..d7cd51b9f4d 100644
--- a/cuda_bindings_12/examples/2_Concepts_and_Techniques/streamOrderedAllocation_test.py
+++ b/cuda_bindings_12/examples/2_Concepts_and_Techniques/streamOrderedAllocation_test.py
@@ -1,4 +1,4 @@
-# SPDX-FileCopyrightText: Copyright (c) 2021-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: Apache-2.0
import ctypes
diff --git a/cuda_bindings_12/examples/3_CUDA_Features/globalToShmemAsyncCopy_test.py b/cuda_bindings_12/examples/3_CUDA_Features/globalToShmemAsyncCopy_test.py
index c9772ce9ea6..37a83047124 100644
--- a/cuda_bindings_12/examples/3_CUDA_Features/globalToShmemAsyncCopy_test.py
+++ b/cuda_bindings_12/examples/3_CUDA_Features/globalToShmemAsyncCopy_test.py
@@ -1,4 +1,4 @@
-# SPDX-FileCopyrightText: Copyright (c) 2021-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: Apache-2.0
import ctypes
diff --git a/cuda_bindings_12/examples/3_CUDA_Features/simpleCudaGraphs_test.py b/cuda_bindings_12/examples/3_CUDA_Features/simpleCudaGraphs_test.py
index 2aa6a604310..91aec2010ba 100644
--- a/cuda_bindings_12/examples/3_CUDA_Features/simpleCudaGraphs_test.py
+++ b/cuda_bindings_12/examples/3_CUDA_Features/simpleCudaGraphs_test.py
@@ -1,4 +1,4 @@
-# SPDX-FileCopyrightText: Copyright (c) 2021-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: Apache-2.0
import ctypes
diff --git a/cuda_bindings_12/examples/4_CUDA_Libraries/conjugateGradientMultiBlockCG_test.py b/cuda_bindings_12/examples/4_CUDA_Libraries/conjugateGradientMultiBlockCG_test.py
index 13ea157bcc6..96aae9e5f58 100644
--- a/cuda_bindings_12/examples/4_CUDA_Libraries/conjugateGradientMultiBlockCG_test.py
+++ b/cuda_bindings_12/examples/4_CUDA_Libraries/conjugateGradientMultiBlockCG_test.py
@@ -1,4 +1,4 @@
-# SPDX-FileCopyrightText: Copyright (c) 2021-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: Apache-2.0
import ctypes
diff --git a/cuda_bindings_12/examples/common/common.py b/cuda_bindings_12/examples/common/common.py
index 0b0239329e8..2701ac18842 100644
--- a/cuda_bindings_12/examples/common/common.py
+++ b/cuda_bindings_12/examples/common/common.py
@@ -1,4 +1,4 @@
-# SPDX-FileCopyrightText: Copyright (c) 2021-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: Apache-2.0
import os
diff --git a/cuda_bindings_12/examples/common/helper_cuda.py b/cuda_bindings_12/examples/common/helper_cuda.py
index 9ea5e038119..3268e4e82e2 100644
--- a/cuda_bindings_12/examples/common/helper_cuda.py
+++ b/cuda_bindings_12/examples/common/helper_cuda.py
@@ -1,4 +1,4 @@
-# SPDX-FileCopyrightText: Copyright (c) 2021-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: Apache-2.0
from common.helper_string import checkCmdLineFlag, getCmdLineArgumentInt
diff --git a/cuda_bindings_12/examples/common/helper_string.py b/cuda_bindings_12/examples/common/helper_string.py
index b6cdd8b2a34..5c48ffc964e 100644
--- a/cuda_bindings_12/examples/common/helper_string.py
+++ b/cuda_bindings_12/examples/common/helper_string.py
@@ -1,4 +1,4 @@
-# SPDX-FileCopyrightText: Copyright (c) 2021-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: Apache-2.0
import sys
diff --git a/cuda_bindings_12/examples/extra/isoFDModelling_test.py b/cuda_bindings_12/examples/extra/isoFDModelling_test.py
index 5438ba48972..3c5d4ee6b3a 100644
--- a/cuda_bindings_12/examples/extra/isoFDModelling_test.py
+++ b/cuda_bindings_12/examples/extra/isoFDModelling_test.py
@@ -1,4 +1,4 @@
-# SPDX-FileCopyrightText: Copyright (c) 2021-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: Apache-2.0
import time
diff --git a/cuda_bindings_12/examples/extra/jit_program_test.py b/cuda_bindings_12/examples/extra/jit_program_test.py
index 55602c92593..f5f4a557666 100644
--- a/cuda_bindings_12/examples/extra/jit_program_test.py
+++ b/cuda_bindings_12/examples/extra/jit_program_test.py
@@ -1,4 +1,4 @@
-# SPDX-FileCopyrightText: Copyright (c) 2021-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: Apache-2.0
import ctypes
diff --git a/cuda_bindings_12/examples/extra/numba_emm_plugin.py b/cuda_bindings_12/examples/extra/numba_emm_plugin.py
index 4cdb65e31cf..35b760b2656 100644
--- a/cuda_bindings_12/examples/extra/numba_emm_plugin.py
+++ b/cuda_bindings_12/examples/extra/numba_emm_plugin.py
@@ -1,4 +1,4 @@
-# SPDX-FileCopyrightText: Copyright (c) 2021-2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: Apache-2.0
diff --git a/cuda_bindings_12/examples/pytest.ini b/cuda_bindings_12/examples/pytest.ini
index e105585d5aa..7dcf4f32863 100644
--- a/cuda_bindings_12/examples/pytest.ini
+++ b/cuda_bindings_12/examples/pytest.ini
@@ -1,4 +1,13 @@
+# SPDX-FileCopyrightText: Copyright (c) 2024-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: Apache-2.0
+
[pytest]
+markers =
+ # Keep this authorship marker registry in sync across all pytest config roots.
+ # Search for "agent_authored(model)" before editing.
+ agent_authored(model): agent-authored test not yet materially human-reviewed
+ human_reviewed: agent-authored test materially reviewed or rewritten by a human
+ human_authored: test authored primarily by a human
python_files = *_test.py
python_functions = main
pythonpath = .
diff --git a/cuda_bindings_12/pyproject.toml b/cuda_bindings_12/pyproject.toml
index 493850ecb4e..aee928a8aef 100644
--- a/cuda_bindings_12/pyproject.toml
+++ b/cuda_bindings_12/pyproject.toml
@@ -3,9 +3,9 @@
[build-system]
requires = [
- "setuptools>=80.0.0",
+ "setuptools>=80.0.0",
"setuptools-scm[simple]>=8",
- "cython>=3.2,<3.3",
+ "cython>=3.2,<3.3",
"pyclibrary>=0.1.7"
]
build-backend = "build_hooks"
@@ -93,10 +93,12 @@ select = [
]
ignore = [
- "UP007",
+ "UP007",
"E741", # ambiguous variable name such as I
- "B007", # rename unsued loop variable to _name
- "UP035" # UP006, UP007, UP035 complain about deprecated Typing. use, but disregard backward compatibility of python version
+ "B007", # rename unsued loop variable to _name
+ "UP035", # UP006, UP007, UP035 complain about deprecated Typing. use, but disregard backward compatibility of python version
+ "B905", # preserve the imported line's zip() behavior
+ "I001", # preserve generated and legacy import ordering
]
exclude = ["cuda/bindings/_version.py"]
@@ -104,6 +106,8 @@ exclude = ["cuda/bindings/_version.py"]
[tool.ruff.lint.per-file-ignores]
"setup.py" = ["F401"]
"__init__.py" = ["F401"]
+"site-packages/_cuda_bindings_redirector.py" = ["F401"]
+"cuda/bindings/_internal/_fast_enum.py" = ["E501"]
"examples/**/*" = [
"E722",
@@ -111,13 +115,7 @@ exclude = ["cuda/bindings/_version.py"]
]
"tests/**/*" = [
- "E722",
- "UP022",
- "E402", # module level import not at top of file
- "F841"] # F841 complains about unused variables, but some assignments have side-effects that could be useful for tests (func calls for example)
-
-"benchmarks/**/*" = [
- "E722",
+ "E722",
"UP022",
"E402", # module level import not at top of file
"F841"] # F841 complains about unused variables, but some assignments have side-effects that could be useful for tests (func calls for example)
diff --git a/cuda_bindings_12/setup.py b/cuda_bindings_12/setup.py
index 25c38941b17..3566d931a64 100644
--- a/cuda_bindings_12/setup.py
+++ b/cuda_bindings_12/setup.py
@@ -1,10 +1,9 @@
# SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: Apache-2.0
-import functools
import os
import pathlib
-import subprocess
+import sys
from warnings import warn
from setuptools import setup
@@ -13,6 +12,7 @@
from setuptools.command.editable_wheel import _TopLevelFinder, editable_wheel
import build_hooks
+from _build_shared import drop_empty_venv_include_dir, make_egg_info_command
if os.environ.get("PARALLEL_LEVEL") is not None:
warn(
@@ -25,29 +25,28 @@
nthreads = int(os.environ.get("CUDA_PYTHON_PARALLEL_LEVEL", "0") or "0")
-def _is_clang(compiler):
- @functools.lru_cache
- def _check(compiler_cxx):
- try:
- output = subprocess.check_output([*compiler_cxx, "--version"]) # noqa: S603
- except subprocess.CalledProcessError:
- return False
- lines = output.decode().splitlines()
- return len(lines) > 0 and "clang" in lines[0]
-
- if not hasattr(compiler, "compiler_cxx"):
- return False
- return _check(tuple(compiler.compiler_cxx))
-
-
class build_ext(_build_ext):
+ def finalize_options(self):
+ super().finalize_options()
+ # A per-build venv path on every -I line defeats sccache's
+ # preprocessor cache; see _build_shared.drop_empty_venv_include_dir().
+ self.include_dirs = drop_empty_venv_include_dir(self.include_dirs)
+ # A conda environment can provide libcudart_static.a in /lib.
+ # Added here rather than to the Extensions: Cython copies Extension
+ # library_dirs into the generated sources, and under build isolation
+ # sys.prefix is a fresh /tmp/build-env- per build, which
+ # defeats sccache's preprocessor cache.
+ prefix_lib = os.path.join(sys.prefix, "lib")
+ if prefix_lib not in self.library_dirs:
+ self.library_dirs.append(prefix_lib)
+
def build_extensions(self):
if nthreads > 0:
self.parallel = nthreads
- if _is_clang(self.compiler):
- for ext in self.extensions:
- ext.extra_compile_args = [a for a in ext.extra_compile_args if a != "-fno-var-tracking-assignments"]
+ if build_hooks.force_build_ext:
+ self.force = True
super().build_extensions()
+ build_hooks.record_build_toolchain()
################################################################################
@@ -123,6 +122,7 @@ def _select_strategy(self, name, tag, build_lib):
"build_ext": build_ext,
"build_py": build_py_with_redirector,
"editable_wheel": editable_wheel_with_redirector,
+ "egg_info": make_egg_info_command("v"),
},
zip_safe=False,
)
diff --git a/cuda_bindings_12/site-packages/_cuda_bindings_redirector.py b/cuda_bindings_12/site-packages/_cuda_bindings_redirector.py
index 0feae7a1374..e9267a0f9c9 100644
--- a/cuda_bindings_12/site-packages/_cuda_bindings_redirector.py
+++ b/cuda_bindings_12/site-packages/_cuda_bindings_redirector.py
@@ -10,12 +10,12 @@
class LazyCudaModule(ModuleType):
-
def __getattr__(self, name):
- if name == '__version__':
+ if name == "__version__":
import warnings
+
warnings.warn(
- "accessing cuda.__version__ is deprecated, " "please switch to use cuda.bindings.__version__ instead",
+ "accessing cuda.__version__ is deprecated, please switch to use cuda.bindings.__version__ instead",
FutureWarning,
stacklevel=2,
)
@@ -27,4 +27,4 @@ def __getattr__(self, name):
# Patch in LazyCudaModule for `cuda`
-sys.modules['cuda'].__class__ = LazyCudaModule
+sys.modules["cuda"].__class__ = LazyCudaModule
diff --git a/cuda_bindings_12/tests/conftest.py b/cuda_bindings_12/tests/conftest.py
new file mode 100644
index 00000000000..e9372b8299c
--- /dev/null
+++ b/cuda_bindings_12/tests/conftest.py
@@ -0,0 +1,28 @@
+# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: Apache-2.0
+
+import importlib
+import pathlib
+import sys
+
+# BEGIN SYNCED PYTEST PLUGIN BOOTSTRAP
+# Keep every block with this marker byte-for-byte identical.
+# Before editing or reviewing, find all copies; see root AGENTS.md.
+# Keep the shared pytest plugin available in wheel-test jobs, which run from
+# the monorepo checkout without installing cuda-python-test-helpers.
+try:
+ import cuda_python_test_helpers._pytest_plugin # noqa: F401
+except ImportError as e:
+ # Don't call .resolve(): resolving symlinks can make parents[2] point
+ # somewhere other than the monorepo root if a sub-directory is symlinked.
+ _test_helpers_root = pathlib.Path(__file__).parents[2] / "cuda_python_test_helpers"
+ if not _test_helpers_root.is_dir():
+ raise RuntimeError(f"cuda-python-test-helpers not installed and not found at {_test_helpers_root}") from e
+ for _k in list(sys.modules):
+ if _k == "cuda_python_test_helpers" or _k.startswith("cuda_python_test_helpers."):
+ del sys.modules[_k]
+ sys.path.insert(0, str(_test_helpers_root))
+ importlib.invalidate_caches()
+
+pytest_plugins = ["cuda_python_test_helpers._pytest_plugin"]
+# END SYNCED PYTEST PLUGIN BOOTSTRAP
diff --git a/cuda_bindings_12/tests/cython/build_tests.bat b/cuda_bindings_12/tests/cython/build_tests.bat
index 9edf9653b76..0ef6abb06f3 100644
--- a/cuda_bindings_12/tests/cython/build_tests.bat
+++ b/cuda_bindings_12/tests/cython/build_tests.bat
@@ -4,6 +4,9 @@ REM SPDX-FileCopyrightText: Copyright (c) 2024-2026 NVIDIA CORPORATION & AFFILIA
REM SPDX-License-Identifier: Apache-2.0
setlocal
- set CL=%CL% /I"%CUDA_HOME%\include"
- cythonize -3 -i -Xfreethreading_compatible=True %~dp0test_*.pyx
-endlocal
+set CL=%CL% /I"%CUDA_HOME%\include"
+REM The Python driver provides Cython's .pxd include path and builds in this
+REM directory so Windows does not duplicate the checkout path in link outputs.
+python "%~dp0build_tests.py"
+set "BUILD_RESULT=%ERRORLEVEL%"
+endlocal & exit /b %BUILD_RESULT%
diff --git a/cuda_bindings_12/tests/cython/build_tests.py b/cuda_bindings_12/tests/cython/build_tests.py
new file mode 100644
index 00000000000..a4a26589ce9
--- /dev/null
+++ b/cuda_bindings_12/tests/cython/build_tests.py
@@ -0,0 +1,109 @@
+# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: Apache-2.0
+"""Build cuda_bindings_12 Cython test extensions in-place.
+
+pixi-build's editable install exposes the `cuda` namespace package via a
+PEP 660 finder hook. Python's import machinery honors the hook, but
+Cython's filesystem .pxd resolver only walks real directories on sys.path,
+so `cimport cuda.bindings.*` fails to locate the .pxd files. We resolve
+the namespace package's source root from `cuda.bindings.__file__` and pass
+it via `include_path=` so cythonize finds the .pxd tree on every platform.
+
+When CUDA_PYTHON_CYTHON_CACHE_DIR is set, cythonize uses the same cache
+namespacing and include-path aliasing as the package build (``_build_shared.py``).
+"""
+
+from __future__ import annotations
+
+import importlib.util
+import os
+import sys
+from pathlib import Path
+
+import Cython
+from Cython.Build import cythonize
+from setuptools import setup
+
+import cuda.bindings
+
+_COMPILER_DIRECTIVES = {"freethreading_compatible": True}
+
+
+def _load_build_shared():
+ # A PEP 517 backend file, not an installed module. Load it by path so we do
+ # not put the package directory on sys.path (that would shadow the
+ # installed package). It needs only the standard library.
+ path = Path(__file__).resolve().parents[2] / "_build_shared.py"
+ spec = importlib.util.spec_from_file_location("cuda_bindings_build_shared", path)
+ module = importlib.util.module_from_spec(spec)
+ spec.loader.exec_module(module)
+ return module
+
+
+build_shared = _load_build_shared()
+
+
+def _bindings_source_root() -> Path:
+ # cuda.bindings.__file__ -> ...//cuda/bindings/__init__.py
+ root = Path(cuda.bindings.__file__).resolve().parents[2]
+ if not (root / "cuda" / "bindings").is_dir():
+ raise RuntimeError(
+ f"cuda.bindings source tree not found at {root}; pixi-build editable install layout may have changed."
+ )
+ return root
+
+
+def _cythonize_tests(pyx_files):
+ cache_path = build_shared._cython_cache_path(
+ "cuda-bindings-12-cython-tests",
+ compiler_directives=_COMPILER_DIRECTIVES,
+ language_level=3,
+ cplus=True,
+ )
+ cythonize_kwargs = {
+ "language_level": 3,
+ "nthreads": 1,
+ "compiler_directives": _COMPILER_DIRECTIVES,
+ "cache": cache_path,
+ }
+ if cache_path is None:
+ return cythonize(
+ pyx_files,
+ include_path=[str(_bindings_source_root())],
+ **cythonize_kwargs,
+ )
+
+ # Distinct alias names so a concurrent package build's .cython-stdlib /
+ # .cython-bindings symlinks are not replaced. Relative aliases resolve
+ # next to _build_shared.py (package root).
+ stdlib_target = Path(Cython.__file__).parent / "Includes"
+ with (
+ build_shared._stable_cython_alias(stdlib_target, Path(".cython-stdlib-tests")) as rel_stdlib,
+ build_shared._stable_cython_alias(_bindings_source_root(), Path(".cython-bindings-tests")) as rel_bindings,
+ ):
+ return cythonize(
+ pyx_files,
+ include_path=[".", rel_bindings, rel_stdlib],
+ **cythonize_kwargs,
+ )
+
+
+def main() -> None:
+ script_dir = Path(__file__).resolve().parent
+ # Avoid appending the absolute checkout path under build/temp: the
+ # concatenated path can exceed Windows' path limit. These files are siblings.
+ os.chdir(script_dir)
+ pyx_files = sorted(p.name for p in script_dir.glob("test_*.pyx"))
+ if not pyx_files:
+ raise SystemExit(f"no test_*.pyx files under {script_dir}")
+
+ ext_modules = _cythonize_tests(pyx_files)
+
+ # pytest imports each extension by bare module name (see test_cython.py),
+ # so build in-place next to its .pyx regardless of the invoking cwd.
+ sys.argv = [sys.argv[0], "build_ext", "--inplace"]
+ setup(name="cuda_bindings_cython_tests", ext_modules=ext_modules)
+
+
+if __name__ == "__main__":
+ main()
diff --git a/cuda_bindings_12/tests/cython/build_tests.sh b/cuda_bindings_12/tests/cython/build_tests.sh
index db8b97fbe7c..0ca0745d2b4 100755
--- a/cuda_bindings_12/tests/cython/build_tests.sh
+++ b/cuda_bindings_12/tests/cython/build_tests.sh
@@ -1,4 +1,5 @@
#!/bin/bash
+set -eo pipefail
# SPDX-FileCopyrightText: Copyright (c) 2024-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: Apache-2.0
@@ -6,12 +7,17 @@
UNAME=$(uname)
if [ "$UNAME" == "Linux" ] ; then
SCRIPTPATH=$(dirname $(realpath "$0"))
- export CPLUS_INCLUDE_PATH=$CUDA_HOME/include:$CPLUS_INCLUDE_PATH
+ export CPLUS_INCLUDE_PATH=$CUDA_HOME/include:${CPLUS_INCLUDE_PATH:-}
elif [[ "$UNAME" == CYGWIN* || "$UNAME" == MINGW* || "$UNAME" == MSYS* ]] ; then
SCRIPTPATH="$(dirname $(cygpath -w $(realpath "$0")))"
- export CL="/I\"${CUDA_HOME}\\include\" ${CL}"
+ export CL="/I\"${CUDA_HOME}\\include\" ${CL:-}"
else
exit 1
fi
-cythonize -3 -i -Xfreethreading_compatible=True ${SCRIPTPATH}/test_*.pyx
+# Use a Python driver so the cuda.bindings source root is resolved at
+# runtime and passed via Cython's include_path -- avoids platform-specific
+# PYTHONPATH separator handling and surfaces import errors as exceptions.
+# nthreads=1 inside the driver mirrors the previous `-j 1` to side-step
+# any process-pool issues and keep builds deterministic.
+python "${SCRIPTPATH}/build_tests.py"
diff --git a/cuda_bindings_12/tests/pytest.ini b/cuda_bindings_12/tests/pytest.ini
index f687618f5be..54922832444 100644
--- a/cuda_bindings_12/tests/pytest.ini
+++ b/cuda_bindings_12/tests/pytest.ini
@@ -4,6 +4,8 @@
[pytest]
norecursedirs = cython
markers =
+ # Keep this authorship marker registry in sync across all pytest config roots.
+ # Search for "agent_authored(model)" before editing.
agent_authored(model): agent-authored test not yet materially human-reviewed
human_reviewed: agent-authored test materially reviewed or rewritten by a human
human_authored: test authored primarily by a human
diff --git a/cuda_bindings_12/tests/test_build_hooks.py b/cuda_bindings_12/tests/test_build_hooks.py
new file mode 100644
index 00000000000..f4f0f77c284
--- /dev/null
+++ b/cuda_bindings_12/tests/test_build_hooks.py
@@ -0,0 +1,358 @@
+# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+#
+# SPDX-License-Identifier: Apache-2.0
+
+"""Tests for cuda_bindings_12/build_hooks.py build infrastructure.
+
+These tests do NOT require cuda.bindings to be built/installed since they test
+build-time infrastructure. Run with --noconftest to avoid loading conftest.py
+which imports cuda.bindings modules:
+
+ pytest tests/test_build_hooks.py -v --noconftest
+
+These tests require Cython to be installed (build_hooks.py imports it).
+"""
+
+import importlib.util
+import os
+import sys
+from pathlib import Path
+
+# build_hooks.py imports Cython and setuptools at the top level; both are
+# declared test dependencies, so a missing install must surface as an
+# ImportError at collection time rather than being hidden by importorskip.
+import Cython # noqa: F401
+import pytest
+import setuptools # noqa: F401
+
+
+def _load_module(name, path, *, register=False):
+ """Load a module from source without permanently modifying sys.path.
+
+ build_hooks.py and _build_shared.py are PEP 517 backend files, not
+ installed modules. We use importlib to load them directly from source to
+ avoid polluting sys.path with the package directory (which contains
+ cuda/ source that could shadow the installed package). With ``register``
+ the module is also entered into sys.modules, which is how build_hooks.py's
+ ``from _build_shared import ...`` finds this copy.
+ """
+ spec = importlib.util.spec_from_file_location(name, path)
+ module = importlib.util.module_from_spec(spec)
+ if register:
+ sys.modules[name] = module
+ spec.loader.exec_module(module)
+ return module
+
+
+# Load the modules once at import time; _build_shared must come first.
+_PACKAGE_ROOT = Path(__file__).parent.parent
+_build_shared = _load_module("_build_shared", _PACKAGE_ROOT / "_build_shared.py", register=True)
+build_hooks = _load_module("build_hooks", _PACKAGE_ROOT / "build_hooks.py")
+
+
+@pytest.fixture(autouse=True)
+def _isolate_toolchain_env():
+ names = ("CUDA_PYTHON_TOOLCHAIN", "CC", "CXX", "LDSHARED", "LDCXXSHARED", "CUDA_PYTHON_CYTHON_CACHE_DIR")
+ original = {name: os.environ[name] for name in names if name in os.environ}
+ for name in names:
+ os.environ.pop(name, None)
+ try:
+ yield
+ finally:
+ for name in names:
+ os.environ.pop(name, None)
+ os.environ.update(original)
+
+
+@pytest.fixture
+def stamp(tmp_path, monkeypatch):
+ """Redirect the toolchain stamp to a scratch path and reset the shared force flag."""
+ scratch = tmp_path / "build" / ".build-toolchain"
+ monkeypatch.setattr(build_hooks, "_BUILD_TOOLCHAIN_STAMP", scratch)
+ monkeypatch.setattr(_build_shared, "force_build_ext", False)
+ monkeypatch.delenv("CUDA_PYTHON_TOOLCHAIN", raising=False)
+ return scratch
+
+
+def _write_stamp(stamp, toolchain):
+ stamp.parent.mkdir(parents=True, exist_ok=True)
+ stamp.write_text(toolchain + "\n")
+
+
+class TestBuildToolchainStamp:
+ """cuda.bindings stamps the toolchain name, through _check_build_toolchain() and record_build_toolchain().
+
+ The stamp-and-force protocol itself is tested by TestBuildKeyStamp.
+ """
+
+ @pytest.mark.agent_authored(model="glm-5.2")
+ def test_same_toolchain_does_not_force(self, stamp):
+ _write_stamp(stamp, "gnu")
+ build_hooks._check_build_toolchain("gnu")
+ assert build_hooks.force_build_ext is False
+
+ @pytest.mark.agent_authored(model="glm-5.2")
+ def test_changed_toolchain_forces_rebuild(self, stamp):
+ _write_stamp(stamp, "gnu")
+ build_hooks._check_build_toolchain("llvm")
+ assert build_hooks.force_build_ext is True
+
+ @pytest.mark.agent_authored(model="glm-5.2")
+ def test_record_writes_stamp(self, stamp):
+ # record_build_toolchain re-derives from env (CUDA_PYTHON_TOOLCHAIN unset →
+ # platform default: gnu on Linux, msvc on Windows).
+ build_hooks.record_build_toolchain()
+ expected = "msvc" if sys.platform == "win32" else "gnu"
+ assert stamp.read_text().strip() == expected
+
+
+# ---------------------------------------------------------------------------
+# Cython cache path helper (workaround for cython/cython#7532)
+#
+# These tests cover the configuration-digest workaround in build_hooks.py.
+# They can be deleted together with the `_cython_cache_path` helper once
+# cython/cython#7532 is resolved in a released Cython version and
+# cuda-python's minimum Cython version includes the fix.
+# See https://github.com/cython/cython/issues/7532
+
+
+_test_helpers_root = Path(__file__).parents[2] / "cuda_python_test_helpers"
+if _test_helpers_root.is_dir() and str(_test_helpers_root) not in sys.path:
+ sys.path.insert(0, str(_test_helpers_root))
+
+from cuda_python_test_helpers.build_shared import (
+ AbiStampPathMixin,
+ BuildKeyStampMixin,
+ CheckToolchainAvailableMixin,
+ DistutilsLinkerIntegrationMixin,
+ DropEmptyVenvIncludeDirMixin,
+ ForceBuildExtReexportMixin,
+ ResolveToolchainMixin,
+ WithCompilerMixin,
+ WithSccacheMixin,
+)
+from cuda_python_test_helpers.cython_cache import POSIX_ONLY_CACHE, CythonAliasMixin, CythonCachePathMixin
+
+
+class TestResolveToolchainShared(ResolveToolchainMixin):
+ build_shared = _build_shared
+
+
+class TestWithSccache(WithSccacheMixin):
+ build_shared = _build_shared
+
+
+class TestWithCompiler(WithCompilerMixin):
+ build_shared = _build_shared
+
+
+class TestDistutilsLinkerIntegration(DistutilsLinkerIntegrationMixin):
+ build_shared = _build_shared
+
+
+class TestCheckToolchainAvailable(CheckToolchainAvailableMixin):
+ build_shared = _build_shared
+
+
+class TestAbiStampPath(AbiStampPathMixin):
+ build_shared = _build_shared
+
+
+class TestBuildKeyStamp(BuildKeyStampMixin):
+ build_shared = _build_shared
+
+
+class TestDropEmptyVenvIncludeDir(DropEmptyVenvIncludeDirMixin):
+ build_shared = _build_shared
+
+
+class TestForceBuildExtReexport(ForceBuildExtReexportMixin):
+ build_hooks = build_hooks
+ build_shared = _build_shared
+
+
+class TestResolveToolchain:
+ """What cuda.bindings chooses in its ``_resolve_toolchain`` wrapper."""
+
+ @pytest.mark.agent_authored(model="claude-sonnet-4-6")
+ def test_linux_flag_set(self, monkeypatch):
+ """c++14 (c++17 costs ~15% on launch benchmarks), plus -Wno-deprecated-declarations; no -Werror."""
+ if sys.platform == "win32":
+ pytest.skip("Linux flags only")
+ monkeypatch.delenv("CUDA_PYTHON_TOOLCHAIN", raising=False)
+ _name, _cc, _cxx, cargs, _largs = build_hooks._resolve_toolchain(debug=False)
+ assert "-std=c++14" in cargs
+ assert "-Wno-deprecated-declarations" in cargs
+ assert "-Werror" not in cargs
+
+ @pytest.mark.agent_authored(model="claude-sonnet-4-6")
+ def test_msvc_flag_set(self, monkeypatch):
+ if sys.platform != "win32":
+ pytest.skip("MSVC flags only on Windows")
+ monkeypatch.delenv("CUDA_PYTHON_TOOLCHAIN", raising=False)
+ _name, _cc, _cxx, cargs, _largs = build_hooks._resolve_toolchain(debug=False)
+ assert "/std:c++14" in cargs
+ assert "/WX" not in cargs
+
+
+class TestCythonCachePath(CythonCachePathMixin):
+ """`_cython_cache_path` tests specific to cuda.bindings.
+
+ Inherits the common tests from CythonCachePathMixin; the mixin
+ covers the package-agnostic behavior. cuda.bindings does not pass
+ ``compile_time_env`` or ``cuda_major``, so the debug-only partition is
+ tested here.
+ """
+
+ build_hooks = build_hooks
+ package = "cuda-bindings-12"
+
+ @POSIX_ONLY_CACHE
+ @pytest.mark.agent_authored(model="grok-4.6")
+ def test_changed_debug_changes_namespace(self, monkeypatch, tmp_path):
+ """``debug`` partitions the namespace."""
+ self._set_env(monkeypatch, str(tmp_path))
+ p1 = build_hooks._cython_cache_path("cuda-bindings-12", debug=False)
+ p2 = build_hooks._cython_cache_path("cuda-bindings-12", debug=True)
+ assert p1 != p2
+
+
+class TestCythonCacheSmokeTest:
+ """Real Cython cache miss/hit through `_cython_cache_path`.
+
+ The actual cythonize exercise lives in
+ ``cuda_python_test_helpers.cython_cache`` so it is shared with
+ ``cuda_core/tests/test_build_hooks.py``.
+ """
+
+ @POSIX_ONLY_CACHE
+ @pytest.mark.agent_authored(model="grok-4.6")
+ def test_cache_miss_then_hit(self, monkeypatch, tmp_path, capsys):
+ from cuda_python_test_helpers.cython_cache import (
+ cython_cache_miss_then_hit,
+ )
+
+ cache_root = tmp_path / "cython-cache"
+ cache_root.mkdir()
+ monkeypatch.setenv("CUDA_PYTHON_CYTHON_CACHE_DIR", str(cache_root))
+
+ cache_path = build_hooks._cython_cache_path(
+ "cuda-bindings-12",
+ compiler_directives={"language_level": 3},
+ language_level=3,
+ cplus=True,
+ )
+ assert cache_path is not None
+ cython_cache_miss_then_hit(cache_path, tmp_path, capsys)
+
+
+class TestCythonAlias(CythonAliasMixin):
+ """`_stable_cython_alias` tests for cuda.bindings."""
+
+ build_hooks = build_hooks
+
+
+# ---------------------------------------------------------------------------
+# CUDA header check
+#
+# A cuda-bindings source tree is generated from one CUDA header set. It
+# compiles only against a toolkit of that major.minor. Against another minor,
+# the C++ compile fails with redefinition errors that do not name the cause.
+# _check_cuda_headers() reads both versions before cythonize and fails early
+# with a message that does. No GPU needed.
+
+
+def _write_cuda_h(tmp_path, cuda_version):
+ include = tmp_path / "include"
+ include.mkdir(exist_ok=True)
+ (include / "cuda.h").write_text(f"#define CUDA_VERSION {cuda_version}\n", encoding="utf-8")
+ return str(tmp_path)
+
+
+class TestCudaHeaderCheck:
+ @pytest.mark.agent_authored(model="claude-fable-5-1")
+ def test_generated_header_version_is_read_from_cydriver_pxd(self):
+ generated = build_hooks._generated_cuda_version()
+ assert generated // 1000 in (12, 13)
+ assert build_hooks._major_minor(13040) == "13.4"
+ assert build_hooks._major_minor(12090) == "12.9"
+
+ @pytest.mark.agent_authored(model="claude-fable-5-1")
+ def test_a_header_of_the_generated_major_minor_passes(self, tmp_path):
+ generated = build_hooks._generated_cuda_version()
+ build_hooks._check_cuda_headers(_write_cuda_h(tmp_path, generated))
+ # Only major.minor matters. The last digit, as in 13041, is a toolkit patch.
+ build_hooks._check_cuda_headers(_write_cuda_h(tmp_path, generated + 1))
+
+ @pytest.mark.agent_authored(model="claude-fable-5-1")
+ @pytest.mark.parametrize("delta", [-10, 10, -1000, 1000])
+ def test_another_header_fails_and_names_both_versions(self, tmp_path, delta):
+ generated = build_hooks._generated_cuda_version()
+ cuda_path = _write_cuda_h(tmp_path, generated + delta)
+ with pytest.raises(RuntimeError) as excinfo:
+ build_hooks._check_cuda_headers(cuda_path)
+ message = str(excinfo.value)
+ needed, found = build_hooks._major_minor(generated), build_hooks._major_minor(generated + delta)
+ assert message.startswith(f"This cuda-bindings source tree needs CUDA {needed} headers, but ")
+ assert os.path.realpath(os.path.join(cuda_path, "include", "cuda.h")) in message # the resolved path
+ assert f" is CUDA {found}. This is a build-time requirement only" in message
+ assert build_hooks._INSTALL_URL in message
+ assert message.endswith(
+ f"Point CUDA_PATH or CUDA_HOME at a CUDA {needed} toolkit, or build from cuda-bindings {found}.x sources."
+ )
+
+ @pytest.mark.agent_authored(model="claude-fable-5-1")
+ def test_an_unreadable_cuda_h_is_a_clear_error(self, tmp_path):
+ with pytest.raises(RuntimeError, match=r"Cannot read CUDA_VERSION from .*cuda\.h"):
+ build_hooks._check_cuda_headers(str(tmp_path)) # no include/cuda.h
+ (tmp_path / "include").mkdir()
+ (tmp_path / "include" / "cuda.h").write_text("/* no CUDA_VERSION macro */\n", encoding="utf-8")
+ with pytest.raises(RuntimeError, match=r"Cannot read CUDA_VERSION from .*cuda\.h"):
+ build_hooks._check_cuda_headers(str(tmp_path))
+
+ @pytest.mark.agent_authored(model="gpt-6")
+ def test_the_build_checks_the_header_before_it_touches_the_tree(self, tmp_path, monkeypatch):
+ cuda_path = _write_cuda_h(tmp_path, build_hooks._generated_cuda_version() + 10)
+ monkeypatch.setattr(build_hooks, "_get_cuda_paths", lambda: [cuda_path])
+
+ def not_reached(*args):
+ raise AssertionError("the header check must run before parsing or generating source files")
+
+ monkeypatch.setattr(build_hooks, "_parse_headers", not_reached)
+ monkeypatch.setattr(build_hooks, "_generate_output", not_reached)
+ with pytest.raises(RuntimeError, match="source tree needs CUDA .* headers"):
+ build_hooks._build_cuda_bindings()
+
+ @pytest.mark.agent_authored(model="gpt-6")
+ def test_split_toolkit_checks_first_available_cuda_header(self, tmp_path, monkeypatch):
+ library_root = tmp_path / "libraries"
+ library_root.mkdir()
+ header_root = tmp_path / "headers"
+ header_root.mkdir()
+ cuda_path = _write_cuda_h(header_root, build_hooks._generated_cuda_version())
+ monkeypatch.setattr(build_hooks, "_get_cuda_paths", lambda: [str(library_root), cuda_path])
+
+ def stop_before_parsing(_required_headers, include_paths):
+ assert include_paths == [str(library_root / "include"), str(header_root / "include")]
+ raise RuntimeError("matching header accepted before parsing")
+
+ monkeypatch.setattr(build_hooks, "_fetch_header_paths", stop_before_parsing)
+ with pytest.raises(RuntimeError, match="matching header accepted before parsing"):
+ build_hooks._build_cuda_bindings()
+
+ @pytest.mark.agent_authored(model="gpt-6")
+ def test_split_toolkit_rejects_mismatched_header_before_later_matching_header(self, tmp_path, monkeypatch):
+ first = tmp_path / "first"
+ second = tmp_path / "second"
+ first.mkdir()
+ second.mkdir()
+ generated = build_hooks._generated_cuda_version()
+ paths = [_write_cuda_h(first, generated + 10), _write_cuda_h(second, generated)]
+ monkeypatch.setattr(build_hooks, "_get_cuda_paths", lambda: paths)
+
+ def not_reached(*args):
+ raise AssertionError("must reject the first header in compiler search order")
+
+ monkeypatch.setattr(build_hooks, "_fetch_header_paths", not_reached)
+ with pytest.raises(RuntimeError, match="source tree needs CUDA .* headers"):
+ build_hooks._build_cuda_bindings()
diff --git a/cuda_bindings_12/tests/test_cuda.py b/cuda_bindings_12/tests/test_cuda.py
index 7a40b9493a5..78dd514b773 100644
--- a/cuda_bindings_12/tests/test_cuda.py
+++ b/cuda_bindings_12/tests/test_cuda.py
@@ -540,6 +540,7 @@ def test_cuda_mem_range_attr():
@pytest.mark.skipif(driverVersionLessThan(11040) or not supportsMemoryPool(), reason="Mempool for graphs not supported")
+@pytest.mark.thread_unsafe(reason="used high memory can be higher if threaded.")
def test_cuda_graphMem_attr():
(err,) = cuda.cuInit(0)
assert err == cuda.CUresult.CUDA_SUCCESS
@@ -755,8 +756,10 @@ def test_device_get_name():
assert err == cuda.CUresult.CUDA_SUCCESS
p = subprocess.check_output(
- ["nvidia-smi", "--query-gpu=name", "--format=csv,noheader"], shell=False, stderr=subprocess.PIPE
- ) # nosec B603, B607
+ ["nvidia-smi", "--query-gpu=name", "--format=csv,noheader"], # noqa: S607
+ shell=False,
+ stderr=subprocess.PIPE,
+ )
delimiter = b"\r\n" if platform.system() == "Windows" else b"\n"
expect = p.split(delimiter)
@@ -801,7 +804,10 @@ def host_fn(userData):
err, stream = cuda.cuStreamCreate(0)
assert err == cuda.CUresult.CUDA_SUCCESS
try:
- (err,) = cuda.cuStreamBeginCapture(stream, cuda.CUstreamCaptureMode.CU_STREAM_CAPTURE_MODE_GLOBAL)
+ # Thread-local mode: pytest-run-parallel runs this test in several threads
+ # at once, and in global mode a call such as cuGraphInstantiate in one
+ # thread fails while another thread is capturing.
+ (err,) = cuda.cuStreamBeginCapture(stream, cuda.CUstreamCaptureMode.CU_STREAM_CAPTURE_MODE_THREAD_LOCAL)
assert err == cuda.CUresult.CUDA_SUCCESS
(err,) = cuda.cuLaunchHostFunc(stream, callback, ctypes.addressof(counter))
assert err == cuda.CUresult.CUDA_SUCCESS
diff --git a/cuda_bindings_12/tests/test_cudart.py b/cuda_bindings_12/tests/test_cudart.py
index 9a496925ca4..2c6dc008aff 100644
--- a/cuda_bindings_12/tests/test_cudart.py
+++ b/cuda_bindings_12/tests/test_cudart.py
@@ -37,7 +37,16 @@ def supportsSparseTexturesDeviceFilter():
def supportsCudaAPI(name):
- return name in dir(cuda) or dir(cudart)
+ return name in dir(cuda) or name in dir(cudart)
+
+
+@pytest.mark.agent_authored(model="claude-opus-5")
+def test_supportsCudaAPI():
+ # Guards the operator precedence: `name in dir(cuda) or dir(cudart)` parses
+ # as `(name in dir(cuda)) or dir(cudart)`, which is truthy for every name.
+ assert supportsCudaAPI("cudaMalloc") is True # runtime module
+ assert supportsCudaAPI("cuInit") is True # driver module
+ assert supportsCudaAPI("this_is_not_a_cuda_api") is False
def test_cudart_memcpy():
@@ -512,7 +521,10 @@ def test_cudart_cudaStreamGetCaptureInfo():
assert status == cudart.cudaStreamCaptureStatus.cudaStreamCaptureStatusNone
# start capture
- (err,) = cudart.cudaStreamBeginCapture(stream, cudart.cudaStreamCaptureMode.cudaStreamCaptureModeGlobal)
+ # Thread-local mode: pytest-run-parallel runs this test in several threads
+ # at once, and in global mode a call such as cudaGraphInstantiate in one
+ # thread fails while another thread is capturing.
+ (err,) = cudart.cudaStreamBeginCapture(stream, cudart.cudaStreamCaptureMode.cudaStreamCaptureModeThreadLocal)
assertSuccess(err)
# validate that stream is capturing now
@@ -1365,7 +1377,10 @@ def host_fn(userData):
err, stream = cudart.cudaStreamCreate()
assertSuccess(err)
try:
- (err,) = cudart.cudaStreamBeginCapture(stream, cudart.cudaStreamCaptureMode.cudaStreamCaptureModeGlobal)
+ # Thread-local mode: pytest-run-parallel runs this test in several threads
+ # at once, and in global mode a call such as cudaGraphInstantiate in one
+ # thread fails while another thread is capturing.
+ (err,) = cudart.cudaStreamBeginCapture(stream, cudart.cudaStreamCaptureMode.cudaStreamCaptureModeThreadLocal)
assertSuccess(err)
(err,) = cudart.cudaLaunchHostFunc(stream, callback, ctypes.addressof(counter))
assertSuccess(err)
diff --git a/cuda_bindings_12/tests/test_cufile.py b/cuda_bindings_12/tests/test_cufile.py
index abd7c7d5f12..6a6f0408770 100644
--- a/cuda_bindings_12/tests/test_cufile.py
+++ b/cuda_bindings_12/tests/test_cufile.py
@@ -120,8 +120,12 @@ def isSupportedFilesystem():
return False
-# Global skip condition for all tests if cuFile library is not available
-pytestmark = pytest.mark.skipif(not cufileLibraryAvailable(), reason="cuFile library not available on this system")
+# These tests open/close the process-global driver inside their bodies, so
+# concurrent copies of a test can close the driver while another copy uses it.
+pytestmark = [
+ pytest.mark.skipif(not cufileLibraryAvailable(), reason="cuFile library not available on this system"),
+ pytest.mark.thread_unsafe(reason="cuFile tests manage process-global driver lifecycle and configuration"),
+]
def test_cufile_success_defined():
@@ -1592,109 +1596,47 @@ def test_batch_io_large_operations():
@pytest.mark.skipif(
cufileVersionLessThan(1140), reason="cuFile parameter APIs require cuFile library version 1.14.0 or later"
)
+@pytest.mark.usefixtures("cufile_env_json")
+@pytest.mark.agent_authored(model="gpt-6-astra")
def test_set_get_parameter_size_t():
- """Test setting and getting size_t parameters with cuFile validation."""
+ """Round-trip parameters without changing later tests' cuFile configuration."""
+ param_val_pairs = (
+ (cufile.SizeTConfigParameter.POLLTHRESHOLD_SIZE_KB, 64),
+ (cufile.SizeTConfigParameter.PROPERTIES_MAX_DIRECT_IO_SIZE_KB, 1024),
+ (cufile.SizeTConfigParameter.PROPERTIES_MAX_DEVICE_CACHE_SIZE_KB, 512),
+ (cufile.SizeTConfigParameter.PROPERTIES_PER_BUFFER_CACHE_SIZE_KB, 128),
+ (cufile.SizeTConfigParameter.PROPERTIES_MAX_DEVICE_PINNED_MEM_SIZE_KB, 2048),
+ (cufile.SizeTConfigParameter.PROPERTIES_IO_BATCHSIZE, 16),
+ (cufile.SizeTConfigParameter.PROPERTIES_BATCH_IO_TIMEOUT_MS, 5000),
+ (cufile.SizeTConfigParameter.EXECUTION_MAX_IO_QUEUE_DEPTH, 32),
+ (cufile.SizeTConfigParameter.EXECUTION_MAX_IO_THREADS, 8),
+ (cufile.SizeTConfigParameter.EXECUTION_MIN_IO_THRESHOLD_SIZE_KB, 4),
+ (cufile.SizeTConfigParameter.EXECUTION_MAX_REQUEST_PARALLELISM, 4),
+ )
- # Initialize CUDA
(err,) = cuda.cuInit(0)
assert err == cuda.CUresult.CUDA_SUCCESS
-
err, device = cuda.cuDeviceGet(0)
assert err == cuda.CUresult.CUDA_SUCCESS
-
err, ctx = cuda.cuDevicePrimaryCtxRetain(device)
assert err == cuda.CUresult.CUDA_SUCCESS
(err,) = cuda.cuCtxSetCurrent(ctx)
assert err == cuda.CUresult.CUDA_SUCCESS
try:
- # Test setting and getting various size_t parameters
-
- # Test poll threshold size (in KB)
- poll_threshold_kb = 64 # 64KB threshold
- cufile.set_parameter_size_t(cufile.SizeTConfigParameter.POLLTHRESHOLD_SIZE_KB, poll_threshold_kb)
- retrieved_value = cufile.get_parameter_size_t(cufile.SizeTConfigParameter.POLLTHRESHOLD_SIZE_KB)
- assert retrieved_value == poll_threshold_kb, (
- f"Poll threshold mismatch: set {poll_threshold_kb}, got {retrieved_value}"
- )
-
- # Test max direct IO size (in KB)
- max_direct_io_kb = 1024 # 1MB max direct IO size
- cufile.set_parameter_size_t(cufile.SizeTConfigParameter.PROPERTIES_MAX_DIRECT_IO_SIZE_KB, max_direct_io_kb)
- retrieved_value = cufile.get_parameter_size_t(cufile.SizeTConfigParameter.PROPERTIES_MAX_DIRECT_IO_SIZE_KB)
- assert retrieved_value == max_direct_io_kb, (
- f"Max direct IO size mismatch: set {max_direct_io_kb}, got {retrieved_value}"
- )
-
- # Test max device cache size (in KB)
- max_cache_kb = 512 # 512KB max cache size
- cufile.set_parameter_size_t(cufile.SizeTConfigParameter.PROPERTIES_MAX_DEVICE_CACHE_SIZE_KB, max_cache_kb)
- retrieved_value = cufile.get_parameter_size_t(cufile.SizeTConfigParameter.PROPERTIES_MAX_DEVICE_CACHE_SIZE_KB)
- assert retrieved_value == max_cache_kb, f"Max cache size mismatch: set {max_cache_kb}, got {retrieved_value}"
-
- # Test per buffer cache size (in KB)
- per_buffer_cache_kb = 128 # 128KB per buffer cache
- cufile.set_parameter_size_t(
- cufile.SizeTConfigParameter.PROPERTIES_PER_BUFFER_CACHE_SIZE_KB, per_buffer_cache_kb
- )
- retrieved_value = cufile.get_parameter_size_t(cufile.SizeTConfigParameter.PROPERTIES_PER_BUFFER_CACHE_SIZE_KB)
- assert retrieved_value == per_buffer_cache_kb, (
- f"Per buffer cache size mismatch: set {per_buffer_cache_kb}, got {retrieved_value}"
- )
-
- # Test max device pinned memory size (in KB)
- max_pinned_kb = 2048 # 2MB max pinned memory
- cufile.set_parameter_size_t(cufile.SizeTConfigParameter.PROPERTIES_MAX_DEVICE_PINNED_MEM_SIZE_KB, max_pinned_kb)
- retrieved_value = cufile.get_parameter_size_t(
- cufile.SizeTConfigParameter.PROPERTIES_MAX_DEVICE_PINNED_MEM_SIZE_KB
- )
- assert retrieved_value == max_pinned_kb, (
- f"Max pinned memory size mismatch: set {max_pinned_kb}, got {retrieved_value}"
- )
-
- # Test IO batch size
- batch_size = 16 # 16 operations per batch
- cufile.set_parameter_size_t(cufile.SizeTConfigParameter.PROPERTIES_IO_BATCHSIZE, batch_size)
- retrieved_value = cufile.get_parameter_size_t(cufile.SizeTConfigParameter.PROPERTIES_IO_BATCHSIZE)
- assert retrieved_value == batch_size, f"IO batch size mismatch: set {batch_size}, got {retrieved_value}"
-
- # Test batch IO timeout (in milliseconds)
- timeout_ms = 5000 # 5 second timeout
- cufile.set_parameter_size_t(cufile.SizeTConfigParameter.PROPERTIES_BATCH_IO_TIMEOUT_MS, timeout_ms)
- retrieved_value = cufile.get_parameter_size_t(cufile.SizeTConfigParameter.PROPERTIES_BATCH_IO_TIMEOUT_MS)
- assert retrieved_value == timeout_ms, f"Batch IO timeout mismatch: set {timeout_ms}, got {retrieved_value}"
-
- # Test execution parameters
- max_io_queue_depth = 32 # Max 32 operations in queue
- cufile.set_parameter_size_t(cufile.SizeTConfigParameter.EXECUTION_MAX_IO_QUEUE_DEPTH, max_io_queue_depth)
- retrieved_value = cufile.get_parameter_size_t(cufile.SizeTConfigParameter.EXECUTION_MAX_IO_QUEUE_DEPTH)
- assert retrieved_value == max_io_queue_depth, (
- f"Max IO queue depth mismatch: set {max_io_queue_depth}, got {retrieved_value}"
- )
-
- max_io_threads = 8 # Max 8 IO threads
- cufile.set_parameter_size_t(cufile.SizeTConfigParameter.EXECUTION_MAX_IO_THREADS, max_io_threads)
- retrieved_value = cufile.get_parameter_size_t(cufile.SizeTConfigParameter.EXECUTION_MAX_IO_THREADS)
- assert retrieved_value == max_io_threads, (
- f"Max IO threads mismatch: set {max_io_threads}, got {retrieved_value}"
- )
-
- min_io_threshold_kb = 4 # 4KB minimum IO threshold
- cufile.set_parameter_size_t(cufile.SizeTConfigParameter.EXECUTION_MIN_IO_THRESHOLD_SIZE_KB, min_io_threshold_kb)
- retrieved_value = cufile.get_parameter_size_t(cufile.SizeTConfigParameter.EXECUTION_MIN_IO_THRESHOLD_SIZE_KB)
- assert retrieved_value == min_io_threshold_kb, (
- f"Min IO threshold mismatch: set {min_io_threshold_kb}, got {retrieved_value}"
- )
-
- max_request_parallelism = 4 # Max 4 parallel requests
- cufile.set_parameter_size_t(
- cufile.SizeTConfigParameter.EXECUTION_MAX_REQUEST_PARALLELISM, max_request_parallelism
- )
- retrieved_value = cufile.get_parameter_size_t(cufile.SizeTConfigParameter.EXECUTION_MAX_REQUEST_PARALLELISM)
- assert retrieved_value == max_request_parallelism, (
- f"Max request parallelism mismatch: set {max_request_parallelism}, got {retrieved_value}"
- )
-
+ # Read merged defaults and JSON config, rather than uninitialized pending values.
+ cufile.driver_open()
+ try:
+ originals = {param: cufile.get_parameter_size_t(param) for param, _ in param_val_pairs}
+ finally:
+ cufile.driver_close()
+
+ for param, value in param_val_pairs:
+ try:
+ cufile.set_parameter_size_t(param, value)
+ assert cufile.get_parameter_size_t(param) == value
+ finally:
+ cufile.set_parameter_size_t(param, originals[param])
finally:
cuda.cuDevicePrimaryCtxRelease(device)
@@ -1702,84 +1644,48 @@ def test_set_get_parameter_size_t():
@pytest.mark.skipif(
cufileVersionLessThan(1140), reason="cuFile parameter APIs require cuFile library version 1.14.0 or later"
)
+@pytest.mark.usefixtures("cufile_env_json")
+@pytest.mark.agent_authored(model="gpt-6-astra")
def test_set_get_parameter_bool():
- """Test setting and getting boolean parameters with cuFile validation."""
+ """Round-trip parameters without changing later tests' cuFile configuration."""
+ param_val_pairs = (
+ (cufile.BoolConfigParameter.PROPERTIES_USE_POLL_MODE, True),
+ (cufile.BoolConfigParameter.PROPERTIES_ALLOW_COMPAT_MODE, False),
+ (cufile.BoolConfigParameter.FORCE_COMPAT_MODE, False),
+ (cufile.BoolConfigParameter.FS_MISC_API_CHECK_AGGRESSIVE, True),
+ (cufile.BoolConfigParameter.EXECUTION_PARALLEL_IO, True),
+ (cufile.BoolConfigParameter.PROFILE_NVTX, False),
+ (cufile.BoolConfigParameter.PROPERTIES_ALLOW_SYSTEM_MEMORY, True),
+ (cufile.BoolConfigParameter.USE_PCIP2PDMA, True),
+ (cufile.BoolConfigParameter.PREFER_IO_URING, False),
+ (cufile.BoolConfigParameter.FORCE_ODIRECT_MODE, True),
+ (cufile.BoolConfigParameter.SKIP_TOPOLOGY_DETECTION, False),
+ (cufile.BoolConfigParameter.STREAM_MEMOPS_BYPASS, True),
+ )
- # Initialize CUDA
(err,) = cuda.cuInit(0)
assert err == cuda.CUresult.CUDA_SUCCESS
-
err, device = cuda.cuDeviceGet(0)
assert err == cuda.CUresult.CUDA_SUCCESS
-
err, ctx = cuda.cuDevicePrimaryCtxRetain(device)
assert err == cuda.CUresult.CUDA_SUCCESS
(err,) = cuda.cuCtxSetCurrent(ctx)
assert err == cuda.CUresult.CUDA_SUCCESS
try:
- # Test setting and getting various boolean parameters
-
- # Test poll mode
- cufile.set_parameter_bool(cufile.BoolConfigParameter.PROPERTIES_USE_POLL_MODE, True)
- retrieved_value = cufile.get_parameter_bool(cufile.BoolConfigParameter.PROPERTIES_USE_POLL_MODE)
- assert retrieved_value is True, f"Poll mode mismatch: set True, got {retrieved_value}"
-
- # Test compatibility mode
- cufile.set_parameter_bool(cufile.BoolConfigParameter.PROPERTIES_ALLOW_COMPAT_MODE, False)
- retrieved_value = cufile.get_parameter_bool(cufile.BoolConfigParameter.PROPERTIES_ALLOW_COMPAT_MODE)
- assert retrieved_value is False, f"Compatibility mode mismatch: set False, got {retrieved_value}"
-
- # Test force compatibility mode
- cufile.set_parameter_bool(cufile.BoolConfigParameter.FORCE_COMPAT_MODE, False)
- retrieved_value = cufile.get_parameter_bool(cufile.BoolConfigParameter.FORCE_COMPAT_MODE)
- assert retrieved_value is False, f"Force compatibility mode mismatch: set False, got {retrieved_value}"
-
- # Test aggressive API check
- cufile.set_parameter_bool(cufile.BoolConfigParameter.FS_MISC_API_CHECK_AGGRESSIVE, True)
- retrieved_value = cufile.get_parameter_bool(cufile.BoolConfigParameter.FS_MISC_API_CHECK_AGGRESSIVE)
- assert retrieved_value is True, f"Aggressive API check mismatch: set True, got {retrieved_value}"
-
- # Test parallel IO
- cufile.set_parameter_bool(cufile.BoolConfigParameter.EXECUTION_PARALLEL_IO, True)
- retrieved_value = cufile.get_parameter_bool(cufile.BoolConfigParameter.EXECUTION_PARALLEL_IO)
- assert retrieved_value is True, f"Parallel IO mismatch: set True, got {retrieved_value}"
-
- # Test NVTX profiling
- cufile.set_parameter_bool(cufile.BoolConfigParameter.PROFILE_NVTX, False)
- retrieved_value = cufile.get_parameter_bool(cufile.BoolConfigParameter.PROFILE_NVTX)
- assert retrieved_value is False, f"NVTX profiling mismatch: set False, got {retrieved_value}"
-
- # Test system memory allowance
- cufile.set_parameter_bool(cufile.BoolConfigParameter.PROPERTIES_ALLOW_SYSTEM_MEMORY, True)
- retrieved_value = cufile.get_parameter_bool(cufile.BoolConfigParameter.PROPERTIES_ALLOW_SYSTEM_MEMORY)
- assert retrieved_value is True, f"System memory allowance mismatch: set True, got {retrieved_value}"
-
- # Test PCI P2P DMA
- cufile.set_parameter_bool(cufile.BoolConfigParameter.USE_PCIP2PDMA, True)
- retrieved_value = cufile.get_parameter_bool(cufile.BoolConfigParameter.USE_PCIP2PDMA)
- assert retrieved_value is True, f"PCI P2P DMA mismatch: set True, got {retrieved_value}"
-
- # Test IO uring preference
- cufile.set_parameter_bool(cufile.BoolConfigParameter.PREFER_IO_URING, False)
- retrieved_value = cufile.get_parameter_bool(cufile.BoolConfigParameter.PREFER_IO_URING)
- assert retrieved_value is False, f"IO uring preference mismatch: set False, got {retrieved_value}"
-
- # Test force O_DIRECT mode
- cufile.set_parameter_bool(cufile.BoolConfigParameter.FORCE_ODIRECT_MODE, True)
- retrieved_value = cufile.get_parameter_bool(cufile.BoolConfigParameter.FORCE_ODIRECT_MODE)
- assert retrieved_value is True, f"Force O_DIRECT mode mismatch: set True, got {retrieved_value}"
-
- # Test topology detection skip
- cufile.set_parameter_bool(cufile.BoolConfigParameter.SKIP_TOPOLOGY_DETECTION, False)
- retrieved_value = cufile.get_parameter_bool(cufile.BoolConfigParameter.SKIP_TOPOLOGY_DETECTION)
- assert retrieved_value is False, f"Topology detection skip mismatch: set False, got {retrieved_value}"
-
- # Test stream memops bypass
- cufile.set_parameter_bool(cufile.BoolConfigParameter.STREAM_MEMOPS_BYPASS, True)
- retrieved_value = cufile.get_parameter_bool(cufile.BoolConfigParameter.STREAM_MEMOPS_BYPASS)
- assert retrieved_value is True, f"Stream memops bypass mismatch: set True, got {retrieved_value}"
-
+ # Read merged defaults and JSON config, rather than uninitialized pending values.
+ cufile.driver_open()
+ try:
+ originals = {param: cufile.get_parameter_bool(param) for param, _ in param_val_pairs}
+ finally:
+ cufile.driver_close()
+
+ for param, value in param_val_pairs:
+ try:
+ cufile.set_parameter_bool(param, value)
+ assert cufile.get_parameter_bool(param) is value
+ finally:
+ cufile.set_parameter_bool(param, originals[param])
finally:
cuda.cuDevicePrimaryCtxRelease(device)
@@ -1787,78 +1693,49 @@ def test_set_get_parameter_bool():
@pytest.mark.skipif(
cufileVersionLessThan(1140), reason="cuFile parameter APIs require cuFile library version 1.14.0 or later"
)
-def test_set_get_parameter_string():
- """Test setting and getting string parameters with cuFile validation."""
+@pytest.mark.usefixtures("cufile_env_json")
+@pytest.mark.agent_authored(model="gpt-6-astra")
+def test_set_get_parameter_string(tmp_path, monkeypatch):
+ """Round-trip string parameters and restore usable logging configuration."""
+ # The environment-path getter requires the variable to exist; preserve the caller's value.
+ monkeypatch.setenv("CUFILE_LOGFILE_PATH", os.environ.get("CUFILE_LOGFILE_PATH", ""))
+ param_val_pairs = (
+ (cufile.StringConfigParameter.LOGGING_LEVEL, "INFO"),
+ (cufile.StringConfigParameter.ENV_LOGFILE_PATH, str(tmp_path / "cufile.log")),
+ (cufile.StringConfigParameter.LOG_DIR, str(tmp_path)),
+ )
+ empty_defaults = {
+ cufile.StringConfigParameter.LOGGING_LEVEL: "DEBUG",
+ cufile.StringConfigParameter.ENV_LOGFILE_PATH: os.path.join(tempfile.gettempdir(), "cufile.log"),
+ cufile.StringConfigParameter.LOG_DIR: tempfile.gettempdir(),
+ }
- # Initialize CUDA
(err,) = cuda.cuInit(0)
assert err == cuda.CUresult.CUDA_SUCCESS
-
err, device = cuda.cuDeviceGet(0)
assert err == cuda.CUresult.CUDA_SUCCESS
-
err, ctx = cuda.cuDevicePrimaryCtxRetain(device)
assert err == cuda.CUresult.CUDA_SUCCESS
(err,) = cuda.cuCtxSetCurrent(ctx)
assert err == cuda.CUresult.CUDA_SUCCESS
try:
- # Test setting and getting various string parameters
- # Note: String parameter tests may have issues with the current implementation
-
- # Test logging level
- logging_level = "INFO"
- try:
- # Convert Python string to null-terminated C string
- logging_level_bytes = logging_level.encode("utf-8") + b"\x00"
- logging_level_buffer = ctypes.create_string_buffer(logging_level_bytes)
- cufile.set_parameter_string(
- cufile.StringConfigParameter.LOGGING_LEVEL, int(ctypes.addressof(logging_level_buffer))
- )
- retrieved_value = cufile.get_parameter_string(cufile.StringConfigParameter.LOGGING_LEVEL, 256)
- logging.info(f"Logging level test: set {logging_level}, got {retrieved_value}")
- # The retrieved value should be a string, so we can compare directly
- assert retrieved_value == logging_level, (
- f"Logging level mismatch: set {logging_level}, got {retrieved_value}"
- )
- except Exception as e:
- logging.error(f"Logging level test failed: {e}")
- # Re-raise the exception to make the test fail
- raise
-
- # Test environment log file path
- logfile_path = tempfile.gettempdir() + "/cufile.log"
+ cufile.driver_open()
try:
- # Convert Python string to null-terminated C string
- logfile_path_bytes = logfile_path.encode("utf-8") + b"\x00"
- logfile_buffer = ctypes.create_string_buffer(logfile_path_bytes)
- cufile.set_parameter_string(
- cufile.StringConfigParameter.ENV_LOGFILE_PATH, int(ctypes.addressof(logfile_buffer))
- )
- retrieved_value = cufile.get_parameter_string(cufile.StringConfigParameter.ENV_LOGFILE_PATH, 256)
- logging.info(f"Log file path test: set {logfile_path}, got {retrieved_value}")
- # The retrieved value should be a string, so we can compare directly
- assert retrieved_value == logfile_path, f"Log file path mismatch: set {logfile_path}, got {retrieved_value}"
- except Exception as e:
- logging.error(f"Log file path test failed: {e}")
- # Re-raise the exception to make the test fail
- raise
-
- # Test log directory
- log_dir = tempfile.gettempdir() + "/cufile_logs"
- try:
- # Convert Python string to null-terminated C string
- log_dir_bytes = log_dir.encode("utf-8") + b"\x00"
- log_dir_buffer = ctypes.create_string_buffer(log_dir_bytes)
- cufile.set_parameter_string(cufile.StringConfigParameter.LOG_DIR, int(ctypes.addressof(log_dir_buffer)))
- retrieved_value = cufile.get_parameter_string(cufile.StringConfigParameter.LOG_DIR, 256)
- logging.info(f"Log directory test: set {log_dir}, got {retrieved_value}")
- # The retrieved value should be a string, so we can compare directly
- assert retrieved_value == log_dir, f"Log directory mismatch: set {log_dir}, got {retrieved_value}"
- except Exception as e:
- logging.error(f"Log directory test failed: {e}")
- # Re-raise the exception to make the test fail
- raise
-
+ originals = {param: cufile.get_parameter_string(param, 256) for param, _ in param_val_pairs}
+ finally:
+ cufile.driver_close()
+
+ for param, value in param_val_pairs:
+ value_buffer = ctypes.create_string_buffer(value.encode("utf-8"))
+ # cuFile returns empty path defaults but rejects them in its setter. Restore
+ # valid temporary paths in that case, as in the current bindings tests.
+ original = originals[param] or empty_defaults[param]
+ original_buffer = ctypes.create_string_buffer(original.encode("utf-8"))
+ try:
+ cufile.set_parameter_string(param, ctypes.addressof(value_buffer))
+ assert cufile.get_parameter_string(param, 256) == value
+ finally:
+ cufile.set_parameter_string(param, ctypes.addressof(original_buffer))
finally:
cuda.cuDevicePrimaryCtxRelease(device)
diff --git a/cuda_bindings_12/tests/test_graphics_apis.py b/cuda_bindings_12/tests/test_graphics_apis.py
index 0cfbac4e870..365bbe1fb2c 100644
--- a/cuda_bindings_12/tests/test_graphics_apis.py
+++ b/cuda_bindings_12/tests/test_graphics_apis.py
@@ -1,12 +1,18 @@
# SPDX-FileCopyrightText: Copyright (c) 2025-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: Apache-2.0
+import os
+import sys
+
import pytest
from cuda.bindings import runtime as cudart
def test_graphics_api_smoketest():
+ if sys.platform.startswith("linux") and not (os.environ.get("DISPLAY") or os.environ.get("WAYLAND_DISPLAY")):
+ pytest.skip("No display available for the OpenGL smoke test")
+
# Due to lazy importing in pyglet, pytest.importorskip doesn't work
try:
import pyglet
diff --git a/cuda_bindings_12/tests/test_kernelParams.py b/cuda_bindings_12/tests/test_kernelParams.py
index 0f11615749f..ed8f02d7473 100644
--- a/cuda_bindings_12/tests/test_kernelParams.py
+++ b/cuda_bindings_12/tests/test_kernelParams.py
@@ -863,20 +863,21 @@ def __init__(self, address, typestr):
ASSERT_DRV(err)
+@pytest.mark.agent_authored(model="gpt-5")
def test_kernelParams_c_int_out_of_range_raises():
- # #363: an out-of-range Python int for a c_int / c_byte kernel argument must
- # raise instead of being silently truncated to fit the declared width.
(err,) = cuda.cuInit(0)
ASSERT_DRV(err)
- err, cuDevice = cuda.cuDeviceGet(0)
+ err, device = cuda.cuDeviceGet(0)
ASSERT_DRV(err)
- err, context = cuda.cuCtxCreate(0, cuDevice)
+ err, context = cuda.cuCtxCreate(0, device)
ASSERT_DRV(err)
+ # #363: an out-of-range Python int for a c_int / c_byte kernel argument must
+ # raise instead of being silently truncated to fit the declared width.
kernelString = """\
extern "C" __global__ void take_int(int i) {}
"""
- module = common_nvrtc(kernelString, cuDevice)
+ module = common_nvrtc(kernelString, device)
err, kernel = cuda.cuModuleGetFunction(module, b"take_int")
ASSERT_DRV(err)
err, stream = cuda.cuStreamCreate(0)
diff --git a/cuda_bindings_12/tests/test_nvfatbin.py b/cuda_bindings_12/tests/test_nvfatbin.py
index 383e5bc4078..8049c9ddfb5 100644
--- a/cuda_bindings_12/tests/test_nvfatbin.py
+++ b/cuda_bindings_12/tests/test_nvfatbin.py
@@ -102,7 +102,7 @@ def nvcc_smoke(tmpdir) -> str:
with open(src, "w") as f:
f.write("")
try:
- subprocess.run( # nosec: S603
+ subprocess.run( # noqa: S603
[nvcc, "-c", str(src), "-o", str(out)],
check=True,
capture_output=True,
@@ -182,7 +182,7 @@ def OBJECT(arch, tmpdir, nvcc_smoke):
# compile a temporary CUDA translation unit.
cmd = [nvcc, "-c", "-arch", arch, "-o", str(tmpdir / "object.o"), str(tmpdir / "object.cu")]
try:
- subprocess.run( # nosec: S603
+ subprocess.run( # noqa: S603
cmd,
check=True,
capture_output=True,
diff --git a/cuda_bindings_12/tests/test_utils.py b/cuda_bindings_12/tests/test_utils.py
index c9afe1f415b..8904d65d2ed 100644
--- a/cuda_bindings_12/tests/test_utils.py
+++ b/cuda_bindings_12/tests/test_utils.py
@@ -112,6 +112,6 @@ def test_get_handle_error(target):
],
)
def test_cyclical_imports(module):
- subprocess.check_call( # nosec B603
+ subprocess.check_call( # noqa: S603
[sys.executable, Path(__file__).parent / "utils" / "check_cyclical_import.py", f"cuda.bindings.{module}"],
)
diff --git a/cuda_bindings_13/MANIFEST.in b/cuda_bindings_13/MANIFEST.in
index 8bc25bbbd77..30c4601806e 100644
--- a/cuda_bindings_13/MANIFEST.in
+++ b/cuda_bindings_13/MANIFEST.in
@@ -2,6 +2,7 @@
# SPDX-License-Identifier: Apache-2.0
recursive-include cuda/ *.pyx *.pxd *.pxi *.pyx.in *.pxd.in *.pxi.in *.h
+include build_hooks.py
# at least with setuptools 75.0.0 this folder was added erroneously
# to the payload, causing file copying to the build environment failed
exclude cuda/bindings cuda?bindings
diff --git a/cuda_bindings_13/tests/conftest.py b/cuda_bindings_13/tests/conftest.py
index 926248450ef..1fcb8b6432f 100644
--- a/cuda_bindings_13/tests/conftest.py
+++ b/cuda_bindings_13/tests/conftest.py
@@ -13,7 +13,11 @@
import cuda.bindings.driver as cuda
-# Keep in sync with cuda_core/tests/conftest.py.
+# BEGIN SYNCED PYTEST PLUGIN BOOTSTRAP
+# Keep every block with this marker byte-for-byte identical.
+# Before editing or reviewing, find all copies; see root AGENTS.md.
+# Keep the shared pytest plugin available in wheel-test jobs, which run from
+# the monorepo checkout without installing cuda-python-test-helpers.
try:
import cuda_python_test_helpers._pytest_plugin # noqa: F401
except ImportError as e:
@@ -29,6 +33,7 @@
importlib.invalidate_caches()
pytest_plugins = ["cuda_python_test_helpers._pytest_plugin"]
+# END SYNCED PYTEST PLUGIN BOOTSTRAP
def pytest_configure(config):
diff --git a/cuda_core/MANIFEST.in b/cuda_core/MANIFEST.in
index 9f34de2b363..f2f22eda4ff 100644
--- a/cuda_core/MANIFEST.in
+++ b/cuda_core/MANIFEST.in
@@ -4,7 +4,8 @@
recursive-include cuda/core *.pyx *.pxd *.pxi *.pyi
recursive-include cuda/core/_cpp *.cpp *.h *.hpp
-recursive-include cuda/core/_include *.h *.hpp
+recursive-include cuda/core/_include *.h *.hpp *.def
+include build_hooks.py
include cuda/core/py.typed
include NOTICE
# package-local shared PEP 517 helper
diff --git a/cuda_core/tests/conftest.py b/cuda_core/tests/conftest.py
index 012de8e7abc..e62f89b8251 100644
--- a/cuda_core/tests/conftest.py
+++ b/cuda_core/tests/conftest.py
@@ -12,7 +12,11 @@
import pytest
-# Keep in sync with cuda_bindings_13/tests/conftest.py.
+# BEGIN SYNCED PYTEST PLUGIN BOOTSTRAP
+# Keep every block with this marker byte-for-byte identical.
+# Before editing or reviewing, find all copies; see root AGENTS.md.
+# Keep the shared pytest plugin available in wheel-test jobs, which run from
+# the monorepo checkout without installing cuda-python-test-helpers.
try:
import cuda_python_test_helpers._pytest_plugin # noqa: F401
except ImportError as e:
@@ -28,6 +32,7 @@
importlib.invalidate_caches()
pytest_plugins = ["cuda_python_test_helpers._pytest_plugin"]
+# END SYNCED PYTEST PLUGIN BOOTSTRAP
from helpers.constants import POOL_SIZE
from helpers.memory import (
diff --git a/cuda_python_test_helpers/cuda_python_test_helpers/_pytest_plugin.py b/cuda_python_test_helpers/cuda_python_test_helpers/_pytest_plugin.py
index 831c8aca4c6..dd06ee0ae60 100644
--- a/cuda_python_test_helpers/cuda_python_test_helpers/_pytest_plugin.py
+++ b/cuda_python_test_helpers/cuda_python_test_helpers/_pytest_plugin.py
@@ -32,7 +32,11 @@ def pytest_collection_modifyitems(config, items): # noqa: ARG001
# Package markers by path
if nodeid.startswith("cuda_pathfinder/tests/") or "/cuda_pathfinder/tests/" in nodeid:
item.add_marker(pytest.mark.pathfinder)
- if nodeid.startswith("cuda_bindings_13/tests/") or "/cuda_bindings_13/tests/" in nodeid:
+ if (
+ nodeid.startswith(("cuda_bindings_12/tests/", "cuda_bindings_13/tests/"))
+ or "/cuda_bindings_12/tests/" in nodeid
+ or "/cuda_bindings_13/tests/" in nodeid
+ ):
item.add_marker(pytest.mark.bindings)
if nodeid.startswith("cuda_core/tests/") or "/cuda_core/tests/" in nodeid:
item.add_marker(pytest.mark.core)
diff --git a/ruff.toml b/ruff.toml
index 120842435f5..5a2778cfb80 100644
--- a/ruff.toml
+++ b/ruff.toml
@@ -140,7 +140,7 @@ inline-quotes = "double"
# CUDA bindings mirror C API naming conventions (CamelCase types, camelCase functions)
# Keep examples opted-in to enforce naming conventions in example-local identifiers.
-"cuda_bindings_13/{cuda,docs,tests}/**" = [
+"{cuda_bindings_12,cuda_bindings_13}/{cuda,docs,tests}/**" = [
"N801", # invalid-class-name
"N802", # invalid-function-name
"N803", # invalid-argument-name
@@ -152,7 +152,7 @@ inline-quotes = "double"
"N802", # invalid-function-name
"N806", # non-lowercase-variable-in-function
]
-"cuda_bindings_13/{build_hooks.py,setup.py}" = ["N801", "N802", "N803", "N806", "N816"]
+"{cuda_bindings_12,cuda_bindings_13}/{build_hooks.py,setup.py}" = ["N801", "N802", "N803", "N806", "N816"]
# scripts and build tooling — print is the expected output method
"qa/**" = ["T201"]
diff --git a/toolshed/check_generated_file_seals.py b/toolshed/check_generated_file_seals.py
index 0cc861a1f4e..cc63fdaabba 100644
--- a/toolshed/check_generated_file_seals.py
+++ b/toolshed/check_generated_file_seals.py
@@ -62,7 +62,7 @@ def load_previously_sealed_paths():
GENERATED_FILE_MARKER_FRAGMENT,
"HEAD",
"--",
- "cuda_bindings_13",
+ ":(glob)cuda_bindings_[0-9]*/**",
],
capture_output=True,
text=True,
diff --git a/toolshed/check_spdx.py b/toolshed/check_spdx.py
index c1de07ff891..1be8a565db7 100644
--- a/toolshed/check_spdx.py
+++ b/toolshed/check_spdx.py
@@ -26,6 +26,7 @@
"benchmarks": "Apache-2.0",
"ci": "Apache-2.0",
"cuda_bindings_13": "Apache-2.0",
+ "cuda_bindings_12": "Apache-2.0",
"cuda_core": "Apache-2.0",
"cuda_pathfinder": "Apache-2.0",
"cuda_python": "Apache-2.0",