Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension


Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
8 changes: 6 additions & 2 deletions .github/labeler.yml
Original file line number Diff line number Diff line change
Expand Up @@ -7,11 +7,15 @@

cuda.bindings:
- changed-files:
- any-glob-to-any-file: 'cuda_bindings/**'
- any-glob-to-any-file:
- 'cuda_bindings/**'
- 'samples/cuda_bindings/**'

cuda.core:
- changed-files:
- any-glob-to-any-file: 'cuda_core/**'
- any-glob-to-any-file:
- 'cuda_core/**'
- 'samples/cuda_core/**'

cuda.pathfinder:
- changed-files:
Expand Down
6 changes: 3 additions & 3 deletions .github/workflows/ci.yml
Original file line number Diff line number Diff line change
Expand Up @@ -150,12 +150,12 @@ jobs:
}

{
echo "bindings=$(has_match '^cuda_bindings/')"
echo "core=$(has_match '^cuda_core/')"
echo "bindings=$(has_match '^(cuda_bindings/|samples/cuda_bindings/)')"
echo "core=$(has_match '^(cuda_core/|samples/cuda_core/)')"
echo "pathfinder=$(has_match '^cuda_pathfinder/')"
echo "python_meta=$(has_match '^cuda_python/')"
echo "test_helpers=$(has_match '^cuda_python_test_helpers/')"
echo "shared=$(has_match '^(\.github/|ci/|scripts/|toolshed/|conftest\.py$|pyproject\.toml$|pixi\.(toml|lock)$|pytest\.ini$|ruff\.toml$)')"
echo "shared=$(has_match '^(\.github/|ci/|scripts/|toolshed/|samples/README\.md$|conftest\.py$|pyproject\.toml$|pixi\.(toml|lock)$|pytest\.ini$|ruff\.toml$)')"
} >> "$GITHUB_OUTPUT"

- name: Compose gating outputs
Expand Down
5 changes: 4 additions & 1 deletion .spdx-ignore
Original file line number Diff line number Diff line change
Expand Up @@ -6,7 +6,10 @@ LICENSE*
.gitattributes
.gitignore
requirements*.txt
cuda_bindings/examples/*

# Samples are synced to NVIDIA/cuda-samples on release and carry either the
# upstream verbose BSD-style copyright header or an SPDX identifier.
samples/**/*.py

# Vendored
cuda_core/cuda/core/_include/dlpack.h
Expand Down
6 changes: 6 additions & 0 deletions ci/tools/run-tests
Original file line number Diff line number Diff line change
Expand Up @@ -102,6 +102,12 @@ elif [[ "${test_module}" == "core" || "${test_module}" == nightly-* ]]; then
# Constrain cuda-toolkit to the requested CTK version to avoid
# pip pulling in a newer nvidia-cuda-runtime that conflicts with it.
pip install "${CORE_WHL[@]}" --group "test-cu${TEST_CUDA_MAJOR}${FREE_THREADING}" "cuda-toolkit==${TEST_CUDA_MAJOR_MINOR}.*"
# Optional deps needed by some cuda.core samples exercised in this suite.
# CuPy already comes from the
# test-cuXX group; install the rest best-effort so those samples run
# instead of being waived. Samples with heavier/unavailable deps stay waived.
echo "Installing optional sample deps (nvtx, pillow)"
pip install nvtx pillow || echo "Warning: optional sample deps install failed; affected samples will be waived"
echo "Installed packages before core tests:"
pip list
echo "Running core tests"
Expand Down
4 changes: 4 additions & 0 deletions conftest.py
Original file line number Diff line number Diff line change
Expand Up @@ -41,6 +41,10 @@ def pytest_collection_modifyitems(config, items): # noqa: ARG001
if nodeid.startswith("tests/integration/") or "/tests/integration/" in nodeid:
item.add_marker(pytest.mark.smoke)

# Each package owns the sample wrapper under its example_tests subtree.
if "example_tests/test_samples.py" in nodeid:
item.add_marker(pytest.mark.samples)

