1+ # Copyright (c) 2025, NVIDIA CORPORATION & AFFILIATES. ALL RIGHTS RESERVED.
2+ #
3+ # SPDX-License-Identifier: Apache-2.0
4+
15import cupy as cp
2- import numpy as np
6+
37from cuda .core .experimental import (
4- Device , LaunchConfig , Program , ProgramOptions , launch ,
5- DeviceMemoryResource , LegacyPinnedMemoryResource , Buffer
8+ Device ,
9+ DeviceMemoryResource ,
10+ LaunchConfig ,
11+ LegacyPinnedMemoryResource ,
12+ Program ,
13+ ProgramOptions ,
14+ launch ,
615)
7- from cuda .core .experimental ._memory import MemoryResource
8- from cuda .core .experimental ._utils .cuda_utils import handle_return
9- from cuda .bindings import driver
16+ from cuda .core .experimental ._dlpack import DLDeviceType
1017
1118# Kernel for memory operations
1219code = """
1320extern "C"
14- __global__ void memory_ops(float* device_data,
21+ __global__ void memory_ops(float* device_data,
1522 float* pinned_data,
1623 size_t N) {
1724 const unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
1825 if (tid < N) {
1926 // Access device memory
2027 device_data[tid] = device_data[tid] + 1.0f;
21-
28+
2229 // Access pinned memory (zero-copy from GPU)
2330 pinned_data[tid] = pinned_data[tid] * 3.0f;
2431 }
4956# 1. Device Memory (GPU-only)
5057device_buffer = device_mr .allocate (total_size , stream = stream )
5158device_array = cp .ndarray (
52- size , dtype = dtype ,
59+ size ,
60+ dtype = dtype ,
5361 memptr = cp .cuda .MemoryPointer (
5462 cp .cuda .UnownedMemory (int (device_buffer .handle ), device_buffer .size , device_buffer ), 0
55- )
63+ ),
5664)
5765
5866# 2. Pinned Memory (CPU memory, GPU accessible)
5967pinned_buffer = pinned_mr .allocate (total_size , stream = stream )
6068pinned_array = cp .ndarray (
61- size , dtype = dtype ,
69+ size ,
70+ dtype = dtype ,
6271 memptr = cp .cuda .MemoryPointer (
6372 cp .cuda .UnownedMemory (int (pinned_buffer .handle ), pinned_buffer .size , pinned_buffer ), 0
64- )
73+ ),
6574)
6675
6776# Initialize data
8190grid = (size + block - 1 ) // block
8291config = LaunchConfig (grid = grid , block = block )
8392
84- launch (stream , config , kernel ,
85- device_buffer , pinned_buffer , cp .uint64 (size ))
93+ launch (stream , config , kernel , device_buffer , pinned_buffer , cp .uint64 (size ))
8694stream .sync ()
8795
8896# Verify kernel operations
113121# Create a new device buffer and copy from pinned
114122new_device_buffer = device_mr .allocate (total_size , stream = stream )
115123new_device_array = cp .ndarray (
116- size , dtype = dtype ,
124+ size ,
125+ dtype = dtype ,
117126 memptr = cp .cuda .MemoryPointer (
118127 cp .cuda .UnownedMemory (int (new_device_buffer .handle ), new_device_buffer .size , new_device_buffer ), 0
119- )
128+ ),
120129)
121130
122131pinned_buffer .copy_to (new_device_buffer , stream = stream )
131140print (f"Pinned buffer DLPack device: { pinned_buffer .__dlpack_device__ ()} " )
132141
133142# Assert DLPack device types
134- from cuda .core .experimental ._memory import DLDeviceType
135-
136143device_dlpack = device_buffer .__dlpack_device__ ()
137144pinned_dlpack = pinned_buffer .__dlpack_device__ ()
138145
142149# Test buffer size properties
143150assert device_buffer .size == total_size , f"Device buffer size mismatch: expected { total_size } , got { device_buffer .size } "
144151assert pinned_buffer .size == total_size , f"Pinned buffer size mismatch: expected { total_size } , got { pinned_buffer .size } "
145- assert new_device_buffer .size == total_size , f"New device buffer size mismatch: expected { total_size } , got { new_device_buffer .size } "
152+ assert new_device_buffer .size == total_size , (
153+ f"New device buffer size mismatch: expected { total_size } , got { new_device_buffer .size } "
154+ )
146155
147156# Test memory resource properties
148157assert device_buffer .memory_resource == device_mr , "Device buffer should use device memory resource"
160169assert pinned_buffer .handle == 0 , "Pinned buffer should be closed"
161170assert new_device_buffer .handle == 0 , "New device buffer should be closed"
162171
163- print ("Memory management example completed!" )
172+ print ("Memory management example completed!" )
0 commit comments