Migrating cuda-python samples from cuda-samples

.github/labeler.yml

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

.spdx-ignore

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

ci/tools/run-tests

@@ -104,6 +104,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}" "${TEST_FT_GROUP[@]}" "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"

cuda_bindings/AGENTS.md

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

cuda_bindings/DESCRIPTION.rst

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

cuda_bindings/docs/source/conf.py

@@ -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",
     ),
 }

cuda_bindings/docs/source/examples.rst

@@ -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/>`.

cuda_bindings/docs/source/overview.rst

@@ -524,7 +524,7 @@ 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#>`_.
 
 For this example, lets use the ``transformKernel`` from
-:cuda-bindings-example:`simple_cubemap_texture.py <0_Introduction/simple_cubemap_texture.py>`.
+:cuda-bindings-example:`simpleCubemapTexture.py <0_Introduction/simpleCubemapTexture/simpleCubemapTexture.py>`.
 The :doc:`examples` page links to more samples covering textures, graphs,
 memory mapping, and multi-GPU workflows.