# Cython tests (any tests/cython subtree)
if (
"/tests/cython/" in nodeid
Expand Down
2 changes: 1 addition & 1 deletion cuda_bindings/AGENTS.md
Original file line number Diff line number Diff line change
Expand Up @@ -39,7 +39,7 @@ subpackage in the `cuda-python` monorepo.
- **Cython tests**:
- build: `tests/cython/build_tests.sh` (or platform equivalent)
- run: `pytest tests/cython/`
- **Examples**: example coverage is pytest-based under `examples/`.
- **Samples**: sample coverage is pytest-based under `../samples/cuda_bindings/`.
- **Benchmarks**: run with `pytest --benchmark-only benchmarks/` when needed.

The `legacy_tests` subdirectory tests the old pre-v2 APIs of `driver`, `runtime`
Expand Down
2 changes: 1 addition & 1 deletion cuda_bindings/DESCRIPTION.rst
Original file line number Diff line number Diff line change
Expand Up @@ -13,7 +13,7 @@ cuda-bindings: Low-level CUDA interfaces

* `Repository <https://github.com/NVIDIA/cuda-python/tree/main/cuda_bindings>`_
* `Documentation <https://nvidia.github.io/cuda-python/cuda-bindings/>`_
* `Examples <https://github.com/NVIDIA/cuda-python/tree/main/cuda_bindings/examples>`_
* `Examples <https://github.com/NVIDIA/cuda-python/tree/main/samples/cuda_bindings>`_
* `Issue tracker <https://github.com/NVIDIA/cuda-python/issues/>`_

For the installation instruction, please refer to the `Installation <https://nvidia.github.io/cuda-python/cuda-bindings/latest/install.html>`_ page.
17 changes: 0 additions & 17 deletions cuda_bindings/cuda/bindings/_example_helpers/__init__.py

This file was deleted.

90 changes: 0 additions & 90 deletions cuda_bindings/cuda/bindings/_example_helpers/common.py

This file was deleted.

48 changes: 0 additions & 48 deletions cuda_bindings/cuda/bindings/_example_helpers/helper_cuda.py

This file was deleted.

15 changes: 0 additions & 15 deletions cuda_bindings/cuda/bindings/_example_helpers/helper_string.py

This file was deleted.

16 changes: 14 additions & 2 deletions cuda_bindings/docs/source/conf.py
Original file line number Diff line number Diff line change
Expand Up @@ -122,11 +122,23 @@ def _html_baseurl():

extlinks = {
"cuda-bindings-example": (
f"https://github.com/NVIDIA/cuda-python/blob/{GITHUB_EXAMPLES_REF}/cuda_bindings/examples/%s",
f"https://github.com/NVIDIA/cuda-python/blob/{GITHUB_EXAMPLES_REF}/samples/cuda_bindings/%s",
"%s",
),
"cuda-bindings-examples": (
f"https://github.com/NVIDIA/cuda-python/tree/{GITHUB_EXAMPLES_REF}/cuda_bindings/examples%s",
f"https://github.com/NVIDIA/cuda-python/tree/{GITHUB_EXAMPLES_REF}/samples/cuda_bindings%s",
"%s",
),
"sample": (
f"https://github.com/NVIDIA/cuda-python/tree/{GITHUB_EXAMPLES_REF}/samples/cuda_bindings/%s",
"%s",
),
"samples": (
f"https://github.com/NVIDIA/cuda-python/tree/{GITHUB_EXAMPLES_REF}/samples/cuda_bindings%s",
"%s",
),
"cuda-core-sample": (
f"https://github.com/NVIDIA/cuda-python/tree/{GITHUB_EXAMPLES_REF}/samples/cuda_core/%s",
"%s",
),
}
Expand Down
81 changes: 45 additions & 36 deletions cuda_bindings/docs/source/examples.rst
Original file line number Diff line number Diff line change
Expand Up @@ -4,62 +4,71 @@
Examples
========

This page links to the ``cuda.bindings`` examples shipped in the
:cuda-bindings-examples:`cuda-python repository </>`.
Use it as a quick index when you want a runnable sample for a specific API area
or CUDA feature.
The canonical, runnable examples for ``cuda.bindings`` live under
:samples:`samples/cuda_bindings/ </>` in the cuda-python repository. Each
sample is a self-contained directory with
its own ``README.md``, ``requirements.txt``, and PEP 723 dependency block,
and every sample is exercised as part of the ``cuda.bindings`` test suite.

The cuda-bindings-flavored samples preserve the same category structure the
``cuda-samples`` repository uses, so they are grouped under
``samples/cuda_bindings/0_Introduction/``,
``samples/cuda_bindings/2_Concepts_and_Techniques/``,
``samples/cuda_bindings/3_CUDA_Features/``,
``samples/cuda_bindings/4_CUDA_Libraries/``, and
``samples/cuda_bindings/extra/``.

Introduction
------------

- :cuda-bindings-example:`clock_nvrtc.py <0_Introduction/clock_nvrtc.py>`
uses NVRTC-compiled CUDA code and the device clock to time a reduction
kernel.
- :cuda-bindings-example:`simple_cubemap_texture.py <0_Introduction/simple_cubemap_texture.py>`
demonstrates cubemap texture sampling and transformation.
- :cuda-bindings-example:`simple_p2p.py <0_Introduction/simple_p2p.py>`
shows peer-to-peer memory access and transfers between multiple GPUs.
- :cuda-bindings-example:`simple_zero_copy.py <0_Introduction/simple_zero_copy.py>`
uses zero-copy mapped host memory for vector addition.
- :cuda-bindings-example:`system_wide_atomics.py <0_Introduction/system_wide_atomics.py>`
demonstrates system-wide atomic operations on managed memory.
- :cuda-bindings-example:`vector_add_drv.py <0_Introduction/vector_add_drv.py>`
uses the CUDA Driver API and unified virtual addressing for vector addition.
- :cuda-bindings-example:`vector_add_mmap.py <0_Introduction/vector_add_mmap.py>`
uses virtual memory management APIs such as ``cuMemCreate`` and
``cuMemMap`` for vector addition.
- :sample:`0_Introduction/clockNvrtc <0_Introduction/clockNvrtc/>`
uses NVRTC-compiled CUDA code and the device clock intrinsic to time a
reduction kernel; also demonstrates dynamic shared memory.
- :sample:`0_Introduction/simpleCubemapTexture <0_Introduction/simpleCubemapTexture/>`
demonstrates cubemap texture sampling, ``cudaMemcpy3D``, and bindless
texture objects.
- :sample:`0_Introduction/systemWideAtomics <0_Introduction/systemWideAtomics/>`
exercises every ``atomic*_system`` intrinsic on managed memory.
- :sample:`0_Introduction/vectorAddDriverApi <0_Introduction/vectorAddDriverApi/>`
is the "hello world" of the raw driver API in Python.
- :sample:`0_Introduction/vectorAddMmap <0_Introduction/vectorAddMmap/>`
uses the Virtual Memory Management API (``cuMemCreate``, ``cuMemMap``,
``cuMemSetAccess``) to stripe an allocation across peer-capable devices.

Concepts and techniques
-----------------------

- :cuda-bindings-example:`stream_ordered_allocation.py <2_Concepts_and_Techniques/stream_ordered_allocation.py>`
demonstrates ``cudaMallocAsync`` and ``cudaFreeAsync`` together with
- :sample:`2_Concepts_and_Techniques/streamOrderedAllocation <2_Concepts_and_Techniques/streamOrderedAllocation/>`
demonstrates ``cudaMallocAsync`` / ``cudaFreeAsync`` together with
memory-pool release thresholds.

CUDA features
-------------

- :cuda-bindings-example:`global_to_shmem_async_copy.py <3_CUDA_Features/global_to_shmem_async_copy.py>`
compares asynchronous global-to-shared-memory copy strategies in matrix
multiplication kernels.
- :cuda-bindings-example:`simple_cuda_graphs.py <3_CUDA_Features/simple_cuda_graphs.py>`
shows both manual CUDA graph construction and stream-capture-based replay.
- :sample:`3_CUDA_Features/globalToShmemAsyncCopy <3_CUDA_Features/globalToShmemAsyncCopy/>`
compares asynchronous global-to-shared-memory copy strategies
(``cuda::memcpy_async`` + ``cuda::pipeline`` / ``cuda::barrier``) in
matrix-multiplication kernels.
- :sample:`3_CUDA_Features/cudaGraphsManualNodes <3_CUDA_Features/cudaGraphsManualNodes/>`
shows both manual CUDA-graph construction (``cudaGraphAdd*Node``) and
stream-capture-based replay.

Libraries and tools
-------------------

- :cuda-bindings-example:`conjugate_gradient_multi_block_cg.py <4_CUDA_Libraries/conjugate_gradient_multi_block_cg.py>`
- :sample:`4_CUDA_Libraries/conjugateGradientMultiBlockCG <4_CUDA_Libraries/conjugateGradientMultiBlockCG/>`
implements a conjugate-gradient solver with cooperative groups and
multi-block synchronization.
- :cuda-bindings-example:`nvidia_smi.py <4_CUDA_Libraries/nvidia_smi.py>`
multi-block grid synchronization.
- :sample:`4_CUDA_Libraries/nvidiaSmi <4_CUDA_Libraries/nvidiaSmi/>`
uses NVML to implement a Python subset of ``nvidia-smi``.

Advanced and interoperability
-----------------------------

- :cuda-bindings-example:`iso_fd_modelling.py <extra/iso_fd_modelling.py>`
runs isotropic finite-difference wave propagation across multiple GPUs with
peer-to-peer halo exchange.
- :cuda-bindings-example:`jit_program.py <extra/jit_program.py>`
JIT-compiles a SAXPY kernel with NVRTC and launches it through the Driver
API.
- :sample:`extra/isoFdModelling <extra/isoFdModelling/>`
runs isotropic finite-difference wave propagation across multiple GPUs
with peer-to-peer halo exchange.
- :sample:`extra/jitProgram <extra/jitProgram/>`
JIT-compiles a SAXPY kernel with NVRTC and launches it through the
Driver API -- the low-level companion to
:cuda-core-sample:`cuda.core's jitLtoLinking sample <jitLtoLinking/>`.
Loading