From 8b0d89349c7520322f310766f41650611938d700 Mon Sep 17 00:00:00 2001 From: 0z5a Date: Thu, 29 Oct 2026 11:27:07 +0100 Subject: [PATCH 1/3] samples: demonstrate Grace Hopper host and managed memory Signed-off-by: 0z5a --- samples/cuda_core/graceHopperMemory/README.md | 55 +++++ .../graceHopperMemory/graceHopperMemory.py | 191 ++++++++++++++++++ .../graceHopperMemory/requirements.txt | 3 + 3 files changed, 249 insertions(+) create mode 100644 samples/cuda_core/graceHopperMemory/README.md create mode 100644 samples/cuda_core/graceHopperMemory/graceHopperMemory.py create mode 100644 samples/cuda_core/graceHopperMemory/requirements.txt diff --git a/samples/cuda_core/graceHopperMemory/README.md b/samples/cuda_core/graceHopperMemory/README.md new file mode 100644 index 00000000000..ac08b3b3a6e --- /dev/null +++ b/samples/cuda_core/graceHopperMemory/README.md @@ -0,0 +1,55 @@ +# Sample: Grace Hopper Memory + +This sample prints the CUDA NUMA and pageable-memory properties and the NVML +C2C status, then runs a kernel directly on a pageable NumPy array using the +host page tables. It compares reads from a host-NUMA pinned buffer, pageable +NumPy memory, device memory and managed memory after prefetch to either the +host NUMA node or the device. Every kernel result is checked. + +The read kernel sums 1 GiB of `float32` values using a grid-stride loop, one +shared-memory reduction and one atomic operation per block. The reported rate +is bytes read divided by host-clock time around a synchronized kernel launch, +the median of ten runs. It is not a tuned peak-bandwidth benchmark. Pageable +and pinned reads receive one untimed warmup. Before **each** managed-memory +measurement, the sample prefetches again to the indicated location and waits +for the prefetch to finish; the host-prefetched kernel time can include +migration to the GPU. + +## Requirements + +- A Grace Hopper system with CUDA pageable access through host page tables, + concurrent managed access and host NUMA memory pools. Systems without these + properties waive with exit code 2 (configurable via + `CUDA_PYTHON_SAMPLE_WAIVER_EXIT_CODE`). +- CUDA Toolkit 13.x, `cuda-python`, `cuda-core` 1.3.0 or newer, and NumPy + 2.2.5 or newer. The required cuda-core APIs are present on the current main + branch; 1.3.0 may not yet be released. +- About 3 GiB of host memory, 1 GiB of device memory, and managed-memory + capacity. Allocate these only on a system where the GPU is available to you. + +Install from `requirements.txt` after the required cuda-core release is +available, or build the matching development version from the repository. + +```bash +cd samples/cuda_core/graceHopperMemory +python graceHopperMemory.py +``` + +## Historical probe measurements + +A separate prototype using the same kernel and input size was measured on one +GH200 480GB system (driver 595.58.03, CUDA 13.4.2, Python 3.12). Three runs, +each reporting the median of ten launches, gave the following ranges: + +| Memory | Read rate (GB/s) | +| --- | ---: | +| Pinned, host NUMA node 0 | 339–345 | +| Pageable NumPy array | 224–225 | +| Device memory | 2386–2417 | +| Managed, first kernel after host prefetch | 5.1–5.4 | +| Managed, after device prefetch | 2380–2393 | + +The host-prefetched managed result includes first-touch migration on that +system; it should not be interpreted as direct C2C bandwidth. The prototype +results are historical measurements, not results from this sample revision. +Rates depend on topology, system load, memory placement and kernel design. diff --git a/samples/cuda_core/graceHopperMemory/graceHopperMemory.py b/samples/cuda_core/graceHopperMemory/graceHopperMemory.py new file mode 100644 index 00000000000..636ce72ab9b --- /dev/null +++ b/samples/cuda_core/graceHopperMemory/graceHopperMemory.py @@ -0,0 +1,191 @@ +# Copyright (c) 2026, NVIDIA CORPORATION. All rights reserved. +# +# Redistribution and use in source and binary forms, with or without +# modification, are permitted provided that the following conditions +# are met: +# * Redistributions of source code must retain the above copyright +# notice, this list of conditions and the following disclaimer. +# * Redistributions in binary form must reproduce the above copyright +# notice, this list of conditions and the following disclaimer in the +# documentation and/or other materials provided with the distribution. +# * Neither the name of NVIDIA CORPORATION nor the names of its +# contributors may be used to endorse or promote products derived +# from this software without specific prior written permission. +# +# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY +# EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +# IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR +# PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR +# CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, +# EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, +# PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR +# PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY +# OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +# (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +# OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + +# /// script +# dependencies = ["cuda-python>=13.0.0", "cuda-core>=1.3.0", "numpy>=2.2.5"] +# /// + +"""Compare GPU reads from host, device, and managed memory on Grace Hopper.""" + +import os +import sys +import time + +import numpy as np + +from cuda.core import ( + Device, + DeviceMemoryResource, + Host, + LaunchConfig, + ManagedMemoryResource, + PinnedMemoryResource, + PinnedMemoryResourceOptions, + Program, + ProgramOptions, + launch, +) + +EXIT_WAIVED = int(os.environ.get("CUDA_PYTHON_SAMPLE_WAIVER_EXIT_CODE", "2")) +N = 1 << 28 # 1 GiB of float32 values +REPS = 10 + +SOURCE = r""" +extern "C" __global__ void read_sum(const float *x, unsigned long long n, double *out) { + __shared__ float part[512]; + float acc = 0.f; + for (unsigned long long i = blockIdx.x * (unsigned long long)blockDim.x + threadIdx.x; i < n; + i += (unsigned long long)gridDim.x * blockDim.x) + acc += x[i]; + part[threadIdx.x] = acc; + __syncthreads(); + for (unsigned s = blockDim.x / 2; s > 0; s >>= 1) { + if (threadIdx.x < s) part[threadIdx.x] += part[threadIdx.x + s]; + __syncthreads(); + } + if (threadIdx.x == 0) atomicAdd(out, (double)part[0]); +} +extern "C" __global__ void scale(float *x, unsigned long long n, float a) { + for (unsigned long long i = blockIdx.x * (unsigned long long)blockDim.x + threadIdx.x; i < n; + i += (unsigned long long)gridDim.x * blockDim.x) + x[i] *= a; +} +extern "C" __global__ void fill(float *x, unsigned long long n, float v) { + for (unsigned long long i = blockIdx.x * (unsigned long long)blockDim.x + threadIdx.x; i < n; + i += (unsigned long long)gridDim.x * blockDim.x) + x[i] = v; +} +extern "C" __global__ void zero(double *out) { *out = 0.0; } +""" + + +def main(): + dev = Device(0) + dev.set_current() + props = dev.properties + print(f"Device: {dev.name}, compute capability {dev.compute_capability}") + for name in ( + "numa_config", + "numa_id", + "host_numa_id", + "pageable_memory_access", + "pageable_memory_access_uses_host_page_tables", + "concurrent_managed_access", + ): + print(f" {name}: {getattr(props, name)}") + print(f" system.Device.is_c2c_enabled: {dev.to_system_device().is_c2c_enabled}") + + if not props.pageable_memory_access_uses_host_page_tables: + print("Host page tables are unavailable for pageable GPU access; waiving sample.") + return EXIT_WAIVED + if ( + not props.concurrent_managed_access + or not props.host_memory_pools_supported + or not props.host_numa_memory_pools_supported + or props.host_numa_id < 0 + ): + print("Managed access, host memory pools, or host NUMA support is unavailable; waiving sample.") + return EXIT_WAIVED + + stream = dev.create_stream() + arch = f"sm_{dev.compute_capability.major}{dev.compute_capability.minor}" + module = Program(SOURCE, code_type="c++", options=ProgramOptions(arch=arch)).compile("cubin") + read_sum, scale, fill, zero = (module.get_kernel(name) for name in ("read_sum", "scale", "fill", "zero")) + config = LaunchConfig(grid=props.multiprocessor_count * 8, block=512) + one = LaunchConfig(grid=1, block=1) + allocations = [] + resources = [] + + def allocate(resource, size): + resources.append(resource) + buffer = resource.allocate(size, stream=stream) + allocations.append(buffer) + return buffer + + try: + result = allocate(DeviceMemoryResource(dev), 8) + host_result = allocate(PinnedMemoryResource(), 8) + + def run_once(ptr, location=None): + if location is not None: + managed.prefetch(location, stream=stream) + stream.sync() + launch(stream, one, zero, result) + stream.sync() + start = time.perf_counter() + launch(stream, config, read_sum, np.uint64(ptr), np.uint64(N), result) + stream.sync() + elapsed = time.perf_counter() - start + result.copy_to(host_result, stream=stream) + stream.sync() + actual = float(np.from_dlpack(host_result).view(np.float64)[0]) + if actual != N: + raise RuntimeError(f"read_sum returned {actual}, expected {N}") + return elapsed + + def time_read(ptr, label, location=None): + if location is None: + run_once(ptr) + durations = sorted(run_once(ptr, location) for _ in range(REPS)) + elapsed = (durations[REPS // 2 - 1] + durations[REPS // 2]) / 2 + print(f" {label:<52} {N * 4 / elapsed / 1e9:8.1f} GB/s ({elapsed * 1e3:.2f} ms)") + + print("[1] In-place kernel on pageable NumPy memory") + array = np.arange(1 << 20, dtype=np.float32) + expected = array * 3.0 + launch(stream, config, scale, np.uint64(array.ctypes.data), np.uint64(array.size), np.float32(3.0)) + stream.sync() + if not np.array_equal(array, expected): + raise RuntimeError("Pageable array result does not match") + print(" Result verified") + + print("[2] Read bandwidth, 1 GiB") + pinned = allocate(PinnedMemoryResource(PinnedMemoryResourceOptions(numa_id=props.host_numa_id)), N * 4) + stream.sync() + np.from_dlpack(pinned).view(np.float32)[:] = 1.0 + time_read(pinned.handle, f"pinned, host NUMA node {props.host_numa_id}") + pageable = np.ones(N, dtype=np.float32) + time_read(pageable.ctypes.data, "pageable NumPy memory") + device_buffer = allocate(DeviceMemoryResource(dev), N * 4) + launch(stream, config, fill, device_buffer, np.uint64(N), np.float32(1.0)) + time_read(device_buffer.handle, "device memory") + + print("[3] Managed memory, first kernel after each prefetch") + managed = allocate(ManagedMemoryResource(), N * 4) + stream.sync() + np.from_dlpack(managed).view(np.float32)[:] = 1.0 + time_read(managed.handle, f"managed, host NUMA node {props.host_numa_id}", Host(numa_id=props.host_numa_id)) + time_read(managed.handle, "managed, GPU", dev) + finally: + for buffer in reversed(allocations): + buffer.close(stream) + stream.sync() + print("Done") + return 0 + + +if __name__ == "__main__": + sys.exit(main()) diff --git a/samples/cuda_core/graceHopperMemory/requirements.txt b/samples/cuda_core/graceHopperMemory/requirements.txt new file mode 100644 index 00000000000..a6b9e43c674 --- /dev/null +++ b/samples/cuda_core/graceHopperMemory/requirements.txt @@ -0,0 +1,3 @@ +cuda-python>=13.0.0 +cuda-core>=1.3.0 +numpy>=2.2.5 From b71ede72d4124ce0225d833d2e7e79c2d4f8d029 Mon Sep 17 00:00:00 2001 From: 0z5a Date: Sun, 1 Nov 2026 23:14:58 +0100 Subject: [PATCH 2/3] samples: simplify Grace Hopper memory sample with CuPy Wrap buffers from the pinned (host NUMA), managed and device memory resources, plus a pageable NumPy array, as CuPy arrays and use standard CuPy ops instead of custom kernels. Drop the bandwidth timing and the NVML C2C query, and verify every result on the CPU. cp.from_dlpack() does not import pinned or CPU memory, so those two are wrapped with cupy.cuda.UnownedMemory. Update the script dependencies to cuda-core[cu13], CuPy and NumPy, and drop cuda-python. Signed-off-by: 0z5a --- samples/cuda_core/graceHopperMemory/README.md | 88 +++---- .../graceHopperMemory/graceHopperMemory.py | 235 ++++++++---------- .../graceHopperMemory/requirements.txt | 4 +- 3 files changed, 146 insertions(+), 181 deletions(-) diff --git a/samples/cuda_core/graceHopperMemory/README.md b/samples/cuda_core/graceHopperMemory/README.md index ac08b3b3a6e..d99a61f0d2b 100644 --- a/samples/cuda_core/graceHopperMemory/README.md +++ b/samples/cuda_core/graceHopperMemory/README.md @@ -1,55 +1,57 @@ # Sample: Grace Hopper Memory -This sample prints the CUDA NUMA and pageable-memory properties and the NVML -C2C status, then runs a kernel directly on a pageable NumPy array using the -host page tables. It compares reads from a host-NUMA pinned buffer, pageable -NumPy memory, device memory and managed memory after prefetch to either the -host NUMA node or the device. Every kernel result is checked. - -The read kernel sums 1 GiB of `float32` values using a grid-stride loop, one -shared-memory reduction and one atomic operation per block. The reported rate -is bytes read divided by host-clock time around a synchronized kernel launch, -the median of ten runs. It is not a tuned peak-bandwidth benchmark. Pageable -and pinned reads receive one untimed warmup. Before **each** managed-memory -measurement, the sample prefetches again to the indicated location and waits -for the prefetch to finish; the host-prefetched kernel time can include -migration to the GPU. +On Grace Hopper, NVLink-C2C makes CPU memory coherently accessible to the GPU, +and the GPU translates addresses through the host page tables. GPU code can +therefore work in place on pinned, managed and ordinary pageable host memory, +and the CPU sees the results without a copy. + +This sample allocates one buffer from each `cuda.core` memory resource, plus a +plain NumPy array: + +- `PinnedMemoryResource` on the GPU's nearest host NUMA node +- `ManagedMemoryResource` +- pageable memory from `np.empty()` +- `DeviceMemoryResource`, for comparison + +It wraps each buffer as a CuPy array and uses standard CuPy operations, with no +custom kernels. For each host-accessible buffer the CPU writes the input, CuPy +doubles it and sums it on the GPU, and the CPU checks the result in the same +memory. Device memory needs an explicit copy in each direction. + +`cp.from_dlpack()` imports the device and managed buffers. It does not import +pinned (`kDLCUDAHost`) or CPU (`kDLCPU`) memory, so the sample wraps those +pointers with `cupy.cuda.UnownedMemory`. ## Requirements -- A Grace Hopper system with CUDA pageable access through host page tables, - concurrent managed access and host NUMA memory pools. Systems without these - properties waive with exit code 2 (configurable via +- A Grace Hopper system: the GPU must access pageable memory through the host + page tables, support concurrent managed access and support host NUMA memory + pools. Other systems waive with exit code 2 (configurable via `CUDA_PYTHON_SAMPLE_WAIVER_EXIT_CODE`). -- CUDA Toolkit 13.x, `cuda-python`, `cuda-core` 1.3.0 or newer, and NumPy - 2.2.5 or newer. The required cuda-core APIs are present on the current main - branch; 1.3.0 may not yet be released. -- About 3 GiB of host memory, 1 GiB of device memory, and managed-memory - capacity. Allocate these only on a system where the GPU is available to you. - -Install from `requirements.txt` after the required cuda-core release is -available, or build the matching development version from the repository. +- `cuda-core` 1.3.0 or newer with CUDA 13, CuPy for CUDA 13 and NumPy 2.2.5 + or newer. ```bash cd samples/cuda_core/graceHopperMemory +pip install -r requirements.txt python graceHopperMemory.py ``` -## Historical probe measurements - -A separate prototype using the same kernel and input size was measured on one -GH200 480GB system (driver 595.58.03, CUDA 13.4.2, Python 3.12). Three runs, -each reporting the median of ten launches, gave the following ranges: - -| Memory | Read rate (GB/s) | -| --- | ---: | -| Pinned, host NUMA node 0 | 339–345 | -| Pageable NumPy array | 224–225 | -| Device memory | 2386–2417 | -| Managed, first kernel after host prefetch | 5.1–5.4 | -| Managed, after device prefetch | 2380–2393 | - -The host-prefetched managed result includes first-touch migration on that -system; it should not be interpreted as direct C2C bandwidth. The prototype -results are historical measurements, not results from this sample revision. -Rates depend on topology, system load, memory placement and kernel design. +## Example Output + +On a GH200 480GB: + +```text +Device: NVIDIA GH200 480GB, compute capability 9.0 + host_numa_id: 0 + pageable_memory_access_uses_host_page_tables: True + concurrent_managed_access: True + host_numa_memory_pools_supported: True +GPU work on host-resident memory: + pinned, host NUMA node 0 sum=17163091968 verified on the CPU without a copy + managed sum=17163091968 verified on the CPU without a copy + pageable NumPy array sum=17163091968 verified on the CPU without a copy +For comparison, device memory needs explicit copies: + device verified after a device-to-host copy +Done +``` diff --git a/samples/cuda_core/graceHopperMemory/graceHopperMemory.py b/samples/cuda_core/graceHopperMemory/graceHopperMemory.py index 636ce72ab9b..2cfb7c4f0e8 100644 --- a/samples/cuda_core/graceHopperMemory/graceHopperMemory.py +++ b/samples/cuda_core/graceHopperMemory/graceHopperMemory.py @@ -25,164 +25,127 @@ # OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. # /// script -# dependencies = ["cuda-python>=13.0.0", "cuda-core>=1.3.0", "numpy>=2.2.5"] +# dependencies = ["cuda-core[cu13]>=1.3.0", "cupy-cuda13x>=14.0.0", "numpy>=2.2.5"] # /// -"""Compare GPU reads from host, device, and managed memory on Grace Hopper.""" +"""Run CuPy directly on host memory on Grace Hopper. + +On Grace Hopper, NVLink-C2C makes CPU memory coherently accessible to the GPU, +and the GPU translates addresses through the host page tables. A kernel can +therefore read and write pinned, managed and even ordinary pageable host memory +in place, and the CPU sees the results without a copy. + +The sample allocates one buffer from each cuda.core memory resource (device, +pinned on the GPU's nearest host NUMA node, and managed) plus a plain NumPy +array. It wraps each one as a CuPy array, runs the same CuPy operations on all +of them and checks the results on the CPU. +""" import os import sys -import time - -import numpy as np - -from cuda.core import ( - Device, - DeviceMemoryResource, - Host, - LaunchConfig, - ManagedMemoryResource, - PinnedMemoryResource, - PinnedMemoryResourceOptions, - Program, - ProgramOptions, - launch, -) + +try: + import cupy as cp + import numpy as np + + from cuda.core import ( + Device, + DeviceMemoryResource, + ManagedMemoryResource, + PinnedMemoryResource, + PinnedMemoryResourceOptions, + ) +except ImportError as e: + print(f"Error: Required package not found: {e}") + print("Install with: pip install -r requirements.txt") + sys.exit(1) EXIT_WAIVED = int(os.environ.get("CUDA_PYTHON_SAMPLE_WAIVER_EXIT_CODE", "2")) -N = 1 << 28 # 1 GiB of float32 values -REPS = 10 - -SOURCE = r""" -extern "C" __global__ void read_sum(const float *x, unsigned long long n, double *out) { - __shared__ float part[512]; - float acc = 0.f; - for (unsigned long long i = blockIdx.x * (unsigned long long)blockDim.x + threadIdx.x; i < n; - i += (unsigned long long)gridDim.x * blockDim.x) - acc += x[i]; - part[threadIdx.x] = acc; - __syncthreads(); - for (unsigned s = blockDim.x / 2; s > 0; s >>= 1) { - if (threadIdx.x < s) part[threadIdx.x] += part[threadIdx.x + s]; - __syncthreads(); - } - if (threadIdx.x == 0) atomicAdd(out, (double)part[0]); -} -extern "C" __global__ void scale(float *x, unsigned long long n, float a) { - for (unsigned long long i = blockIdx.x * (unsigned long long)blockDim.x + threadIdx.x; i < n; - i += (unsigned long long)gridDim.x * blockDim.x) - x[i] *= a; -} -extern "C" __global__ void fill(float *x, unsigned long long n, float v) { - for (unsigned long long i = blockIdx.x * (unsigned long long)blockDim.x + threadIdx.x; i < n; - i += (unsigned long long)gridDim.x * blockDim.x) - x[i] = v; -} -extern "C" __global__ void zero(double *out) { *out = 0.0; } -""" +N = 1 << 24 # 64 MiB of float32 values per buffer + + +def host_as_cupy(ptr, nbytes, owner, device_id): + # cp.from_dlpack() does not import host memory (kDLCUDAHost or kDLCPU), so + # wrap the pointer directly. owner keeps the memory alive. + memory = cp.cuda.UnownedMemory(ptr, nbytes, owner, device_id=device_id) + return cp.ndarray((nbytes // 4,), cp.float32, cp.cuda.MemoryPointer(memory, 0)) + + +def run(label, gpu_array, cpu_array): + """Fill on the CPU, compute on the GPU, check on the CPU.""" + cpu_array[:] = np.arange(N, dtype=np.float32) % 1024 + gpu_array *= 2 + total = int(gpu_array.sum(dtype=cp.float64)) + cp.cuda.get_current_stream().synchronize() + + expected = 2 * (np.arange(N, dtype=np.float32) % 1024) + if total != int(expected.sum(dtype=np.float64)) or not np.array_equal(cpu_array, expected): + raise RuntimeError(f"{label}: result mismatch") + print(f" {label:<38} sum={total} verified on the CPU without a copy") def main(): dev = Device(0) dev.set_current() props = dev.properties - print(f"Device: {dev.name}, compute capability {dev.compute_capability}") - for name in ( - "numa_config", - "numa_id", - "host_numa_id", - "pageable_memory_access", - "pageable_memory_access_uses_host_page_tables", - "concurrent_managed_access", + cc = dev.compute_capability + print(f"Device: {dev.name}, compute capability {cc.major}.{cc.minor}") + print(f" host_numa_id: {props.host_numa_id}") + print(f" pageable_memory_access_uses_host_page_tables: {props.pageable_memory_access_uses_host_page_tables}") + print(f" concurrent_managed_access: {props.concurrent_managed_access}") + print(f" host_numa_memory_pools_supported: {props.host_numa_memory_pools_supported}") + + if not ( + props.pageable_memory_access_uses_host_page_tables + and props.concurrent_managed_access + and props.host_numa_memory_pools_supported + and props.host_numa_id >= 0 ): - print(f" {name}: {getattr(props, name)}") - print(f" system.Device.is_c2c_enabled: {dev.to_system_device().is_c2c_enabled}") - - if not props.pageable_memory_access_uses_host_page_tables: - print("Host page tables are unavailable for pageable GPU access; waiving sample.") - return EXIT_WAIVED - if ( - not props.concurrent_managed_access - or not props.host_memory_pools_supported - or not props.host_numa_memory_pools_supported - or props.host_numa_id < 0 - ): - print("Managed access, host memory pools, or host NUMA support is unavailable; waiving sample.") + print("This sample requires a Grace Hopper class system; waiving.") return EXIT_WAIVED stream = dev.create_stream() - arch = f"sm_{dev.compute_capability.major}{dev.compute_capability.minor}" - module = Program(SOURCE, code_type="c++", options=ProgramOptions(arch=arch)).compile("cubin") - read_sum, scale, fill, zero = (module.get_kernel(name) for name in ("read_sum", "scale", "fill", "zero")) - config = LaunchConfig(grid=props.multiprocessor_count * 8, block=512) - one = LaunchConfig(grid=1, block=1) - allocations = [] - resources = [] - - def allocate(resource, size): - resources.append(resource) - buffer = resource.allocate(size, stream=stream) - allocations.append(buffer) - return buffer + nbytes = N * 4 + device_buffer = DeviceMemoryResource(dev).allocate(nbytes, stream=stream) + pinned_buffer = PinnedMemoryResource(PinnedMemoryResourceOptions(numa_id=props.host_numa_id)).allocate( + nbytes, stream=stream + ) + managed_buffer = ManagedMemoryResource().allocate(nbytes, stream=stream) + pageable = np.empty(N, dtype=np.float32) + # Allocations are ordered on stream; CuPy uses its own stream. + stream.sync() try: - result = allocate(DeviceMemoryResource(dev), 8) - host_result = allocate(PinnedMemoryResource(), 8) - - def run_once(ptr, location=None): - if location is not None: - managed.prefetch(location, stream=stream) - stream.sync() - launch(stream, one, zero, result) - stream.sync() - start = time.perf_counter() - launch(stream, config, read_sum, np.uint64(ptr), np.uint64(N), result) - stream.sync() - elapsed = time.perf_counter() - start - result.copy_to(host_result, stream=stream) - stream.sync() - actual = float(np.from_dlpack(host_result).view(np.float64)[0]) - if actual != N: - raise RuntimeError(f"read_sum returned {actual}, expected {N}") - return elapsed - - def time_read(ptr, label, location=None): - if location is None: - run_once(ptr) - durations = sorted(run_once(ptr, location) for _ in range(REPS)) - elapsed = (durations[REPS // 2 - 1] + durations[REPS // 2]) / 2 - print(f" {label:<52} {N * 4 / elapsed / 1e9:8.1f} GB/s ({elapsed * 1e3:.2f} ms)") - - print("[1] In-place kernel on pageable NumPy memory") - array = np.arange(1 << 20, dtype=np.float32) - expected = array * 3.0 - launch(stream, config, scale, np.uint64(array.ctypes.data), np.uint64(array.size), np.float32(3.0)) - stream.sync() - if not np.array_equal(array, expected): - raise RuntimeError("Pageable array result does not match") - print(" Result verified") - - print("[2] Read bandwidth, 1 GiB") - pinned = allocate(PinnedMemoryResource(PinnedMemoryResourceOptions(numa_id=props.host_numa_id)), N * 4) - stream.sync() - np.from_dlpack(pinned).view(np.float32)[:] = 1.0 - time_read(pinned.handle, f"pinned, host NUMA node {props.host_numa_id}") - pageable = np.ones(N, dtype=np.float32) - time_read(pageable.ctypes.data, "pageable NumPy memory") - device_buffer = allocate(DeviceMemoryResource(dev), N * 4) - launch(stream, config, fill, device_buffer, np.uint64(N), np.float32(1.0)) - time_read(device_buffer.handle, "device memory") - - print("[3] Managed memory, first kernel after each prefetch") - managed = allocate(ManagedMemoryResource(), N * 4) - stream.sync() - np.from_dlpack(managed).view(np.float32)[:] = 1.0 - time_read(managed.handle, f"managed, host NUMA node {props.host_numa_id}", Host(numa_id=props.host_numa_id)) - time_read(managed.handle, "managed, GPU", dev) + print("GPU work on host-resident memory:") + run( + f"pinned, host NUMA node {props.host_numa_id}", + host_as_cupy(pinned_buffer.handle, nbytes, pinned_buffer, dev.device_id), + np.from_dlpack(pinned_buffer).view(np.float32), + ) + run( + "managed", + cp.from_dlpack(managed_buffer).view(cp.float32), + np.from_dlpack(managed_buffer).view(np.float32), + ) + run( + "pageable NumPy array", + host_as_cupy(pageable.ctypes.data, nbytes, pageable, dev.device_id), + pageable, + ) + + print("For comparison, device memory needs explicit copies:") + device_array = cp.from_dlpack(device_buffer).view(cp.float32) + device_array.set(np.arange(N, dtype=np.float32) % 1024) + device_array *= 2 + if not np.array_equal(device_array.get(), 2 * (np.arange(N, dtype=np.float32) % 1024)): + raise RuntimeError("device: result mismatch") + print(f" {'device':<38} verified after a device-to-host copy") finally: - for buffer in reversed(allocations): + cp.cuda.get_current_stream().synchronize() + for buffer in (managed_buffer, pinned_buffer, device_buffer): buffer.close(stream) stream.sync() + print("Done") return 0 diff --git a/samples/cuda_core/graceHopperMemory/requirements.txt b/samples/cuda_core/graceHopperMemory/requirements.txt index a6b9e43c674..e3b31cbbfec 100644 --- a/samples/cuda_core/graceHopperMemory/requirements.txt +++ b/samples/cuda_core/graceHopperMemory/requirements.txt @@ -1,3 +1,3 @@ -cuda-python>=13.0.0 -cuda-core>=1.3.0 +cuda-core[cu13]>=1.3.0 +cupy-cuda13x>=14.0.0 numpy>=2.2.5 From a7058a32e3541634cbc1600ceb5f5c654edb9693 Mon Sep 17 00:00:00 2001 From: 0z5a Date: Wed, 4 Nov 2026 15:52:09 +0100 Subject: [PATCH 3/3] fix(samples): use released cuda-core for Grace Hopper memory cuda-core 1.2.2 supplies the APIs used by this sample. Lower the minimum version consistently in the script metadata, requirements, and README; verify the clean requirements and sample on GH200. Signed-off-by: 0z5a --- samples/cuda_core/graceHopperMemory/README.md | 2 +- samples/cuda_core/graceHopperMemory/graceHopperMemory.py | 2 +- samples/cuda_core/graceHopperMemory/requirements.txt | 2 +- 3 files changed, 3 insertions(+), 3 deletions(-) diff --git a/samples/cuda_core/graceHopperMemory/README.md b/samples/cuda_core/graceHopperMemory/README.md index d99a61f0d2b..9bc8a4fcd21 100644 --- a/samples/cuda_core/graceHopperMemory/README.md +++ b/samples/cuda_core/graceHopperMemory/README.md @@ -28,7 +28,7 @@ pointers with `cupy.cuda.UnownedMemory`. page tables, support concurrent managed access and support host NUMA memory pools. Other systems waive with exit code 2 (configurable via `CUDA_PYTHON_SAMPLE_WAIVER_EXIT_CODE`). -- `cuda-core` 1.3.0 or newer with CUDA 13, CuPy for CUDA 13 and NumPy 2.2.5 +- `cuda-core` 1.2.2 or newer with CUDA 13, CuPy for CUDA 13 and NumPy 2.2.5 or newer. ```bash diff --git a/samples/cuda_core/graceHopperMemory/graceHopperMemory.py b/samples/cuda_core/graceHopperMemory/graceHopperMemory.py index 2cfb7c4f0e8..e1922989be5 100644 --- a/samples/cuda_core/graceHopperMemory/graceHopperMemory.py +++ b/samples/cuda_core/graceHopperMemory/graceHopperMemory.py @@ -25,7 +25,7 @@ # OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. # /// script -# dependencies = ["cuda-core[cu13]>=1.3.0", "cupy-cuda13x>=14.0.0", "numpy>=2.2.5"] +# dependencies = ["cuda-core[cu13]>=1.2.2", "cupy-cuda13x>=14.0.0", "numpy>=2.2.5"] # /// """Run CuPy directly on host memory on Grace Hopper. diff --git a/samples/cuda_core/graceHopperMemory/requirements.txt b/samples/cuda_core/graceHopperMemory/requirements.txt index e3b31cbbfec..73f17af1648 100644 --- a/samples/cuda_core/graceHopperMemory/requirements.txt +++ b/samples/cuda_core/graceHopperMemory/requirements.txt @@ -1,3 +1,3 @@ -cuda-core[cu13]>=1.3.0 +cuda-core[cu13]>=1.2.2 cupy-cuda13x>=14.0.0 numpy>=2.2.5