From 582cf852a00e1c75217f853ba7c9798034cd5c53 Mon Sep 17 00:00:00 2001 From: "Ralf W. Grosse-Kunstleve" Date: Tue, 20 Oct 2026 21:40:07 -0700 Subject: [PATCH 01/10] Integrate the CUDA 12 source with shared build and test helpers Signed-off-by: Ralf W. Grosse-Kunstleve --- .github/labeler.yml | 1 + .gitignore | 19 + benchmarks/cuda_bindings/runner/runtime.py | 5 +- cuda_bindings_12/DESCRIPTION.rst | 8 +- cuda_bindings_12/LICENSE | 25 ++ cuda_bindings_12/MANIFEST.in | 5 +- cuda_bindings_12/README.md | 6 +- cuda_bindings_12/_build_shared.py | 1 + cuda_bindings_12/benchmarks/pytest.ini | 6 + cuda_bindings_12/build_hooks.py | 245 ++++++++++-- .../cuda/bindings/_internal/utils.pxd | 6 +- .../cuda/bindings/_internal/utils.pyx | 9 +- cuda_bindings_12/cuda/bindings/_lib/utils.pxd | 1 + cuda_bindings_12/cuda/bindings/_lib/utils.pxi | 6 +- .../cuda/bindings/_lib/windll.pxd | 2 +- .../cuda/bindings/utils/__init__.py | 5 +- cuda_bindings_12/docs/README.md | 7 +- cuda_bindings_12/docs/build_docs.sh | 14 +- .../_static/javascripts/version_dropdown.js | 30 +- .../_templates/sidebar/variant-selector.html | 15 +- cuda_bindings_12/docs/source/contribute.md | 2 +- cuda_bindings_12/docs/source/install.md | 8 +- cuda_bindings_12/docs/source/motivation.md | 14 +- cuda_bindings_12/docs/source/overview.md | 114 +++--- cuda_bindings_12/docs/source/release.rst | 29 +- .../docs/source/release/11.7.1-notes.md | 1 + .../docs/source/release/11.8.0-notes.md | 2 +- .../docs/source/release/12.9.1-notes.rst | 51 +++ .../docs/source/release/12.9.2-notes.rst | 21 + .../docs/source/release/12.9.3-notes.rst | 29 ++ .../docs/source/release/12.9.4-notes.rst | 23 ++ .../docs/source/release/12.9.5-notes.rst | 29 ++ .../docs/source/release/12.9.6-notes.rst | 64 ++++ .../docs/source/release/12.9.7-notes.rst | 53 +++ .../docs/source/release/12.X.Y-notes.rst | 2 + cuda_bindings_12/docs/source/support.rst | 8 +- .../docs/source/tips_and_tricks.rst | 6 +- cuda_bindings_12/docs/versions.json | 7 - .../0_Introduction/clock_nvrtc_test.py | 2 +- .../simpleCubemapTexture_test.py | 2 +- .../examples/0_Introduction/simpleP2P_test.py | 2 +- .../0_Introduction/simpleZeroCopy_test.py | 2 +- .../0_Introduction/systemWideAtomics_test.py | 2 +- .../0_Introduction/vectorAddDrv_test.py | 2 +- .../0_Introduction/vectorAddMMAP_test.py | 2 +- .../streamOrderedAllocation_test.py | 2 +- .../globalToShmemAsyncCopy_test.py | 2 +- .../3_CUDA_Features/simpleCudaGraphs_test.py | 2 +- .../conjugateGradientMultiBlockCG_test.py | 2 +- cuda_bindings_12/examples/common/common.py | 2 +- .../examples/common/helper_cuda.py | 2 +- .../examples/common/helper_string.py | 2 +- .../examples/extra/isoFDModelling_test.py | 2 +- .../examples/extra/jit_program_test.py | 2 +- .../examples/extra/numba_emm_plugin.py | 2 +- cuda_bindings_12/examples/pytest.ini | 9 + cuda_bindings_12/setup.py | 40 +- .../_cuda_bindings_redirector.py | 8 +- cuda_bindings_12/tests/conftest.py | 28 ++ cuda_bindings_12/tests/cython/build_tests.bat | 9 +- cuda_bindings_12/tests/cython/build_tests.py | 109 ++++++ cuda_bindings_12/tests/cython/build_tests.sh | 12 +- cuda_bindings_12/tests/nvml/test_cuda.py | 26 +- cuda_bindings_12/tests/nvml/test_init.py | 14 +- cuda_bindings_12/tests/nvml/test_pynvml.py | 18 - cuda_bindings_12/tests/nvml/util.py | 2 +- cuda_bindings_12/tests/pytest.ini | 2 + cuda_bindings_12/tests/test_build_hooks.py | 358 ++++++++++++++++++ cuda_bindings_12/tests/test_cuda.py | 7 +- cuda_bindings_12/tests/test_cudart.py | 11 +- cuda_bindings_12/tests/test_cufile.py | 317 +++++----------- cuda_bindings_12/tests/test_graphics_apis.py | 6 + cuda_bindings_12/tests/test_kernelParams.py | 11 +- cuda_bindings_12/tests/test_nvfatbin.py | 4 +- cuda_bindings_12/tests/test_nvjitlink.py | 7 + cuda_bindings_12/tests/test_nvvm.py | 2 +- cuda_bindings_12/tests/test_utils.py | 23 +- cuda_bindings_13/MANIFEST.in | 1 + cuda_bindings_13/tests/conftest.py | 7 +- cuda_core/MANIFEST.in | 3 +- cuda_core/tests/conftest.py | 7 +- 81 files changed, 1454 insertions(+), 500 deletions(-) create mode 120000 cuda_bindings_12/_build_shared.py create mode 100644 cuda_bindings_12/docs/source/release/12.9.1-notes.rst create mode 100644 cuda_bindings_12/docs/source/release/12.9.2-notes.rst create mode 100644 cuda_bindings_12/docs/source/release/12.9.3-notes.rst create mode 100644 cuda_bindings_12/docs/source/release/12.9.4-notes.rst create mode 100644 cuda_bindings_12/docs/source/release/12.9.5-notes.rst create mode 100644 cuda_bindings_12/docs/source/release/12.9.6-notes.rst create mode 100644 cuda_bindings_12/docs/source/release/12.9.7-notes.rst delete mode 100644 cuda_bindings_12/docs/versions.json create mode 100644 cuda_bindings_12/tests/conftest.py create mode 100644 cuda_bindings_12/tests/cython/build_tests.py create mode 100644 cuda_bindings_12/tests/test_build_hooks.py diff --git a/.github/labeler.yml b/.github/labeler.yml index 344b3d312e7..989c998fed7 100644 --- a/.github/labeler.yml +++ b/.github/labeler.yml @@ -8,6 +8,7 @@ cuda.bindings: - changed-files: - any-glob-to-any-file: + - 'cuda_bindings_12/**' - 'cuda_bindings_13/**' - 'samples/cuda_bindings/**' diff --git a/.gitignore b/.gitignore index 6010fe02f68..ea620d25a43 100644 --- a/.gitignore +++ b/.gitignore @@ -40,6 +40,25 @@ cuda_bindings_13/cuda/bindings/_internal/runtime.pyx cuda_bindings_13/cuda/bindings/_internal/runtime_ptds.pyx cuda_bindings_13/cuda/bindings/utils/_get_handle.pyx +# CUDA 12 bindings source generation +cuda_bindings_12/cuda/bindings/_bindings/cyruntime.pxd +cuda_bindings_12/cuda/bindings/_bindings/cyruntime.pyx +cuda_bindings_12/cuda/bindings/_bindings/cyruntime_ptds.pxd +cuda_bindings_12/cuda/bindings/_bindings/cyruntime_ptds.pyx +cuda_bindings_12/cuda/bindings/_internal/cufile.pyx +cuda_bindings_12/cuda/bindings/_internal/driver.pyx +cuda_bindings_12/cuda/bindings/_internal/nvfatbin.pyx +cuda_bindings_12/cuda/bindings/_internal/nvjitlink.pyx +cuda_bindings_12/cuda/bindings/_internal/nvml.pyx +cuda_bindings_12/cuda/bindings/_internal/nvrtc.pyx +cuda_bindings_12/cuda/bindings/_internal/nvvm.pyx +cuda_bindings_12/cuda/bindings/cyruntime.pxd +cuda_bindings_12/cuda/bindings/cyruntime.pyx +cuda_bindings_12/cuda/bindings/cyruntime_functions.pxi +cuda_bindings_12/cuda/bindings/cyruntime_types.pxi +cuda_bindings_12/cuda/bindings/runtime.pxd +cuda_bindings_12/cuda/bindings/runtime.pyx + # Version files from setuptools_scm _version.py diff --git a/benchmarks/cuda_bindings/runner/runtime.py b/benchmarks/cuda_bindings/runner/runtime.py index 153d92f259d..7e3ba127ff3 100644 --- a/benchmarks/cuda_bindings/runner/runtime.py +++ b/benchmarks/cuda_bindings/runner/runtime.py @@ -30,7 +30,10 @@ def ensure_context() -> int: assert_drv(err) _device = device - err, ctx = cuda.cuCtxCreate(None, 0, device) + if cuda.CUDA_VERSION < 13000: + err, ctx = cuda.cuCtxCreate(0, device) + else: + err, ctx = cuda.cuCtxCreate(None, 0, device) assert_drv(err) _ctx = ctx return ctx diff --git a/cuda_bindings_12/DESCRIPTION.rst b/cuda_bindings_12/DESCRIPTION.rst index ca0297523e0..52fbdf848e8 100644 --- a/cuda_bindings_12/DESCRIPTION.rst +++ b/cuda_bindings_12/DESCRIPTION.rst @@ -5,11 +5,11 @@ cuda-bindings: Low-level CUDA interfaces **************************************** -`cuda.bindings `_ is a standard set of low-level interfaces, providing full coverage of and 1:1 access to the CUDA host APIs from Python. Checkout the `Overview `_ for the workflow and performance results. +`cuda.bindings `_ is a standard set of low-level interfaces, providing full coverage of and 1:1 access to the CUDA host APIs from Python. Checkout the `Overview `_ for the workflow and performance results. -* `Repository `_ +* `Repository `_ * `Documentation `_ -* `Examples `_ +* `Examples `_ * `Issue tracker `_ -For the installation instruction, please refer to the `Installation `_ page. +For the installation instruction, please refer to the `Installation `_ page. diff --git a/cuda_bindings_12/LICENSE b/cuda_bindings_12/LICENSE index d6f74778be8..f3fe76ecadf 100644 --- a/cuda_bindings_12/LICENSE +++ b/cuda_bindings_12/LICENSE @@ -176,3 +176,28 @@ Copyright (c) 2024-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. of your accepting any such warranty or additional liability. END OF TERMS AND CONDITIONS + + APPENDIX: How to apply the Apache License to your work. + + To apply the Apache License to your work, attach the following + boilerplate notice, with the fields enclosed by brackets "[]" + replaced with your own identifying information. (Don't include + the brackets!) The text should be enclosed in the appropriate + comment syntax for the file format. We also recommend that a + file or class name and description of purpose be included on the + same "printed page" as the copyright notice for easier + identification within third-party archives. + + Copyright [yyyy] [name of copyright owner] + + Licensed under the Apache License, Version 2.0 (the "License"); + you may not use this file except in compliance with the License. + You may obtain a copy of the License at + + http://www.apache.org/licenses/LICENSE-2.0 + + Unless required by applicable law or agreed to in writing, software + distributed under the License is distributed on an "AS IS" BASIS, + WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + See the License for the specific language governing permissions and + limitations under the License. diff --git a/cuda_bindings_12/MANIFEST.in b/cuda_bindings_12/MANIFEST.in index d381e04d59e..30c4601806e 100644 --- a/cuda_bindings_12/MANIFEST.in +++ b/cuda_bindings_12/MANIFEST.in @@ -1,8 +1,11 @@ # SPDX-FileCopyrightText: Copyright (c) 2024-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: Apache-2.0 -recursive-include cuda/ *.pyx *.pxd *.pxi +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 exclude cuda/bindings/_bindings cuda?bindings?_bindings +# package-local shared PEP 517 helper +include _build_shared.py diff --git a/cuda_bindings_12/README.md b/cuda_bindings_12/README.md index a0657706d06..1b606ba9a75 100644 --- a/cuda_bindings_12/README.md +++ b/cuda_bindings_12/README.md @@ -1,10 +1,10 @@ # `cuda.bindings`: Low-level CUDA interfaces -`cuda.bindings` is a standard set of low-level interfaces, providing full coverage of and access to the CUDA host APIs from Python. Checkout the [Overview page](https://nvidia.github.io/cuda-python/cuda-bindings/latest/overview.html) for the workflow and performance results. +`cuda.bindings` is a standard set of low-level interfaces, providing full coverage of and access to the CUDA host APIs from Python. Checkout the [Overview page](https://nvidia.github.io/cuda-python/cuda-bindings/12.9.7/overview.html) for the workflow and performance results. ## Installing -Please refer to the [Installation page](https://nvidia.github.io/cuda-python/cuda-bindings/latest/install.html) for instructions and required/optional dependencies. +Please refer to the [Installation page](https://nvidia.github.io/cuda-python/cuda-bindings/12.9.7/install.html) for instructions and required/optional dependencies. ## Developing @@ -31,7 +31,7 @@ To run these tests: ### Cython Unit Tests -Cython tests are located in `tests/cython` and need to be built. These builds have the same CUDA Toolkit header requirements as [Installing from Source](https://nvidia.github.io/cuda-python/cuda-bindings/latest/install.html#requirements) where the major.minor version must match `cuda.bindings`. To build them: +Cython tests are located in `tests/cython` and need to be built. These builds have the same CUDA Toolkit header requirements as [Installing from Source](https://nvidia.github.io/cuda-python/cuda-bindings/12.9.7/install.html#requirements) where the major.minor version must match `cuda.bindings`. To build them: 1. Setup environment variable `CUDA_HOME` with the path to the CUDA Toolkit installation. 2. Run `build_tests` script located in `test/cython` appropriate to your platform. This will both cythonize the tests and build them. diff --git a/cuda_bindings_12/_build_shared.py b/cuda_bindings_12/_build_shared.py new file mode 120000 index 00000000000..e4ab5964749 --- /dev/null +++ b/cuda_bindings_12/_build_shared.py @@ -0,0 +1 @@ +../cuda_bindings_13/_build_shared.py \ No newline at end of file diff --git a/cuda_bindings_12/benchmarks/pytest.ini b/cuda_bindings_12/benchmarks/pytest.ini index 99da0054320..f57e8b1e1cc 100644 --- a/cuda_bindings_12/benchmarks/pytest.ini +++ b/cuda_bindings_12/benchmarks/pytest.ini @@ -4,3 +4,9 @@ [pytest] required_plugins = pytest-benchmark addopts = --benchmark-skip +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/build_hooks.py b/cuda_bindings_12/build_hooks.py index f09ee822a50..4cca4063475 100644 --- a/cuda_bindings_12/build_hooks.py +++ b/cuda_bindings_12/build_hooks.py @@ -12,15 +12,29 @@ import functools import glob import os +import re import shutil import sys import sysconfig import tempfile +from pathlib import Path from warnings import warn from setuptools import build_meta as _build_meta from setuptools.extension import Extension +import _build_shared +from _build_shared import ( + _abi_stamp_path, + _check_toolchain_available, + _cython_cache_path, + _resolve_toolchain_name, + _stable_cython_alias, + check_build_key, + record_build_key, + resolve_toolchain, +) + # Metadata hooks delegate directly to setuptools -- no CUDA needed. prepare_metadata_for_build_editable = _build_meta.prepare_metadata_for_build_editable prepare_metadata_for_build_wheel = _build_meta.prepare_metadata_for_build_wheel @@ -29,9 +43,29 @@ get_requires_for_build_wheel = _build_meta.get_requires_for_build_wheel get_requires_for_build_editable = _build_meta.get_requires_for_build_editable + +def __getattr__(name): + # setup.py reads ``build_hooks.force_build_ext``; the flag itself lives in + # _build_shared, where check_build_key() sets it. + if name == "force_build_ext": + return _build_shared.force_build_ext + raise AttributeError(f"module {__name__!r} has no attribute {name!r}") + + +# Note: There is no support guarantee for environment variables like +# CUDA_PYTHON_TOOLCHAIN, CUDA_PYTHON_CYTHON_CACHE_DIR, etc. They may be +# removed or changed in the future. + # Populated by _build_cuda_bindings(); consumed by setup.py. _extensions = None +# The generated sources declare the types and functions of one CUDA header set, +# and cydriver.pxd records which. The install docs state the build rule that results. +_CYDRIVER_PXD = Path(__file__).resolve().parent / "cuda" / "bindings" / "cydriver.pxd" +_GENERATED_VERSION_RE = re.compile(r"^cdef enum:\s*CUDA_VERSION\s*=\s*(\d+)\s*$") +_CUDA_H_VERSION_RE = re.compile(r"^#\s*define\s+CUDA_VERSION\s+(\d+)\s*$") +_INSTALL_URL = "https://nvidia.github.io/cuda-python/cuda-bindings/12.9.10/install.html#installing-from-source" + @functools.cache def _get_cuda_paths() -> list[str]: @@ -43,6 +77,130 @@ def _get_cuda_paths() -> list[str]: return CUDA_HOME +# ----------------------------------------------------------------------- +# CUDA header check + + +def _cuda_h_path(cuda_path: str) -> str: + """The cuda.h under cuda_path, with symlinks such as /usr/local/cuda resolved for messages.""" + return os.path.realpath(os.path.join(cuda_path, "include", "cuda.h")) + + +def _read_version_macro(path: str, pattern: re.Pattern) -> int | None: + """The integer on the first line of ``path`` that matches ``pattern``, or None if no line does.""" + with open(path, encoding="utf-8") as f: + for line in f: + m = pattern.match(line) + if m: + return int(m.group(1)) + return None + + +def _read_cuda_h_version(cuda_path: str) -> int: + """The CUDA_VERSION macro of the cuda.h under cuda_path, for example 13040 for 13.4.""" + cuda_h = _cuda_h_path(cuda_path) + try: + version = _read_version_macro(cuda_h, _CUDA_H_VERSION_RE) + except OSError: + version = None + if version is None: + raise RuntimeError( + f"Cannot read CUDA_VERSION from {cuda_h}. " + "Ensure CUDA_PATH or CUDA_HOME points to a CUDA Toolkit with include/cuda.h." + ) + return version + + +def _generated_cuda_version() -> int: + """The CUDA_VERSION of the headers that this source tree was generated from. + + Read from cuda/bindings/cydriver.pxd. + """ + version = _read_version_macro(str(_CYDRIVER_PXD), _GENERATED_VERSION_RE) + if version is None: + raise RuntimeError(f"Cannot read CUDA_VERSION from {_CYDRIVER_PXD}") + return version + + +def _major_minor(cuda_version: int) -> str: + """13040 -> \"13.4\".""" + return f"{cuda_version // 1000}.{cuda_version // 10 % 100}" + + +def _check_cuda_headers(cuda_path: str) -> None: + """Reject a toolkit whose cuda.h is not the major.minor that this source tree was generated from. + + Against another minor, the C++ compile fails with a long list of + redefinition and undeclared-type errors that do not name the cause. See + https://github.com/NVIDIA/cuda-python/issues/2783. Runs before cythonize, + which is the first step that touches the source tree. + """ + generated = _generated_cuda_version() + needed, found = _major_minor(generated), _major_minor(_read_cuda_h_version(cuda_path)) + if found != needed: + raise RuntimeError( + f"This cuda-bindings source tree needs CUDA {needed} headers, but {_cuda_h_path(cuda_path)} is " + f"CUDA {found}. This is a build-time requirement only: at run time this cuda-bindings build can be " + f"used with any CUDA {generated // 1000}.x toolkit, see {_INSTALL_URL}. Point CUDA_PATH or CUDA_HOME " + f"at a CUDA {needed} toolkit, or build from cuda-bindings {found}.x sources." + ) + + +def _tweak_flags(name, extra_compile_args, extra_link_args): + """cuda-bindings flags that do not belong in the shared set.""" + if name != "msvc": + # cudaMemcpy*Array* and cudaGetDriverEntryPoint are deprecated but still + # supported; suppress the resulting warnings so a future -Werror build + # is not broken by Cython-generated calls we cannot control. + extra_compile_args = [*extra_compile_args, "-Wno-deprecated-declarations"] + return extra_compile_args, extra_link_args + + +def _resolve_toolchain(debug=False, compile_for_coverage=False): + """Resolve the C/C++ toolchain from CUDA_PYTHON_TOOLCHAIN (cuda.bindings flags). + + See _build_shared.resolve_toolchain() for the return value and the + environment handling. What is specific to cuda.bindings is declared here. + """ + return resolve_toolchain( + # c++14: raising to c++17 costs a measured ~15% on launch_{256,512}_args + # from gcc's c++17 variadic-template expansion. + cxx_std=14, + debug=debug, + compile_for_coverage=compile_for_coverage, + tweak=_tweak_flags, + ) + + +# ----------------------------------------------------------------------- +# Toolchain stamp + +# Records the toolchain of the last completed build for this extension ABI, +# so setup.py can force build_ext when it changes. Written by +# record_build_toolchain(). +_BUILD_TOOLCHAIN_STAMP = _abi_stamp_path(".build-toolchain") + + +def _check_build_toolchain(toolchain): + """Set force_build_ext when the toolchain changed since the last build. + + Setuptools' freshness check does not include the extension flags, so a + stale .so compiled by a previous toolchain would otherwise be packaged. + """ + check_build_key(_BUILD_TOOLCHAIN_STAMP, toolchain, "Toolchain") + + +def record_build_toolchain() -> None: + """Stamp the toolchain of the build that just completed. + + setup.py calls this after build_ext succeeds, so that a build which failed + partway through does not claim outputs it never produced. Re-derives the + toolchain name from the environment rather than caching it in a global. + """ + name, *_ = _resolve_toolchain_name() + record_build_key(_BUILD_TOOLCHAIN_STAMP, name) + + # ----------------------------------------------------------------------- # Header parsing helpers (called only from _build_cuda_bindings) @@ -296,18 +454,24 @@ def _prep_extensions(sources, libraries, include_dirs, library_dirs, extra_compi # Main build function -def _build_cuda_bindings(strip=False): +def _build_cuda_bindings(debug=False): """Build all cuda-bindings extensions. All CUDA-dependent logic (header parsing, code generation, cythonization) is deferred to this function so that metadata queries do not require a CUDA toolkit installation. """ + import Cython from Cython.Build import cythonize + from Cython.Compiler import Options as _CythonOptions global _extensions cuda_paths = _get_cuda_paths() + # CUDA 12 supports split toolkit installations. Check the first cuda.h + # found in the compiler's include-search order, allowing library-only roots. + cuda_header_root = next((path for path in cuda_paths if os.path.isfile(_cuda_h_path(path))), cuda_paths[0]) + _check_cuda_headers(cuda_header_root) if os.environ.get("PARALLEL_LEVEL") is not None: warn( @@ -322,6 +486,14 @@ def _build_cuda_bindings(strip=False): parser_caching = bool(os.environ.get("CUDA_PYTHON_PARSER_CACHING", False)) compile_for_coverage = bool(int(os.environ.get("CUDA_PYTHON_COVERAGE", "0"))) + toolchain, _cc, _cxx, extra_compile_args, extra_link_args = _resolve_toolchain( + debug=debug, compile_for_coverage=compile_for_coverage + ) + _check_toolchain_available(toolchain) + extra_cythonize_kwargs = {} + if debug and sys.platform != "win32": + extra_cythonize_kwargs["gdb_debug"] = True + # Parse CUDA headers include_path_list = [os.path.join(path, "include") for path in cuda_paths] header_dict = _fetch_header_paths(_REQUIRED_HEADERS, include_path_list) @@ -370,33 +542,11 @@ def _build_cuda_bindings(strip=False): include_dirs = [ os.path.dirname(sysconfig.get_path("include")), ] + include_path_list - library_dirs = [sysconfig.get_path("platlib"), os.path.join(os.sys.prefix, "lib")] + # Add the build-environment library path in setup.py instead of generated Cython sources. + library_dirs = [] cudalib_subdirs = [r"lib\x64"] if sys.platform == "win32" else ["lib64", "lib"] library_dirs.extend(os.path.join(prefix, subdir) for prefix in cuda_paths for subdir in cudalib_subdirs) - extra_compile_args = [] - extra_link_args = [] - extra_cythonize_kwargs = {} - if sys.platform != "win32": - extra_compile_args += [ - "-std=c++14", - "-fpermissive", - "-Wno-deprecated-declarations", - "-fno-var-tracking-assignments", - ] - if "--debug" in sys.argv: - extra_cythonize_kwargs["gdb_debug"] = True - extra_compile_args += ["-g", "-O0"] - extra_compile_args += ["-D _GLIBCXX_ASSERTIONS"] - else: - extra_compile_args += ["-O3"] - if strip and sys.platform == "linux": - extra_link_args += ["-Wl,--strip-all"] - if compile_for_coverage: - # CYTHON_TRACE_NOGIL indicates to trace nogil functions. It is not - # related to free-threading builds. - extra_compile_args += ["-DCYTHON_TRACE_NOGIL=1", "-DCYTHON_USE_SYS_MONITORING=0"] - # Rename architecture-specific files dst_files = _rename_architecture_specific_files() @@ -433,28 +583,59 @@ def _cleanup_dst_files(): ) # Cythonize + _CythonOptions.warning_errors = True cython_directives = dict(language_level=3, embedsignature=True, binding=True, freethreading_compatible=True) if compile_for_coverage: cython_directives["linetrace"] = True - _extensions = cythonize( - extensions, - nthreads=nthreads, - build_dir="build/cython", + # Force a full rebuild when the toolchain changed since the last successful + # build, so a stale .so from a previous toolchain is never packaged. + _check_build_toolchain(toolchain) + + cache_path = _cython_cache_path( + "cuda-bindings-12", compiler_directives=cython_directives, - **extra_cythonize_kwargs, + language_level=3, + cplus=True, + debug=debug, ) + def _do_cythonize(cython_include_path): + global _extensions + _extensions = cythonize( + extensions, + nthreads=nthreads, + build_dir="." if compile_for_coverage else "build/cython", + compiler_directives=cython_directives, + include_path=cython_include_path, + cache=cache_path, + **extra_cythonize_kwargs, + ) + + if cache_path is not None: + # Alias Cython's bundled .pxd declarations under a stable worktree-relative + # path so Cython's cache fingerprint sees the same path on every run + # despite PEP 517 build environments landing under randomized temp prefixes. + stdlib_target = Path(Cython.__file__).parent / "Includes" + stdlib_alias = Path(__file__).parent / ".cython-stdlib" + with _stable_cython_alias(stdlib_target, stdlib_alias) as rel_stdlib: + _do_cythonize([".", rel_stdlib]) + else: + _do_cythonize(["."]) + # ----------------------------------------------------------------------- # PEP 517 build hooks def build_wheel(wheel_directory, config_settings=None, metadata_directory=None): - _build_cuda_bindings(strip=True) + debug = config_settings.get("debug", False) if config_settings else False + _build_cuda_bindings(debug=debug) return _build_meta.build_wheel(wheel_directory, config_settings, metadata_directory) def build_editable(wheel_directory, config_settings=None, metadata_directory=None): - _build_cuda_bindings(strip=False) + debug_default = sys.platform != "win32" # Debug builds not supported on Windows + debug = config_settings.get("debug", debug_default) if config_settings else debug_default + _build_cuda_bindings(debug=debug) return _build_meta.build_editable(wheel_directory, config_settings, metadata_directory) diff --git a/cuda_bindings_12/cuda/bindings/_internal/utils.pxd b/cuda_bindings_12/cuda/bindings/_internal/utils.pxd index dfbb85bd85c..f4a76de327e 100644 --- a/cuda_bindings_12/cuda/bindings/_internal/utils.pxd +++ b/cuda_bindings_12/cuda/bindings/_internal/utils.pxd @@ -24,7 +24,7 @@ cdef extern from * nogil: if (own_data) manager_.reset(data); else - raw_data_ = data; + raw_data_ = data; } nullable_unique_ptr(const nullable_unique_ptr&) = delete; @@ -39,7 +39,7 @@ cdef extern from * nogil: { manager_ = std::move(other.manager_); raw_data_ = nullptr; // just in case - } + } else { manager_.reset(nullptr); // just in case @@ -55,7 +55,7 @@ cdef extern from * nogil: { manager_ = std::move(other.manager_); raw_data_ = nullptr; // just in case - } + } else { manager_.reset(nullptr); // just in case diff --git a/cuda_bindings_12/cuda/bindings/_internal/utils.pyx b/cuda_bindings_12/cuda/bindings/_internal/utils.pyx index df17a9e47df..ce9012113df 100644 --- a/cuda_bindings_12/cuda/bindings/_internal/utils.pyx +++ b/cuda_bindings_12/cuda/bindings/_internal/utils.pyx @@ -98,7 +98,14 @@ cdef int get_nested_resource_ptr(nested_resource[ResT] &in_out_ptr, object obj, nested_ptr.reset(nested_vec, True) for i, obj_i in enumerate(obj): if ResT is char: - obj_i_bytes = ((obj_i)).encode() + obj_i_type = type(obj_i) + if obj_i_type is str: + obj_i_bytes = obj_i.encode("utf-8") + elif obj_i_type is bytes: + obj_i_bytes = obj_i + else: + raise TypeError( + f"Expected str or bytes, got {obj_i_type.__name__}") str_len = (len(obj_i_bytes)) + 1 # including null termination deref(nested_res_vec)[i].resize(str_len) obj_i_ptr = (obj_i_bytes) diff --git a/cuda_bindings_12/cuda/bindings/_lib/utils.pxd b/cuda_bindings_12/cuda/bindings/_lib/utils.pxd index 706c6d3fc7a..c36ddfb0d12 100644 --- a/cuda_bindings_12/cuda/bindings/_lib/utils.pxd +++ b/cuda_bindings_12/cuda/bindings/_lib/utils.pxd @@ -152,6 +152,7 @@ cdef class _HelperCUcoredumpSettings: cdef cydriver.CUcoredumpSettings_enum _attrib cdef bint _is_getter cdef size_t _size + cdef object _references # keeps caller bytes alive so _charstar stays valid # Return values cdef cpp_bool _bool diff --git a/cuda_bindings_12/cuda/bindings/_lib/utils.pxi b/cuda_bindings_12/cuda/bindings/_lib/utils.pxi index c11d0d0c197..eebdd17b41e 100644 --- a/cuda_bindings_12/cuda/bindings/_lib/utils.pxi +++ b/cuda_bindings_12/cuda/bindings/_lib/utils.pxi @@ -734,6 +734,8 @@ cdef class _HelperCUcoredumpSettings: self._cptr = self._charstar self._size = 1024 else: + # Keep a reference so the borrowed _charstar buffer stays alive. + self._references = init_value self._charstar = init_value self._cptr = self._charstar self._size = len(init_value) @@ -756,7 +758,9 @@ cdef class _HelperCUcoredumpSettings: raise TypeError('Unsupported attribute: {}'.format(attr.name)) def __dealloc__(self): - pass + # Only the getter owns heap storage; the setter borrows caller bytes. + if self._is_getter: + free(self._charstar) @property def cptr(self): diff --git a/cuda_bindings_12/cuda/bindings/_lib/windll.pxd b/cuda_bindings_12/cuda/bindings/_lib/windll.pxd index 294a1a9fd90..b5fd5c4db90 100644 --- a/cuda_bindings_12/cuda/bindings/_lib/windll.pxd +++ b/cuda_bindings_12/cuda/bindings/_lib/windll.pxd @@ -14,7 +14,7 @@ cdef extern from "windows.h" nogil: ctypedef const char *LPCSTR ctypedef int BOOL - cdef DWORD LOAD_LIBRARY_SEARCH_SYSTEM32 = 0x00000800 + const DWORD LOAD_LIBRARY_SEARCH_SYSTEM32 HMODULE _LoadLibraryExW "LoadLibraryExW"( LPCWSTR lpLibFileName, diff --git a/cuda_bindings_12/cuda/bindings/utils/__init__.py b/cuda_bindings_12/cuda/bindings/utils/__init__.py index 62d083dc7c6..d6136da14b0 100644 --- a/cuda_bindings_12/cuda/bindings/utils/__init__.py +++ b/cuda_bindings_12/cuda/bindings/utils/__init__.py @@ -26,6 +26,9 @@ def get_cuda_native_handle(obj: Any) -> int: """ obj_type = type(obj) try: - return _handle_getters[obj_type](obj) + getter = _handle_getters[obj_type] except KeyError: raise TypeError("Unknown type: " + str(obj_type)) from None + # Deliberately outside the try: a KeyError raised by the getter itself is a + # bug in that getter, not an unregistered type. + return getter(obj) diff --git a/cuda_bindings_12/docs/README.md b/cuda_bindings_12/docs/README.md index a5e65842a8e..8b445f41e65 100644 --- a/cuda_bindings_12/docs/README.md +++ b/cuda_bindings_12/docs/README.md @@ -1,11 +1,10 @@ # Build the documentation 1. Install the `cuda-bindings` package of the version that we need to document. -2. Ensure the version is included in the [`versions.json`](./versions.json). +2. Ensure the version is included in the canonical + [`nv-versions.json`](../../cuda_bindings_13/docs/nv-versions.json). 3. Build the docs with `./build_docs.sh`. 4. The html artifacts should be available under both `./build/html/latest` and `./build/html/`. -Alternatively, we can build all the docs at once by running [`cuda_python/docs/build_all_docs.sh`](../../cuda_python/docs/build_all_docs.sh). - To publish the docs with the built version, it is important to note that the html files of older versions -should be kept intact, in order for the version selection (through `versions.json`) to work. +should be kept intact, in order for the version selection (through `nv-versions.json`) to work. diff --git a/cuda_bindings_12/docs/build_docs.sh b/cuda_bindings_12/docs/build_docs.sh index 826e9fbb6fd..351b1a27c45 100755 --- a/cuda_bindings_12/docs/build_docs.sh +++ b/cuda_bindings_12/docs/build_docs.sh @@ -16,13 +16,12 @@ fi # SPHINX_CUDA_BINDINGS_VER is used to create a subdir under build/html # (the Makefile file for sphinx-build also honors it if defined). -# If there's a post release (ex: .post1) we don't want it to show up in the -# version selector or directory structure. +# Post releases reuse their base release directory, while prerelease and +# development releases remain distinct so they cannot overwrite stable docs. if [[ -z "${SPHINX_CUDA_BINDINGS_VER}" ]]; then export SPHINX_CUDA_BINDINGS_VER=$(python -c "from importlib.metadata import version; \ - ver = '.'.join(str(version('cuda-bindings')).split('.')[:3]); \ - print(ver)" \ - | awk -F'+' '{print $1}') + ver = str(version('cuda-bindings')).split('+', 1)[0]; \ + print(ver if '.dev' in ver else ver.split('.post', 1)[0])") fi # build the docs (in parallel) @@ -33,8 +32,9 @@ SPHINXOPTS="-j 4 -d build/.doctrees" make html # obsecure Sphinx errors #SPHINXOPTS="-v" make html -# to support version dropdown menu -cp ./versions.json build/html +# Keep the CUDA 12.9 release pages on the monorepo's canonical bindings +# selector so publishing this line cannot remove newer CUDA 13 entries. +cp ../../cuda_bindings_13/docs/nv-versions.json build/html/nv-versions.json # to have a redirection page (to the latest docs) cp source/_templates/main.html build/html/index.html diff --git a/cuda_bindings_12/docs/source/_static/javascripts/version_dropdown.js b/cuda_bindings_12/docs/source/_static/javascripts/version_dropdown.js index aa0ce2bdc67..172b25f5729 100644 --- a/cuda_bindings_12/docs/source/_static/javascripts/version_dropdown.js +++ b/cuda_bindings_12/docs/source/_static/javascripts/version_dropdown.js @@ -1,22 +1,7 @@ // SPDX-FileCopyrightText: Copyright (c) 2024-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. // SPDX-License-Identifier: Apache-2.0 -function change_current_version(event) { - event.preventDefault(); - - var selectedVersion = event.target.textContent; - var currentVersion = document.getElementById('currentVersion'); - - // need to update both the on-screen state and the internal (persistent) storage - currentVersion.textContent = selectedVersion; - sessionStorage.setItem("currentVersion", selectedVersion); - - // Navigate to the clicked URL - window.location.href = event.target.href; -} - - -function add_version_dropdown(jsonLoc, targetLoc, currentVersion) { +function add_version_dropdown(jsonLoc, currentVersion) { var otherVersionsDiv = document.getElementById('otherVersions'); fetch(jsonLoc) @@ -26,17 +11,18 @@ function add_version_dropdown(jsonLoc, targetLoc, currentVersion) { .then(function(data) { var versions = data; - if (Object.keys(versions).length >= 1) { + if (Array.isArray(versions) && versions.length >= 1) { var dlElement = document.createElement('dl'); var dtElement = document.createElement('dt'); dtElement.textContent = 'Versions'; dlElement.appendChild(dtElement); - for (var ver in versions) { - var url = versions[ver]; + for (var index = 0; index < versions.length; index++) { + var ver = versions[index].version; + var url = versions[index].url; var ddElement = document.createElement('dd'); var aElement = document.createElement('a'); - aElement.setAttribute('href', targetLoc + url); + aElement.setAttribute('href', url); aElement.textContent = ver; if (ver === currentVersion) { @@ -46,8 +32,6 @@ function add_version_dropdown(jsonLoc, targetLoc, currentVersion) { } ddElement.appendChild(aElement); - // Attach event listeners to version links - ddElement.addEventListener('click', change_current_version); dlElement.appendChild(ddElement); } @@ -56,6 +40,6 @@ function add_version_dropdown(jsonLoc, targetLoc, currentVersion) { } }) .catch(function(error) { - console.error('Error fetching version.json:', error); + console.error('Error fetching nv-versions.json:', error); }); } diff --git a/cuda_bindings_12/docs/source/_templates/sidebar/variant-selector.html b/cuda_bindings_12/docs/source/_templates/sidebar/variant-selector.html index b041194c501..92539a95163 100644 --- a/cuda_bindings_12/docs/source/_templates/sidebar/variant-selector.html +++ b/cuda_bindings_12/docs/source/_templates/sidebar/variant-selector.html @@ -1,7 +1,7 @@
cuda-bindings - v: {{ version }} + v: {{ release }}
@@ -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..714180e0d89 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.14 * 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..bc92e60ce95 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.14 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/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/nvml/test_cuda.py b/cuda_bindings_12/tests/nvml/test_cuda.py index d05e2634a12..aaab23e55fe 100644 --- a/cuda_bindings_12/tests/nvml/test_cuda.py +++ b/cuda_bindings_12/tests/nvml/test_cuda.py @@ -1,6 +1,8 @@ # SPDX-FileCopyrightText: Copyright (c) 2025-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: Apache-2.0 +import pytest + import cuda.bindings.driver as cuda from cuda.bindings import nvml @@ -15,10 +17,15 @@ def get_nvml_device_names(): for idx in range(num_devices): handle = nvml.device_get_handle_by_index_v2(idx) name = nvml.device_get_name(handle) - info = nvml.device_get_pci_info_v3(handle) - assert isinstance(info.bus, int) + try: + info = nvml.device_get_pci_info_v3(handle) + except nvml.NotSupportedError: + bus_id = -1 + else: + bus_id = info.bus + assert isinstance(bus_id, int) assert isinstance(name, str) - result.append({"name": name, "id": info.bus}) + result.append({"name": name, "id": bus_id}) return result @@ -54,4 +61,15 @@ def test_cuda_device_order(): cuda_devices = get_cuda_device_names() nvml_devices = get_nvml_device_names() - assert cuda_devices == nvml_devices, "CUDA and NVML device lists do not match" + for kind in ("Orin", "Thor"): + if any(kind in device["name"] for device in nvml_devices): + pytest.skip(f"Skipping test on {kind}, which has non-standard device naming") + + def compare(cuda_device, nvml_device): + return cuda_device["name"] == nvml_device["name"] and ( + nvml_device["id"] == -1 or cuda_device["id"] == nvml_device["id"] + ) + + assert len(cuda_devices) <= len(nvml_devices) + for cuda_device in cuda_devices: + assert any(compare(cuda_device, nvml_device) for nvml_device in nvml_devices) diff --git a/cuda_bindings_12/tests/nvml/test_init.py b/cuda_bindings_12/tests/nvml/test_init.py index 94b489ab45f..bcda0591bb6 100644 --- a/cuda_bindings_12/tests/nvml/test_init.py +++ b/cuda_bindings_12/tests/nvml/test_init.py @@ -6,7 +6,15 @@ import pytest -from cuda.bindings import nvml +from cuda.bindings import driver, nvml + + +def _driver_version_at_least_13_4(): + (status,) = driver.cuInit(0) + assert status == driver.CUresult.CUDA_SUCCESS + status, version = driver.cuDriverGetVersion() + assert status == driver.CUresult.CUDA_SUCCESS + return version >= 13040 def assert_nvml_is_initialized(): @@ -36,10 +44,14 @@ def test_devices_are_the_same_architecture(all_devices): @pytest.mark.skipif(sys.platform == "win32", reason="Test not supported on Windows") +@pytest.mark.thread_unsafe(reason="nvml init affects other threads") def test_init_ref_count(): """ Verifies that we can call NVML shutdown and init(2) multiple times, and that ref counting works """ + if _driver_version_at_least_13_4(): + pytest.skip("Init behavior changed in CUDA 13.4") + with pytest.raises(nvml.UninitializedError): nvml.shutdown() diff --git a/cuda_bindings_12/tests/nvml/test_pynvml.py b/cuda_bindings_12/tests/nvml/test_pynvml.py index a96b2c87575..e3ea533dd83 100644 --- a/cuda_bindings_12/tests/nvml/test_pynvml.py +++ b/cuda_bindings_12/tests/nvml/test_pynvml.py @@ -4,8 +4,6 @@ # A set of tests ported from https://github.com/gpuopenanalytics/pynvml/blob/11.5.3/pynvml/tests/test_nvml.py import os -import time - import pytest from cuda.bindings import nvml @@ -147,22 +145,6 @@ def test_device_get_power_usage(ngpus, handles): assert power_mwatts >= 0.0 -def test_device_get_total_energy_consumption(ngpus, handles): - for i in range(ngpus): - with unsupported_before(handles[i], nvml.DeviceArch.VOLTA): - energy_mjoules1 = nvml.device_get_total_energy_consumption(handles[i]) - - for j in range(10): # idle for 150 ms - time.sleep(0.015) # and check for increase every 15 ms - with unsupported_before(handles[i], nvml.DeviceArch.VOLTA): - energy_mjoules2 = nvml.device_get_total_energy_consumption(handles[i]) - assert energy_mjoules2 >= energy_mjoules1 - if energy_mjoules2 > energy_mjoules1: - break - else: - raise AssertionError("energy did not increase across 150 ms interval") - - # [Skipping] pynvml.nvmlDeviceGetGpuOperationMode # [Skipping] pynvml.nvmlDeviceGetCurrentGpuOperationMode # [Skipping] pynvml.nvmlDeviceGetPendingGpuOperationMode diff --git a/cuda_bindings_12/tests/nvml/util.py b/cuda_bindings_12/tests/nvml/util.py index 038fe58d8be..06cb1c2277a 100644 --- a/cuda_bindings_12/tests/nvml/util.py +++ b/cuda_bindings_12/tests/nvml/util.py @@ -43,5 +43,5 @@ def supports_ecc(device): def supports_nvlink(device): fields = nvml.FieldValue(1) - fields[0].field_id = nvml.FI.DEV_NVLINK_GET_STATE + fields[0].field_id = nvml.FieldId.DEV_NVLINK_GET_STATE return nvml.device_get_field_values(device, fields)[0].nvml_return == nvml.Return.SUCCESS 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..b2217832592 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) diff --git a/cuda_bindings_12/tests/test_cudart.py b/cuda_bindings_12/tests/test_cudart.py index 9a496925ca4..7d87bfa17b1 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(): 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_nvjitlink.py b/cuda_bindings_12/tests/test_nvjitlink.py index 19dc040e47d..bec58b7a048 100644 --- a/cuda_bindings_12/tests/test_nvjitlink.py +++ b/cuda_bindings_12/tests/test_nvjitlink.py @@ -100,6 +100,13 @@ def test_create_and_destroy(option): nvjitlink.destroy(handle) +@pytest.mark.human_authored +def test_create_and_destroy_bytes_options(): + handle = nvjitlink.create(1, [b"-arch=sm_80"]) + assert handle != 0 + nvjitlink.destroy(handle) + + @pytest.mark.parametrize("option", ARCHITECTURES) def test_complete_empty(option): handle = nvjitlink.create(1, [f"-arch={option}"]) diff --git a/cuda_bindings_12/tests/test_nvvm.py b/cuda_bindings_12/tests/test_nvvm.py index bd6f49f03f1..3cea5a0f817 100644 --- a/cuda_bindings_12/tests/test_nvvm.py +++ b/cuda_bindings_12/tests/test_nvvm.py @@ -246,7 +246,7 @@ def test_get_buffer_empty(get_size, get_buffer): assert buffer == b"\x00" -@pytest.mark.parametrize("options", [[], ["-opt=0"], ["-opt=3", "-g"]]) +@pytest.mark.parametrize("options", [[], ["-opt=0"], ["-opt=3", "-g"], [b"-opt=0"]]) def test_compile_program_with_minimal_nvvm_ir(minimal_nvvmir, options): with nvvm_program() as prog: nvvm.add_module_to_program(prog, minimal_nvvmir, len(minimal_nvvmir), "FileNameHere.ll") diff --git a/cuda_bindings_12/tests/test_utils.py b/cuda_bindings_12/tests/test_utils.py index c9afe1f415b..33487c7f529 100644 --- a/cuda_bindings_12/tests/test_utils.py +++ b/cuda_bindings_12/tests/test_utils.py @@ -92,6 +92,27 @@ def test_get_handle_error(target): handle = get_cuda_native_handle(target) +@pytest.mark.agent_authored(model="claude-opus-5") +def test_get_handle_does_not_report_a_registered_type_as_unknown(monkeypatch): + """A KeyError from inside a handle getter is a bug in that getter. + + Reporting it as "Unknown type" is wrong twice over: the type *is* + registered, and `from None` hides the traceback that would say otherwise. + """ + from cuda.bindings.utils import _handle_getters + + class Registered: + pass + + def getter(_obj): + raise KeyError("lookup inside the getter failed") + + monkeypatch.setitem(_handle_getters, Registered, getter) + + with pytest.raises(KeyError, match="lookup inside the getter failed"): + get_cuda_native_handle(Registered()) + + @pytest.mark.parametrize( "module", [ @@ -112,6 +133,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 ( From d8815d7e43dfd2699e2cc1e09058ac91f820e9ae Mon Sep 17 00:00:00 2001 From: "Ralf W. Grosse-Kunstleve" Date: Tue, 20 Oct 2026 22:19:39 -0700 Subject: [PATCH 02/10] test(bindings): isolate CUDA 12 graph capture by thread Signed-off-by: Ralf W. Grosse-Kunstleve --- cuda_bindings_12/tests/test_cuda.py | 5 ++++- cuda_bindings_12/tests/test_cudart.py | 10 ++++++++-- 2 files changed, 12 insertions(+), 3 deletions(-) diff --git a/cuda_bindings_12/tests/test_cuda.py b/cuda_bindings_12/tests/test_cuda.py index b2217832592..78dd514b773 100644 --- a/cuda_bindings_12/tests/test_cuda.py +++ b/cuda_bindings_12/tests/test_cuda.py @@ -804,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 7d87bfa17b1..2c6dc008aff 100644 --- a/cuda_bindings_12/tests/test_cudart.py +++ b/cuda_bindings_12/tests/test_cudart.py @@ -521,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 @@ -1374,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) From a2983531e9a876ef7e674ad1ab8f32f242e3d16f Mon Sep 17 00:00:00 2001 From: "Ralf W. Grosse-Kunstleve" Date: Tue, 20 Oct 2026 22:19:44 -0700 Subject: [PATCH 03/10] test: classify CUDA 12 bindings alongside CUDA 13 Signed-off-by: Ralf W. Grosse-Kunstleve --- .../cuda_python_test_helpers/_pytest_plugin.py | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) 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) From 60bc60a38f01d4de913c4fb581ee49e4bacd9da1 Mon Sep 17 00:00:00 2001 From: "Ralf W. Grosse-Kunstleve" Date: Wed, 21 Oct 2026 06:56:29 -0700 Subject: [PATCH 04/10] Restore the security scan for the integrated CUDA 12 sources Signed-off-by: Ralf W. Grosse-Kunstleve --- .github/workflows/bandit.yml | 4 +--- 1 file changed, 1 insertion(+), 3 deletions(-) diff --git a/.github/workflows/bandit.yml b/.github/workflows/bandit.yml index 926619a799a..e414a2d7d3e 100644 --- a/.github/workflows/bandit.yml +++ b/.github/workflows/bandit.yml @@ -42,10 +42,8 @@ jobs: echo "codes=$(uvx toml2json ./ruff.toml | jq -r '.lint.ignore | map(select(test("^S\\d+"))) | join(",")')" >> "$GITHUB_OUTPUT" - name: Perform Bandit Analysis using Ruff uses: astral-sh/ruff-action@278981a28ce3188b1e39527901f38254bf3aac89 # v4.1.0 - # Preserve the inactive CUDA 12 snapshot; the source-integration layer - # adapts its security annotations and removes this temporary exclusion. with: - args: "check --select S --ignore ${{ steps.ignore-codes.outputs.codes }} --extend-exclude cuda_bindings_12 --output-format sarif --output-file results.sarif" + args: "check --select S --ignore ${{ steps.ignore-codes.outputs.codes }} --output-format sarif --output-file results.sarif" - name: Upload SARIF file uses: github/codeql-action/upload-sarif@24c54180a607b1449ed407dd24f251e4e9147c8d # v4.38.3 with: From 3107da317da09938ad64da996dda085b4bd06159 Mon Sep 17 00:00:00 2001 From: "Ralf W. Grosse-Kunstleve" Date: Wed, 21 Oct 2026 09:43:29 -0700 Subject: [PATCH 05/10] Remove superseded CUDA 12 package benchmarks Signed-off-by: Ralf W. Grosse-Kunstleve --- cuda_bindings_12/README.md | 7 +- cuda_bindings_12/benchmarks/conftest.py | 91 ----- cuda_bindings_12/benchmarks/kernels.py | 159 --------- cuda_bindings_12/benchmarks/pytest.ini | 12 - cuda_bindings_12/benchmarks/test_cupy.py | 199 ----------- .../benchmarks/test_launch_latency.py | 336 ------------------ cuda_bindings_12/benchmarks/test_numba.py | 52 --- .../benchmarks/test_pointer_attributes.py | 112 ------ cuda_bindings_12/pyproject.toml | 6 - 9 files changed, 2 insertions(+), 972 deletions(-) delete mode 100644 cuda_bindings_12/benchmarks/conftest.py delete mode 100644 cuda_bindings_12/benchmarks/kernels.py delete mode 100644 cuda_bindings_12/benchmarks/pytest.ini delete mode 100644 cuda_bindings_12/benchmarks/test_cupy.py delete mode 100755 cuda_bindings_12/benchmarks/test_launch_latency.py delete mode 100644 cuda_bindings_12/benchmarks/test_numba.py delete mode 100644 cuda_bindings_12/benchmarks/test_pointer_attributes.py diff --git a/cuda_bindings_12/README.md b/cuda_bindings_12/README.md index 1b606ba9a75..49fc1ca6470 100644 --- a/cuda_bindings_12/README.md +++ b/cuda_bindings_12/README.md @@ -60,8 +60,5 @@ To run these samples: ### Benchmark -Allows for analyzing binding performance using plugin [pytest-benchmark](https://github.com/ionelmc/pytest-benchmark). - -To run these benchmarks: -* `python -m pytest --benchmark-only benchmarks/` against editable installations -* `pytest --benchmark-only benchmarks/` against installed packages +See the shared [cuda.bindings benchmark harness](../benchmarks/cuda_bindings/README.md) +for performance benchmarks and usage instructions. diff --git a/cuda_bindings_12/benchmarks/conftest.py b/cuda_bindings_12/benchmarks/conftest.py deleted file mode 100644 index 0719dca47b8..00000000000 --- a/cuda_bindings_12/benchmarks/conftest.py +++ /dev/null @@ -1,91 +0,0 @@ -# SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. -# SPDX-License-Identifier: Apache-2.0 - -import numpy as np -import pytest - -from cuda import cuda, cudart, nvrtc - - -def ASSERT_DRV(err): - if isinstance(err, cuda.CUresult): - if err != cuda.CUresult.CUDA_SUCCESS: - raise RuntimeError(f"Cuda Error: {err}") - elif isinstance(err, cudart.cudaError_t): - if err != cudart.cudaError_t.cudaSuccess: - raise RuntimeError(f"Cudart Error: {err}") - elif isinstance(err, nvrtc.nvrtcResult): - if err != nvrtc.nvrtcResult.NVRTC_SUCCESS: - raise RuntimeError(f"Nvrtc Error: {err}") - else: - raise RuntimeError(f"Unknown error type: {err}") - - -@pytest.fixture(scope="function") -def init_cuda(): - # Initialize - (err,) = cuda.cuInit(0) - ASSERT_DRV(err) - err, device = cuda.cuDeviceGet(0) - ASSERT_DRV(err) - err, ctx = cuda.cuCtxCreate(0, device) - ASSERT_DRV(err) - - # create stream - err, stream = cuda.cuStreamCreate(cuda.CUstream_flags.CU_STREAM_NON_BLOCKING.value) - ASSERT_DRV(err) - - yield device, ctx, stream - - (err,) = cuda.cuStreamDestroy(stream) - ASSERT_DRV(err) - (err,) = cuda.cuCtxDestroy(ctx) - ASSERT_DRV(err) - - -@pytest.fixture(scope="function") -def load_module(): - module = None - - def _load_module(kernel_string, device): - nonlocal module - # Get module - err, major = cuda.cuDeviceGetAttribute( - cuda.CUdevice_attribute.CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, device - ) - ASSERT_DRV(err) - err, minor = cuda.cuDeviceGetAttribute( - cuda.CUdevice_attribute.CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, device - ) - ASSERT_DRV(err) - - err, prog = nvrtc.nvrtcCreateProgram(str.encode(kernel_string), b"kernelString.cu", 0, [], []) - ASSERT_DRV(err) - opts = [b"--fmad=false", bytes("--gpu-architecture=sm_" + str(major) + str(minor), "ascii")] - (err,) = nvrtc.nvrtcCompileProgram(prog, 2, opts) - - err_log, logSize = nvrtc.nvrtcGetProgramLogSize(prog) - ASSERT_DRV(err_log) - log = b" " * logSize - (err_log,) = nvrtc.nvrtcGetProgramLog(prog, log) - ASSERT_DRV(err_log) - result = log.decode() - if len(result) > 1: - print(result) - - ASSERT_DRV(err) - err, cubinSize = nvrtc.nvrtcGetCUBINSize(prog) - ASSERT_DRV(err) - cubin = b" " * cubinSize - (err,) = nvrtc.nvrtcGetCUBIN(prog, cubin) - ASSERT_DRV(err) - cubin = np.char.array(cubin) - err, module = cuda.cuModuleLoadData(cubin) - ASSERT_DRV(err) - - return module - - yield _load_module - - (err,) = cuda.cuModuleUnload(module) - ASSERT_DRV(err) diff --git a/cuda_bindings_12/benchmarks/kernels.py b/cuda_bindings_12/benchmarks/kernels.py deleted file mode 100644 index 89f1e1a0a87..00000000000 --- a/cuda_bindings_12/benchmarks/kernels.py +++ /dev/null @@ -1,159 +0,0 @@ -# SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. -# SPDX-License-Identifier: Apache-2.0 - -kernel_string = """\ -#define ITEM_PARAM(x, T) T x -#define REP1(x, T) , ITEM_PARAM(x, T) -#define REP2(x, T) REP1(x##0, T) REP1(x##1, T) -#define REP4(x, T) REP2(x##0, T) REP2(x##1, T) -#define REP8(x, T) REP4(x##0, T) REP4(x##1, T) -#define REP16(x, T) REP8(x##0, T) REP8(x##1, T) -#define REP32(x, T) REP16(x##0, T) REP16(x##1, T) -#define REP64(x, T) REP32(x##0, T) REP32(x##1, T) -#define REP128(x, T) REP64(x##0, T) REP64(x##1, T) -#define REP256(x, T) REP128(x##0, T) REP128(x##1, T) - -template -struct KernelFunctionParam -{ - unsigned char p[maxBytes]; -}; - -extern "C" __global__ void small_kernel(float *f) -{ - *f = 0.0f; -} - -extern "C" __global__ void empty_kernel() -{ - return; -} - -extern "C" __global__ -void small_kernel_512_args( - ITEM_PARAM(F, int*) - REP1(A, int*) - REP2(A, int*) - REP4(A, int*) - REP8(A, int*) - REP16(A, int*) - REP32(A, int*) - REP64(A, int*) - REP128(A, int*) - REP256(A, int*)) -{ - *F = 0; -} - -extern "C" __global__ -void small_kernel_512_bools( - ITEM_PARAM(F, bool) - REP1(A, bool) - REP2(A, bool) - REP4(A, bool) - REP8(A, bool) - REP16(A, bool) - REP32(A, bool) - REP64(A, bool) - REP128(A, bool) - REP256(A, bool)) -{ - return; -} - -extern "C" __global__ -void small_kernel_512_ints( - ITEM_PARAM(F, int) - REP1(A, int) - REP2(A, int) - REP4(A, int) - REP8(A, int) - REP16(A, int) - REP32(A, int) - REP64(A, int) - REP128(A, int) - REP256(A, int)) -{ - return; -} - -extern "C" __global__ -void small_kernel_512_doubles( - ITEM_PARAM(F, double) - REP1(A, double) - REP2(A, double) - REP4(A, double) - REP8(A, double) - REP16(A, double) - REP32(A, double) - REP64(A, double) - REP128(A, double) - REP256(A, double)) -{ - return; -} - -extern "C" __global__ -void small_kernel_512_chars( - ITEM_PARAM(F, char) - REP1(A, char) - REP2(A, char) - REP4(A, char) - REP8(A, char) - REP16(A, char) - REP32(A, char) - REP64(A, char) - REP128(A, char) - REP256(A, char)) -{ - return; -} - -extern "C" __global__ -void small_kernel_512_longlongs( - ITEM_PARAM(F, long long) - REP1(A, long long) - REP2(A, long long) - REP4(A, long long) - REP8(A, long long) - REP16(A, long long) - REP32(A, long long) - REP64(A, long long) - REP128(A, long long) - REP256(A, long long)) -{ - return; -} - -extern "C" __global__ -void small_kernel_256_args( - ITEM_PARAM(F, int*) - REP1(A, int*) - REP2(A, int*) - REP4(A, int*) - REP8(A, int*) - REP16(A, int*) - REP32(A, int*) - REP64(A, int*) - REP128(A, int*)) -{ - *F = 0; -} - -extern "C" __global__ -void small_kernel_16_args( - ITEM_PARAM(F, int*) - REP1(A, int*) - REP2(A, int*) - REP4(A, int*) - REP8(A, int*)) -{ - *F = 0; -} - -extern "C" __global__ void small_kernel_2048B(KernelFunctionParam<2048> param) -{ - // Do not touch param to prevent compiler from copying - // the whole structure from const bank to lmem. -} -""" diff --git a/cuda_bindings_12/benchmarks/pytest.ini b/cuda_bindings_12/benchmarks/pytest.ini deleted file mode 100644 index f57e8b1e1cc..00000000000 --- a/cuda_bindings_12/benchmarks/pytest.ini +++ /dev/null @@ -1,12 +0,0 @@ -# SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. -# SPDX-License-Identifier: Apache-2.0 - -[pytest] -required_plugins = pytest-benchmark -addopts = --benchmark-skip -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/benchmarks/test_cupy.py b/cuda_bindings_12/benchmarks/test_cupy.py deleted file mode 100644 index cd6c6740350..00000000000 --- a/cuda_bindings_12/benchmarks/test_cupy.py +++ /dev/null @@ -1,199 +0,0 @@ -# SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. -# SPDX-License-Identifier: Apache-2.0 - -import ctypes - -import pytest - -try: - import cupy - - skip_tests = False -except ImportError: - skip_tests = True - -from kernels import kernel_string - - -def launch(kernel, args=()): - kernel((1,), (1,), args) - - -# Measure launch latency with no parmaeters -@pytest.mark.skipif(skip_tests, reason="cupy is not installed") -@pytest.mark.benchmark(group="cupy") -def test_launch_latency_empty_kernel(benchmark): - module = cupy.RawModule(code=kernel_string) - kernel = module.get_function("empty_kernel") - - stream = cupy.cuda.stream.Stream(non_blocking=True) - - with stream: - benchmark(launch, kernel) - stream.synchronize() - - -# Measure launch latency with a single parameter -@pytest.mark.skipif(skip_tests, reason="cupy is not installed") -@pytest.mark.benchmark(group="cupy") -def test_launch_latency_small_kernel(benchmark): - module = cupy.RawModule(code=kernel_string) - kernel = module.get_function("small_kernel") - cupy.cuda.set_allocator() - arg = cupy.cuda.alloc(ctypes.sizeof(ctypes.c_float)) - - stream = cupy.cuda.stream.Stream(non_blocking=True) - - with stream: - benchmark(launch, kernel, (arg,)) - stream.synchronize() - - -# Measure launch latency with many parameters using builtin parameter packing -@pytest.mark.skipif(skip_tests, reason="cupy is not installed") -@pytest.mark.benchmark(group="cupy") -def test_launch_latency_small_kernel_512_args(benchmark): - module = cupy.RawModule(code=kernel_string) - kernel = module.get_function("small_kernel_512_args") - cupy.cuda.set_allocator() - - args = [] - for _ in range(512): - args.append(cupy.cuda.alloc(ctypes.sizeof(ctypes.c_int))) - args = tuple(args) - - stream = cupy.cuda.stream.Stream(non_blocking=True) - - with stream: - benchmark(launch, kernel, args) - stream.synchronize() - - -# Measure launch latency with many parameters using builtin parameter packing -@pytest.mark.skipif(skip_tests, reason="cupy is not installed") -@pytest.mark.benchmark(group="cupy") -def test_launch_latency_small_kernel_512_bools(benchmark): - module = cupy.RawModule(code=kernel_string) - kernel = module.get_function("small_kernel_512_bools") - cupy.cuda.set_allocator() - - args = [True] * 512 - args = tuple(args) - - stream = cupy.cuda.stream.Stream(non_blocking=True) - - with stream: - benchmark(launch, kernel, args) - stream.synchronize() - - -# Measure launch latency with many parameters using builtin parameter packing -@pytest.mark.skipif(skip_tests, reason="cupy is not installed") -@pytest.mark.benchmark(group="cupy") -def test_launch_latency_small_kernel_512_doubles(benchmark): - module = cupy.RawModule(code=kernel_string) - kernel = module.get_function("small_kernel_512_doubles") - cupy.cuda.set_allocator() - - args = [1.2345] * 512 - args = tuple(args) - - stream = cupy.cuda.stream.Stream(non_blocking=True) - - with stream: - benchmark(launch, kernel, args) - stream.synchronize() - - -# Measure launch latency with many parameters using builtin parameter packing -@pytest.mark.skipif(skip_tests, reason="cupy is not installed") -@pytest.mark.benchmark(group="cupy") -def test_launch_latency_small_kernel_512_ints(benchmark): - module = cupy.RawModule(code=kernel_string) - kernel = module.get_function("small_kernel_512_ints") - cupy.cuda.set_allocator() - - args = [123] * 512 - args = tuple(args) - - stream = cupy.cuda.stream.Stream(non_blocking=True) - - with stream: - benchmark(launch, kernel, args) - stream.synchronize() - - -# Measure launch latency with many parameters using builtin parameter packing -@pytest.mark.skipif(skip_tests, reason="cupy is not installed") -@pytest.mark.benchmark(group="cupy") -def test_launch_latency_small_kernel_512_bytes(benchmark): - module = cupy.RawModule(code=kernel_string) - kernel = module.get_function("small_kernel_512_chars") - cupy.cuda.set_allocator() - - args = [127] * 512 - args = tuple(args) - - stream = cupy.cuda.stream.Stream(non_blocking=True) - - with stream: - benchmark(launch, kernel, args) - stream.synchronize() - - -# Measure launch latency with many parameters using builtin parameter packing -@pytest.mark.skipif(skip_tests, reason="cupy is not installed") -@pytest.mark.benchmark(group="cupy") -def test_launch_latency_small_kernel_512_longlongs(benchmark): - module = cupy.RawModule(code=kernel_string) - kernel = module.get_function("small_kernel_512_longlongs") - cupy.cuda.set_allocator() - - args = [9223372036854775806] * 512 - args = tuple(args) - - stream = cupy.cuda.stream.Stream(non_blocking=True) - - with stream: - benchmark(launch, kernel, args) - stream.synchronize() - - -# Measure launch latency with many parameters using builtin parameter packing -@pytest.mark.skipif(skip_tests, reason="cupy is not installed") -@pytest.mark.benchmark(group="cupy") -def test_launch_latency_small_kernel_256_args(benchmark): - module = cupy.RawModule(code=kernel_string) - kernel = module.get_function("small_kernel_256_args") - cupy.cuda.set_allocator() - - args = [] - for _ in range(256): - args.append(cupy.cuda.alloc(ctypes.sizeof(ctypes.c_int))) - args = tuple(args) - - stream = cupy.cuda.stream.Stream(non_blocking=True) - - with stream: - benchmark(launch, kernel, args) - stream.synchronize() - - -# Measure launch latency with many parameters using builtin parameter packing -@pytest.mark.skipif(skip_tests, reason="cupy is not installed") -@pytest.mark.benchmark(group="cupy") -def test_launch_latency_small_kernel_16_args(benchmark): - module = cupy.RawModule(code=kernel_string) - kernel = module.get_function("small_kernel_16_args") - cupy.cuda.set_allocator() - - args = [] - for _ in range(16): - args.append(cupy.cuda.alloc(ctypes.sizeof(ctypes.c_int))) - args = tuple(args) - - stream = cupy.cuda.stream.Stream(non_blocking=True) - - with stream: - benchmark(launch, kernel, args) - stream.synchronize() diff --git a/cuda_bindings_12/benchmarks/test_launch_latency.py b/cuda_bindings_12/benchmarks/test_launch_latency.py deleted file mode 100755 index 541db98d556..00000000000 --- a/cuda_bindings_12/benchmarks/test_launch_latency.py +++ /dev/null @@ -1,336 +0,0 @@ -# SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. -# SPDX-License-Identifier: Apache-2.0 - -import ctypes - -import pytest -from conftest import ASSERT_DRV -from kernels import kernel_string - -from cuda import cuda - - -def launch(kernel, stream, args=(), arg_types=()): - cuda.cuLaunchKernel( - kernel, - 1, - 1, - 1, # grid dim - 1, - 1, - 1, # block dim - 0, - stream, # shared mem and stream - (args, arg_types), - 0, - ) # arguments - - -def launch_packed(kernel, stream, params): - cuda.cuLaunchKernel( - kernel, - 1, - 1, - 1, # grid dim - 1, - 1, - 1, # block dim - 0, - stream, # shared mem and stream - params, - 0, - ) # arguments - - -# Measure launch latency with no parmaeters -@pytest.mark.benchmark(group="launch-latency") -def test_launch_latency_empty_kernel(benchmark, init_cuda, load_module): - device, ctx, stream = init_cuda - module = load_module(kernel_string, device) - - err, func = cuda.cuModuleGetFunction(module, b"empty_kernel") - ASSERT_DRV(err) - - benchmark(launch, func, stream) - - cuda.cuCtxSynchronize() - - -# Measure launch latency with a single parameter -@pytest.mark.benchmark(group="launch-latency") -def test_launch_latency_small_kernel(benchmark, init_cuda, load_module): - device, ctx, stream = init_cuda - module = load_module(kernel_string, device) - - err, func = cuda.cuModuleGetFunction(module, b"small_kernel") - ASSERT_DRV(err) - - err, f = cuda.cuMemAlloc(ctypes.sizeof(ctypes.c_float)) - ASSERT_DRV(err) - - benchmark(launch, func, stream, args=(f,), arg_types=(None,)) - - cuda.cuCtxSynchronize() - - (err,) = cuda.cuMemFree(f) - ASSERT_DRV(err) - - -# Measure launch latency with many parameters using builtin parameter packing -@pytest.mark.benchmark(group="launch-latency") -def test_launch_latency_small_kernel_512_args(benchmark, init_cuda, load_module): - device, ctx, stream = init_cuda - module = load_module(kernel_string, device) - - err, func = cuda.cuModuleGetFunction(module, b"small_kernel_512_args") - ASSERT_DRV(err) - - args = [] - arg_types = [None] * 512 - for _ in arg_types: - err, p = cuda.cuMemAlloc(ctypes.sizeof(ctypes.c_int)) - ASSERT_DRV(err) - args.append(p) - - args = tuple(args) - arg_types = tuple(arg_types) - - benchmark(launch, func, stream, args=args, arg_types=arg_types) - - cuda.cuCtxSynchronize() - - for p in args: - (err,) = cuda.cuMemFree(p) - ASSERT_DRV(err) - - -@pytest.mark.benchmark(group="launch-latency") -def test_launch_latency_small_kernel_512_bools(benchmark, init_cuda, load_module): - device, ctx, stream = init_cuda - module = load_module(kernel_string, device) - - err, func = cuda.cuModuleGetFunction(module, b"small_kernel_512_bools") - ASSERT_DRV(err) - - args = [True] * 512 - arg_types = [ctypes.c_bool] * 512 - - args = tuple(args) - arg_types = tuple(arg_types) - - benchmark(launch, func, stream, args=args, arg_types=arg_types) - - cuda.cuCtxSynchronize() - - -@pytest.mark.benchmark(group="launch-latency") -def test_launch_latency_small_kernel_512_doubles(benchmark, init_cuda, load_module): - device, ctx, stream = init_cuda - module = load_module(kernel_string, device) - - err, func = cuda.cuModuleGetFunction(module, b"small_kernel_512_doubles") - ASSERT_DRV(err) - - args = [1.2345] * 512 - arg_types = [ctypes.c_double] * 512 - - args = tuple(args) - arg_types = tuple(arg_types) - - benchmark(launch, func, stream, args=args, arg_types=arg_types) - - cuda.cuCtxSynchronize() - - -@pytest.mark.benchmark(group="launch-latency") -def test_launch_latency_small_kernel_512_ints(benchmark, init_cuda, load_module): - device, ctx, stream = init_cuda - module = load_module(kernel_string, device) - - err, func = cuda.cuModuleGetFunction(module, b"small_kernel_512_ints") - ASSERT_DRV(err) - - args = [123] * 512 - arg_types = [ctypes.c_int] * 512 - - args = tuple(args) - arg_types = tuple(arg_types) - - benchmark(launch, func, stream, args=args, arg_types=arg_types) - - cuda.cuCtxSynchronize() - - -@pytest.mark.benchmark(group="launch-latency") -def test_launch_latency_small_kernel_512_bytes(benchmark, init_cuda, load_module): - device, ctx, stream = init_cuda - module = load_module(kernel_string, device) - - err, func = cuda.cuModuleGetFunction(module, b"small_kernel_512_chars") - ASSERT_DRV(err) - - args = [127] * 512 - arg_types = [ctypes.c_byte] * 512 - - args = tuple(args) - arg_types = tuple(arg_types) - - benchmark(launch, func, stream, args=args, arg_types=arg_types) - - cuda.cuCtxSynchronize() - - -@pytest.mark.benchmark(group="launch-latency") -def test_launch_latency_small_kernel_512_longlongs(benchmark, init_cuda, load_module): - device, ctx, stream = init_cuda - module = load_module(kernel_string, device) - - err, func = cuda.cuModuleGetFunction(module, b"small_kernel_512_longlongs") - ASSERT_DRV(err) - - args = [9223372036854775806] * 512 - arg_types = [ctypes.c_longlong] * 512 - - args = tuple(args) - arg_types = tuple(arg_types) - - benchmark(launch, func, stream, args=args, arg_types=arg_types) - - cuda.cuCtxSynchronize() - - -# Measure launch latency with many parameters using builtin parameter packing -@pytest.mark.benchmark(group="launch-latency") -def test_launch_latency_small_kernel_256_args(benchmark, init_cuda, load_module): - device, ctx, stream = init_cuda - module = load_module(kernel_string, device) - - err, func = cuda.cuModuleGetFunction(module, b"small_kernel_256_args") - ASSERT_DRV(err) - - args = [] - arg_types = [None] * 256 - for _ in arg_types: - err, p = cuda.cuMemAlloc(ctypes.sizeof(ctypes.c_int)) - ASSERT_DRV(err) - args.append(p) - - args = tuple(args) - arg_types = tuple(arg_types) - - benchmark(launch, func, stream, args=args, arg_types=arg_types) - - cuda.cuCtxSynchronize() - - for p in args: - (err,) = cuda.cuMemFree(p) - ASSERT_DRV(err) - - -# Measure launch latency with many parameters using builtin parameter packing -@pytest.mark.benchmark(group="launch-latency") -def test_launch_latency_small_kernel_16_args(benchmark, init_cuda, load_module): - device, ctx, stream = init_cuda - module = load_module(kernel_string, device) - - err, func = cuda.cuModuleGetFunction(module, b"small_kernel_16_args") - ASSERT_DRV(err) - - args = [] - arg_types = [None] * 16 - for _ in arg_types: - err, p = cuda.cuMemAlloc(ctypes.sizeof(ctypes.c_int)) - ASSERT_DRV(err) - args.append(p) - - args = tuple(args) - arg_types = tuple(arg_types) - - benchmark(launch, func, stream, args=args, arg_types=arg_types) - - cuda.cuCtxSynchronize() - - for p in args: - (err,) = cuda.cuMemFree(p) - ASSERT_DRV(err) - - -# Measure launch latency with many parameters, excluding parameter packing -@pytest.mark.benchmark(group="launch-latency") -def test_launch_latency_small_kernel_512_args_ctypes(benchmark, init_cuda, load_module): - device, ctx, stream = init_cuda - module = load_module(kernel_string, device) - - err, func = cuda.cuModuleGetFunction(module, b"small_kernel_512_args") - ASSERT_DRV(err) - - vals = [] - val_ps = [] - for i in range(512): - err, p = cuda.cuMemAlloc(ctypes.sizeof(ctypes.c_int)) - ASSERT_DRV(err) - vals.append(p) - val_ps.append(ctypes.c_void_p(int(vals[i]))) - - packagedParams = (ctypes.c_void_p * 512)() - for i in range(512): - packagedParams[i] = ctypes.addressof(val_ps[i]) - - benchmark(launch_packed, func, stream, packagedParams) - - cuda.cuCtxSynchronize() - - for p in vals: - (err,) = cuda.cuMemFree(p) - ASSERT_DRV(err) - - -def pack_and_launch(kernel, stream, params): - packed_params = (ctypes.c_void_p * len(params))() - ptrs = [0] * len(params) - for i in range(len(params)): - ptrs[i] = ctypes.c_void_p(int(params[i])) - packed_params[i] = ctypes.addressof(ptrs[i]) - - cuda.cuLaunchKernel(kernel, 1, 1, 1, 1, 1, 1, 0, stream, packed_params, 0) - - -# Measure launch latency plus parameter packing using ctypes -@pytest.mark.benchmark(group="launch-latency") -def test_launch_latency_small_kernel_512_args_ctypes_with_packing(benchmark, init_cuda, load_module): - device, ctx, stream = init_cuda - module = load_module(kernel_string, device) - - err, func = cuda.cuModuleGetFunction(module, b"small_kernel_512_args") - ASSERT_DRV(err) - - vals = [] - for i in range(512): - err, p = cuda.cuMemAlloc(ctypes.sizeof(ctypes.c_int)) - ASSERT_DRV(err) - vals.append(p) - - benchmark(pack_and_launch, func, stream, vals) - - cuda.cuCtxSynchronize() - - for p in vals: - (err,) = cuda.cuMemFree(p) - ASSERT_DRV(err) - - -# Measure launch latency with a single large struct parameter -@pytest.mark.benchmark(group="launch-latency") -def test_launch_latency_small_kernel_2048B(benchmark, init_cuda, load_module): - device, ctx, stream = init_cuda - module = load_module(kernel_string, device) - - err, func = cuda.cuModuleGetFunction(module, b"small_kernel_2048B") - ASSERT_DRV(err) - - class struct_2048B(ctypes.Structure): - _fields_ = [("values", ctypes.c_uint8 * 2048)] - - benchmark(launch, func, stream, args=(struct_2048B(),), arg_types=(None,)) - - cuda.cuCtxSynchronize() diff --git a/cuda_bindings_12/benchmarks/test_numba.py b/cuda_bindings_12/benchmarks/test_numba.py deleted file mode 100644 index 4f708bf69d7..00000000000 --- a/cuda_bindings_12/benchmarks/test_numba.py +++ /dev/null @@ -1,52 +0,0 @@ -# SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. -# SPDX-License-Identifier: Apache-2.0 - -import numpy as np -import pytest - -try: - from numba import cuda - - skip_tests = False -except ImportError: - skip_tests = True - - -def launch_empty(kernel, stream): - kernel[1, 1, stream]() - - -def launch(kernel, stream, arg): - kernel[1, 1, stream](arg) - - -# Measure launch latency with no parmaeters -@pytest.mark.skipif(skip_tests, reason="Numba is not installed") -@pytest.mark.benchmark(group="numba", min_rounds=1000) -def test_launch_latency_empty_kernel(benchmark): - stream = cuda.stream() - - @cuda.jit - def empty_kernel(): - return - - benchmark(launch_empty, empty_kernel, stream) - - cuda.synchronize() - - -# Measure launch latency with a single parameter -@pytest.mark.skipif(skip_tests, reason="Numba is not installed") -@pytest.mark.benchmark(group="numba", min_rounds=1000) -def test_launch_latency_small_kernel(benchmark): - stream = cuda.stream() - - arg = cuda.device_array(1, dtype=np.float32, stream=stream) - - @cuda.jit - def small_kernel(array): - array[0] = 0.0 - - benchmark(launch, small_kernel, stream, arg) - - cuda.synchronize() diff --git a/cuda_bindings_12/benchmarks/test_pointer_attributes.py b/cuda_bindings_12/benchmarks/test_pointer_attributes.py deleted file mode 100644 index 136d4be19f0..00000000000 --- a/cuda_bindings_12/benchmarks/test_pointer_attributes.py +++ /dev/null @@ -1,112 +0,0 @@ -# SPDX-FileCopyrightText: Copyright (c) 2021-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. -# SPDX-License-Identifier: Apache-2.0 - -import random - -import pytest -from conftest import ASSERT_DRV - -from cuda import cuda - -random.seed(0) - -idx = 0 - - -def query_attribute(attribute, ptrs): - global idx - ptr = ptrs[idx] - idx = (idx + 1) % len(ptrs) - - cuda.cuPointerGetAttribute(attribute, ptr) - - -def query_attributes(attributes, ptrs): - global idx - ptr = ptrs[idx] - idx = (idx + 1) % len(ptrs) - - cuda.cuPointerGetAttributes(len(attributes), attributes, ptr) - - -@pytest.mark.benchmark(group="pointer-attributes") -# Measure cuPointerGetAttribute in the same way as C benchmarks -def test_pointer_get_attribute(benchmark, init_cuda): - _ = init_cuda - - ptrs = [] - for _ in range(500): - err, ptr = cuda.cuMemAlloc(1 << 18) - ASSERT_DRV(err) - ptrs.append(ptr) - - random.shuffle(ptrs) - - benchmark(query_attribute, cuda.CUpointer_attribute.CU_POINTER_ATTRIBUTE_MEMORY_TYPE, ptrs) - - for p in ptrs: - (err,) = cuda.cuMemFree(p) - ASSERT_DRV(err) - - -@pytest.mark.benchmark(group="pointer-attributes") -# Measure cuPointerGetAttributes with all attributes -def test_pointer_get_attributes_all(benchmark, init_cuda): - _ = init_cuda - - ptrs = [] - for _ in range(500): - err, ptr = cuda.cuMemAlloc(1 << 18) - ASSERT_DRV(err) - ptrs.append(ptr) - - random.shuffle(ptrs) - - attributes = [ - cuda.CUpointer_attribute.CU_POINTER_ATTRIBUTE_CONTEXT, - cuda.CUpointer_attribute.CU_POINTER_ATTRIBUTE_MEMORY_TYPE, - cuda.CUpointer_attribute.CU_POINTER_ATTRIBUTE_HOST_POINTER, - cuda.CUpointer_attribute.CU_POINTER_ATTRIBUTE_P2P_TOKENS, - cuda.CUpointer_attribute.CU_POINTER_ATTRIBUTE_SYNC_MEMOPS, - cuda.CUpointer_attribute.CU_POINTER_ATTRIBUTE_BUFFER_ID, - cuda.CUpointer_attribute.CU_POINTER_ATTRIBUTE_IS_MANAGED, - cuda.CUpointer_attribute.CU_POINTER_ATTRIBUTE_DEVICE_ORDINAL, - cuda.CUpointer_attribute.CU_POINTER_ATTRIBUTE_IS_LEGACY_CUDA_IPC_CAPABLE, - cuda.CUpointer_attribute.CU_POINTER_ATTRIBUTE_RANGE_START_ADDR, - cuda.CUpointer_attribute.CU_POINTER_ATTRIBUTE_RANGE_SIZE, - cuda.CUpointer_attribute.CU_POINTER_ATTRIBUTE_MAPPED, - cuda.CUpointer_attribute.CU_POINTER_ATTRIBUTE_ALLOWED_HANDLE_TYPES, - cuda.CUpointer_attribute.CU_POINTER_ATTRIBUTE_IS_GPU_DIRECT_RDMA_CAPABLE, - cuda.CUpointer_attribute.CU_POINTER_ATTRIBUTE_ACCESS_FLAGS, - cuda.CUpointer_attribute.CU_POINTER_ATTRIBUTE_MEMPOOL_HANDLE, - ] - - benchmark(query_attributes, attributes, ptrs) - - for p in ptrs: - (err,) = cuda.cuMemFree(p) - ASSERT_DRV(err) - - -@pytest.mark.benchmark(group="pointer-attributes") -# Measure cuPointerGetAttributes with a single attribute -def test_pointer_get_attributes_single(benchmark, init_cuda): - _ = init_cuda - - ptrs = [] - for _ in range(500): - err, ptr = cuda.cuMemAlloc(1 << 18) - ASSERT_DRV(err) - ptrs.append(ptr) - - random.shuffle(ptrs) - - attributes = [ - cuda.CUpointer_attribute.CU_POINTER_ATTRIBUTE_MEMORY_TYPE, - ] - - benchmark(query_attributes, attributes, ptrs) - - for p in ptrs: - (err,) = cuda.cuMemFree(p) - ASSERT_DRV(err) diff --git a/cuda_bindings_12/pyproject.toml b/cuda_bindings_12/pyproject.toml index 493850ecb4e..42428ce67b5 100644 --- a/cuda_bindings_12/pyproject.toml +++ b/cuda_bindings_12/pyproject.toml @@ -116,12 +116,6 @@ exclude = ["cuda/bindings/_version.py"] "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", - "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) - [tool.setuptools_scm] root = ".." version_file = "cuda/bindings/_version.py" From 370ee2d2c28d18e9eadf4d2f232475104df92c0d Mon Sep 17 00:00:00 2001 From: "Ralf W. Grosse-Kunstleve" Date: Wed, 21 Oct 2026 09:43:30 -0700 Subject: [PATCH 06/10] Document CUDA 12 support for Python 3.15 Signed-off-by: Ralf W. Grosse-Kunstleve --- cuda_bindings_12/docs/source/install.md | 2 +- cuda_bindings_12/docs/source/support.rst | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/cuda_bindings_12/docs/source/install.md b/cuda_bindings_12/docs/source/install.md index 714180e0d89..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.10 - 3.14 +* 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 diff --git a/cuda_bindings_12/docs/source/support.rst b/cuda_bindings_12/docs/source/support.rst index bc92e60ce95..379bdb9d9ab 100644 --- a/cuda_bindings_12/docs/source/support.rst +++ b/cuda_bindings_12/docs/source/support.rst @@ -19,7 +19,7 @@ The ``cuda.bindings`` module has the following support policy: 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.10 - 3.14 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 From 4bd441391d91fd3b2be11b30a728d3442273d774 Mon Sep 17 00:00:00 2001 From: "Ralf W. Grosse-Kunstleve" Date: Wed, 21 Oct 2026 09:46:54 -0700 Subject: [PATCH 07/10] Enable pre-commit checks for the integrated CUDA 12 source Signed-off-by: Ralf W. Grosse-Kunstleve --- .pre-commit-config.yaml | 20 ++++++++++---------- cuda_bindings_12/pyproject.toml | 16 ++++++++++------ ruff.toml | 4 ++-- toolshed/check_generated_file_seals.py | 2 +- toolshed/check_spdx.py | 1 + 5 files changed, 24 insertions(+), 19 deletions(-) diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml index 38761afeedc..21b41bc52bf 100644 --- a/.pre-commit-config.yaml +++ b/.pre-commit-config.yaml @@ -14,9 +14,6 @@ ci: # Please update the rev: SHAs below with this command: # pre-commit autoupdate --freeze -# This unchanged maintenance-source import is inactive until its integration layer. -exclude: ^cuda_bindings_12/ - repos: # Runs first so a leaked credential blocks the commit before any formatter runs. # Self-installing: the hook downloads a pinned, checksum-verified trufflehog on @@ -31,7 +28,7 @@ repos: hooks: - id: ruff-check args: [--fix, --show-fixes] - exclude: (^cuda_bindings_13/cuda/bindings/_internal/_fast_enum\.py$)|(.*\.pyi$) + exclude: (^cuda_bindings_(?:12|13)/cuda/bindings/_internal/_fast_enum\.py$)|(.*\.pyi$) - id: ruff-format exclude: .*\.pyi$ @@ -71,7 +68,7 @@ repos: name: Check generated-file seals entry: python ./toolshed/check_generated_file_seals.py language: python - files: ^cuda_bindings_13/ + files: ^cuda_bindings_(?:12|13)/ types: [text] - id: check-pixi-cuda-version @@ -96,6 +93,9 @@ repos: - "_" # fake script name, because bash considers $0 (the first argument) to be the script name language: system files: '^.*/docs/source/.*\.md$' + # The imported CUDA 12.9 documentation retains its historical MyST + # sources; new documentation elsewhere must continue to use reST. + exclude: '^cuda_bindings_12/docs/source/' - id: stubgen-pyx-cuda-core name: Generate .pyi stubs for cuda_core @@ -155,7 +155,7 @@ repos: rev: "3e8a8703264a2f4a69428a0aa4dcb512790b2c8c" # frozen: v6.0.0 hooks: - id: check-added-large-files - exclude: cuda_bindings_13/cuda/bindings/nvml.pyx + exclude: '^(?:cuda_bindings_13/cuda/bindings/nvml\.pyx|cuda_bindings_12/cuda/bindings/(?:driver\.pyx|runtime\.pyx\.in|nvml\.pyx))$' - id: check-case-conflict - id: check-docstring-first - id: check-merge-conflict @@ -164,14 +164,14 @@ repos: - id: check-yaml - id: debug-statements - id: end-of-file-fixer - exclude: &gen_exclude '^(?:cuda_python/README\.md|(?:.*/)?CLAUDE\.md|(?:.*/)?\.git_archival\.txt|cuda_bindings_13/cuda/bindings/.*\.in?|cuda_bindings_13/docs/source/module/.*\.rst?|.*\.pyi)$' + exclude: &gen_exclude '^(?:cuda_python/README\.md|(?:.*/)?CLAUDE\.md|(?:.*/)?\.git_archival\.txt|cuda_bindings_(?:12|13)/cuda/bindings/.*\.in?|cuda_bindings_(?:12|13)/docs/source/module/.*\.rst?|.*\.pyi)$' - id: mixed-line-ending - id: trailing-whitespace exclude: | (?x)^(?: cuda_python/README\.md| - cuda_bindings_13/cuda/bindings/.*\.in?| - cuda_bindings_13/docs/source/module/.*\.rst?| + cuda_bindings_(?:12|13)/cuda/bindings/.*\.in?| + cuda_bindings_(?:12|13)/docs/source/module/.*\.rst?| .*\.patch$ )$ @@ -213,7 +213,7 @@ repos: hooks: - id: cython-lint args: [--no-pycodestyle] - exclude: ^cuda_bindings_13/ + exclude: ^cuda_bindings_(?:12|13)/ additional_dependencies: - Cython==3.2.9 diff --git a/cuda_bindings_12/pyproject.toml b/cuda_bindings_12/pyproject.toml index 42428ce67b5..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,7 +115,7 @@ exclude = ["cuda/bindings/_version.py"] ] "tests/**/*" = [ - "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/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", From a2a4158bd51f85df740902892bccb2b61582f763 Mon Sep 17 00:00:00 2001 From: "Ralf W. Grosse-Kunstleve" Date: Wed, 21 Oct 2026 10:22:35 -0700 Subject: [PATCH 08/10] Defer CUDA 12 runtime backports to a focused maintenance layer Signed-off-by: Ralf W. Grosse-Kunstleve --- .../cuda/bindings/_internal/utils.pyx | 9 +------- cuda_bindings_12/cuda/bindings/_lib/utils.pxd | 1 - cuda_bindings_12/cuda/bindings/_lib/utils.pxi | 6 +----- .../cuda/bindings/utils/__init__.py | 5 +---- cuda_bindings_12/tests/test_nvjitlink.py | 7 ------- cuda_bindings_12/tests/test_nvvm.py | 2 +- cuda_bindings_12/tests/test_utils.py | 21 ------------------- 7 files changed, 4 insertions(+), 47 deletions(-) diff --git a/cuda_bindings_12/cuda/bindings/_internal/utils.pyx b/cuda_bindings_12/cuda/bindings/_internal/utils.pyx index ce9012113df..df17a9e47df 100644 --- a/cuda_bindings_12/cuda/bindings/_internal/utils.pyx +++ b/cuda_bindings_12/cuda/bindings/_internal/utils.pyx @@ -98,14 +98,7 @@ cdef int get_nested_resource_ptr(nested_resource[ResT] &in_out_ptr, object obj, nested_ptr.reset(nested_vec, True) for i, obj_i in enumerate(obj): if ResT is char: - obj_i_type = type(obj_i) - if obj_i_type is str: - obj_i_bytes = obj_i.encode("utf-8") - elif obj_i_type is bytes: - obj_i_bytes = obj_i - else: - raise TypeError( - f"Expected str or bytes, got {obj_i_type.__name__}") + obj_i_bytes = ((obj_i)).encode() str_len = (len(obj_i_bytes)) + 1 # including null termination deref(nested_res_vec)[i].resize(str_len) obj_i_ptr = (obj_i_bytes) diff --git a/cuda_bindings_12/cuda/bindings/_lib/utils.pxd b/cuda_bindings_12/cuda/bindings/_lib/utils.pxd index c36ddfb0d12..706c6d3fc7a 100644 --- a/cuda_bindings_12/cuda/bindings/_lib/utils.pxd +++ b/cuda_bindings_12/cuda/bindings/_lib/utils.pxd @@ -152,7 +152,6 @@ cdef class _HelperCUcoredumpSettings: cdef cydriver.CUcoredumpSettings_enum _attrib cdef bint _is_getter cdef size_t _size - cdef object _references # keeps caller bytes alive so _charstar stays valid # Return values cdef cpp_bool _bool diff --git a/cuda_bindings_12/cuda/bindings/_lib/utils.pxi b/cuda_bindings_12/cuda/bindings/_lib/utils.pxi index eebdd17b41e..c11d0d0c197 100644 --- a/cuda_bindings_12/cuda/bindings/_lib/utils.pxi +++ b/cuda_bindings_12/cuda/bindings/_lib/utils.pxi @@ -734,8 +734,6 @@ cdef class _HelperCUcoredumpSettings: self._cptr = self._charstar self._size = 1024 else: - # Keep a reference so the borrowed _charstar buffer stays alive. - self._references = init_value self._charstar = init_value self._cptr = self._charstar self._size = len(init_value) @@ -758,9 +756,7 @@ cdef class _HelperCUcoredumpSettings: raise TypeError('Unsupported attribute: {}'.format(attr.name)) def __dealloc__(self): - # Only the getter owns heap storage; the setter borrows caller bytes. - if self._is_getter: - free(self._charstar) + pass @property def cptr(self): diff --git a/cuda_bindings_12/cuda/bindings/utils/__init__.py b/cuda_bindings_12/cuda/bindings/utils/__init__.py index d6136da14b0..62d083dc7c6 100644 --- a/cuda_bindings_12/cuda/bindings/utils/__init__.py +++ b/cuda_bindings_12/cuda/bindings/utils/__init__.py @@ -26,9 +26,6 @@ def get_cuda_native_handle(obj: Any) -> int: """ obj_type = type(obj) try: - getter = _handle_getters[obj_type] + return _handle_getters[obj_type](obj) except KeyError: raise TypeError("Unknown type: " + str(obj_type)) from None - # Deliberately outside the try: a KeyError raised by the getter itself is a - # bug in that getter, not an unregistered type. - return getter(obj) diff --git a/cuda_bindings_12/tests/test_nvjitlink.py b/cuda_bindings_12/tests/test_nvjitlink.py index bec58b7a048..19dc040e47d 100644 --- a/cuda_bindings_12/tests/test_nvjitlink.py +++ b/cuda_bindings_12/tests/test_nvjitlink.py @@ -100,13 +100,6 @@ def test_create_and_destroy(option): nvjitlink.destroy(handle) -@pytest.mark.human_authored -def test_create_and_destroy_bytes_options(): - handle = nvjitlink.create(1, [b"-arch=sm_80"]) - assert handle != 0 - nvjitlink.destroy(handle) - - @pytest.mark.parametrize("option", ARCHITECTURES) def test_complete_empty(option): handle = nvjitlink.create(1, [f"-arch={option}"]) diff --git a/cuda_bindings_12/tests/test_nvvm.py b/cuda_bindings_12/tests/test_nvvm.py index 3cea5a0f817..bd6f49f03f1 100644 --- a/cuda_bindings_12/tests/test_nvvm.py +++ b/cuda_bindings_12/tests/test_nvvm.py @@ -246,7 +246,7 @@ def test_get_buffer_empty(get_size, get_buffer): assert buffer == b"\x00" -@pytest.mark.parametrize("options", [[], ["-opt=0"], ["-opt=3", "-g"], [b"-opt=0"]]) +@pytest.mark.parametrize("options", [[], ["-opt=0"], ["-opt=3", "-g"]]) def test_compile_program_with_minimal_nvvm_ir(minimal_nvvmir, options): with nvvm_program() as prog: nvvm.add_module_to_program(prog, minimal_nvvmir, len(minimal_nvvmir), "FileNameHere.ll") diff --git a/cuda_bindings_12/tests/test_utils.py b/cuda_bindings_12/tests/test_utils.py index 33487c7f529..8904d65d2ed 100644 --- a/cuda_bindings_12/tests/test_utils.py +++ b/cuda_bindings_12/tests/test_utils.py @@ -92,27 +92,6 @@ def test_get_handle_error(target): handle = get_cuda_native_handle(target) -@pytest.mark.agent_authored(model="claude-opus-5") -def test_get_handle_does_not_report_a_registered_type_as_unknown(monkeypatch): - """A KeyError from inside a handle getter is a bug in that getter. - - Reporting it as "Unknown type" is wrong twice over: the type *is* - registered, and `from None` hides the traceback that would say otherwise. - """ - from cuda.bindings.utils import _handle_getters - - class Registered: - pass - - def getter(_obj): - raise KeyError("lookup inside the getter failed") - - monkeypatch.setitem(_handle_getters, Registered, getter) - - with pytest.raises(KeyError, match="lookup inside the getter failed"): - get_cuda_native_handle(Registered()) - - @pytest.mark.parametrize( "module", [ From baa6222128f7c1049cc4f06d07d94b75e2315591 Mon Sep 17 00:00:00 2001 From: "Ralf W. Grosse-Kunstleve" Date: Wed, 21 Oct 2026 10:22:40 -0700 Subject: [PATCH 09/10] Defer CUDA 12 NVML test maintenance to CI integration Signed-off-by: Ralf W. Grosse-Kunstleve --- cuda_bindings_12/tests/nvml/test_cuda.py | 26 ++++------------------ cuda_bindings_12/tests/nvml/test_init.py | 14 +----------- cuda_bindings_12/tests/nvml/test_pynvml.py | 18 +++++++++++++++ cuda_bindings_12/tests/nvml/util.py | 2 +- 4 files changed, 24 insertions(+), 36 deletions(-) diff --git a/cuda_bindings_12/tests/nvml/test_cuda.py b/cuda_bindings_12/tests/nvml/test_cuda.py index aaab23e55fe..d05e2634a12 100644 --- a/cuda_bindings_12/tests/nvml/test_cuda.py +++ b/cuda_bindings_12/tests/nvml/test_cuda.py @@ -1,8 +1,6 @@ # SPDX-FileCopyrightText: Copyright (c) 2025-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: Apache-2.0 -import pytest - import cuda.bindings.driver as cuda from cuda.bindings import nvml @@ -17,15 +15,10 @@ def get_nvml_device_names(): for idx in range(num_devices): handle = nvml.device_get_handle_by_index_v2(idx) name = nvml.device_get_name(handle) - try: - info = nvml.device_get_pci_info_v3(handle) - except nvml.NotSupportedError: - bus_id = -1 - else: - bus_id = info.bus - assert isinstance(bus_id, int) + info = nvml.device_get_pci_info_v3(handle) + assert isinstance(info.bus, int) assert isinstance(name, str) - result.append({"name": name, "id": bus_id}) + result.append({"name": name, "id": info.bus}) return result @@ -61,15 +54,4 @@ def test_cuda_device_order(): cuda_devices = get_cuda_device_names() nvml_devices = get_nvml_device_names() - for kind in ("Orin", "Thor"): - if any(kind in device["name"] for device in nvml_devices): - pytest.skip(f"Skipping test on {kind}, which has non-standard device naming") - - def compare(cuda_device, nvml_device): - return cuda_device["name"] == nvml_device["name"] and ( - nvml_device["id"] == -1 or cuda_device["id"] == nvml_device["id"] - ) - - assert len(cuda_devices) <= len(nvml_devices) - for cuda_device in cuda_devices: - assert any(compare(cuda_device, nvml_device) for nvml_device in nvml_devices) + assert cuda_devices == nvml_devices, "CUDA and NVML device lists do not match" diff --git a/cuda_bindings_12/tests/nvml/test_init.py b/cuda_bindings_12/tests/nvml/test_init.py index bcda0591bb6..94b489ab45f 100644 --- a/cuda_bindings_12/tests/nvml/test_init.py +++ b/cuda_bindings_12/tests/nvml/test_init.py @@ -6,15 +6,7 @@ import pytest -from cuda.bindings import driver, nvml - - -def _driver_version_at_least_13_4(): - (status,) = driver.cuInit(0) - assert status == driver.CUresult.CUDA_SUCCESS - status, version = driver.cuDriverGetVersion() - assert status == driver.CUresult.CUDA_SUCCESS - return version >= 13040 +from cuda.bindings import nvml def assert_nvml_is_initialized(): @@ -44,14 +36,10 @@ def test_devices_are_the_same_architecture(all_devices): @pytest.mark.skipif(sys.platform == "win32", reason="Test not supported on Windows") -@pytest.mark.thread_unsafe(reason="nvml init affects other threads") def test_init_ref_count(): """ Verifies that we can call NVML shutdown and init(2) multiple times, and that ref counting works """ - if _driver_version_at_least_13_4(): - pytest.skip("Init behavior changed in CUDA 13.4") - with pytest.raises(nvml.UninitializedError): nvml.shutdown() diff --git a/cuda_bindings_12/tests/nvml/test_pynvml.py b/cuda_bindings_12/tests/nvml/test_pynvml.py index e3ea533dd83..a96b2c87575 100644 --- a/cuda_bindings_12/tests/nvml/test_pynvml.py +++ b/cuda_bindings_12/tests/nvml/test_pynvml.py @@ -4,6 +4,8 @@ # A set of tests ported from https://github.com/gpuopenanalytics/pynvml/blob/11.5.3/pynvml/tests/test_nvml.py import os +import time + import pytest from cuda.bindings import nvml @@ -145,6 +147,22 @@ def test_device_get_power_usage(ngpus, handles): assert power_mwatts >= 0.0 +def test_device_get_total_energy_consumption(ngpus, handles): + for i in range(ngpus): + with unsupported_before(handles[i], nvml.DeviceArch.VOLTA): + energy_mjoules1 = nvml.device_get_total_energy_consumption(handles[i]) + + for j in range(10): # idle for 150 ms + time.sleep(0.015) # and check for increase every 15 ms + with unsupported_before(handles[i], nvml.DeviceArch.VOLTA): + energy_mjoules2 = nvml.device_get_total_energy_consumption(handles[i]) + assert energy_mjoules2 >= energy_mjoules1 + if energy_mjoules2 > energy_mjoules1: + break + else: + raise AssertionError("energy did not increase across 150 ms interval") + + # [Skipping] pynvml.nvmlDeviceGetGpuOperationMode # [Skipping] pynvml.nvmlDeviceGetCurrentGpuOperationMode # [Skipping] pynvml.nvmlDeviceGetPendingGpuOperationMode diff --git a/cuda_bindings_12/tests/nvml/util.py b/cuda_bindings_12/tests/nvml/util.py index 06cb1c2277a..038fe58d8be 100644 --- a/cuda_bindings_12/tests/nvml/util.py +++ b/cuda_bindings_12/tests/nvml/util.py @@ -43,5 +43,5 @@ def supports_ecc(device): def supports_nvlink(device): fields = nvml.FieldValue(1) - fields[0].field_id = nvml.FieldId.DEV_NVLINK_GET_STATE + fields[0].field_id = nvml.FI.DEV_NVLINK_GET_STATE return nvml.device_get_field_values(device, fields)[0].nvml_return == nvml.Return.SUCCESS From 2f01e9869af332436e6265b9c1645b91e369b28f Mon Sep 17 00:00:00 2001 From: "Ralf W. Grosse-Kunstleve" Date: Wed, 21 Oct 2026 11:10:40 -0700 Subject: [PATCH 10/10] Pin CUDA 12 documentation links to a published snapshot Signed-off-by: Ralf W. Grosse-Kunstleve --- cuda_bindings_12/DESCRIPTION.rst | 7 +++++-- cuda_bindings_12/README.md | 3 +++ cuda_bindings_12/build_hooks.py | 4 +++- 3 files changed, 11 insertions(+), 3 deletions(-) diff --git a/cuda_bindings_12/DESCRIPTION.rst b/cuda_bindings_12/DESCRIPTION.rst index 52fbdf848e8..31e5edfc6a4 100644 --- a/cuda_bindings_12/DESCRIPTION.rst +++ b/cuda_bindings_12/DESCRIPTION.rst @@ -5,10 +5,13 @@ cuda-bindings: Low-level CUDA interfaces **************************************** -`cuda.bindings `_ is a standard set of low-level interfaces, providing full coverage of and 1:1 access to the CUDA host APIs from Python. Checkout the `Overview `_ for the workflow and performance results. +.. Keep documentation URLs aligned with README.md and build_hooks.py. + Use a published CUDA 12 snapshot; latest follows the current CUDA major. + +`cuda.bindings `_ is a standard set of low-level interfaces, providing full coverage of and 1:1 access to the CUDA host APIs from Python. Checkout the `Overview `_ for the workflow and performance results. * `Repository `_ -* `Documentation `_ +* `Documentation `_ * `Examples `_ * `Issue tracker `_ diff --git a/cuda_bindings_12/README.md b/cuda_bindings_12/README.md index 49fc1ca6470..b0e6f29c14a 100644 --- a/cuda_bindings_12/README.md +++ b/cuda_bindings_12/README.md @@ -2,6 +2,9 @@ `cuda.bindings` is a standard set of low-level interfaces, providing full coverage of and access to the CUDA host APIs from Python. Checkout the [Overview page](https://nvidia.github.io/cuda-python/cuda-bindings/12.9.7/overview.html) for the workflow and performance results. + + ## Installing Please refer to the [Installation page](https://nvidia.github.io/cuda-python/cuda-bindings/12.9.7/install.html) for instructions and required/optional dependencies. diff --git a/cuda_bindings_12/build_hooks.py b/cuda_bindings_12/build_hooks.py index 4cca4063475..7258d7abf53 100644 --- a/cuda_bindings_12/build_hooks.py +++ b/cuda_bindings_12/build_hooks.py @@ -64,7 +64,9 @@ def __getattr__(name): _CYDRIVER_PXD = Path(__file__).resolve().parent / "cuda" / "bindings" / "cydriver.pxd" _GENERATED_VERSION_RE = re.compile(r"^cdef enum:\s*CUDA_VERSION\s*=\s*(\d+)\s*$") _CUDA_H_VERSION_RE = re.compile(r"^#\s*define\s+CUDA_VERSION\s+(\d+)\s*$") -_INSTALL_URL = "https://nvidia.github.io/cuda-python/cuda-bindings/12.9.10/install.html#installing-from-source" +# Keep this published CUDA 12 snapshot aligned with README.md and DESCRIPTION.rst; +# latest follows the current CUDA major. Refresh when newer CUDA 12 docs are published. +_INSTALL_URL = "https://nvidia.github.io/cuda-python/cuda-bindings/12.9.7/install.html#installing-from-source" @functools.cache