Skip to content

fix(cuda.core): grant the current device access to PinnedMemoryResource pools #3405

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
saiyambharara wants to merge 1 commit into
base: main
Choose a base branch
from
Open

fix(cuda.core): grant the current device access to PinnedMemoryResource pools #3405

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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: 7 additions & 1 deletion cuda_core/cuda/core/_memory/_pinned_memory_resource.pyi
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -73,6 +73,12 @@ class PinnedMemoryResource(_MemPool):
is unsupported and stream-ordered allocation is not needed, use
:class:`LegacyPinnedMemoryResource`.

The driver makes a host memory pool accessible only from the host by
default, so the resource grants the device that is current when it is
created read and write access to the pool. Kernels on other devices cannot
access its buffers. Without options, the grant applies to the driver's
current pool, which the whole process shares.

To create an IPC-Enabled memory resource (MR) that is capable of sharing
allocations between processes, specify ``ipc_enabled=True`` in the initializer
option. When IPC is enabled and ``numa_id`` is not specified, the NUMA node
Expand Down Expand Up @@ -167,7 +173,7 @@ class PinnedMemoryResource(_MemPool):
"""The host NUMA node ID used for pool placement, or -1 for OS-managed placement."""
@property
def is_device_accessible(self) -> bool:
"""Return True. This memory resource provides device-accessible buffers."""
"""Return True. The device that is current when the resource is created can read and write its buffers."""
@property
def is_host_accessible(self) -> bool:
"""Return True. This memory resource provides host-accessible buffers."""
Expand Down
38 changes: 30 additions & 8 deletions cuda_core/cuda/core/_memory/_pinned_memory_resource.pyx
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -15,6 +15,8 @@ from cuda.core._memory._memory_pool cimport (
)
from cuda.core._memory cimport _ipc
from cuda.core._memory._ipc cimport IPCAllocationHandle
from cuda.core._memory._location cimport cumemlocation_from_id
from cuda.core._rt cimport as_cu
from cuda.core._stream cimport Stream, Stream_accept
from cuda.core._utils.cuda_utils cimport (
check_or_create_options,
Expand Down Expand Up @@ -99,6 +101,12 @@ cdef class PinnedMemoryResource(_MemPool):
is unsupported and stream-ordered allocation is not needed, use
:class:`LegacyPinnedMemoryResource`.

The driver makes a host memory pool accessible only from the host by
default, so the resource grants the device that is current when it is
created read and write access to the pool. Kernels on other devices cannot
access its buffers. Without options, the grant applies to the driver's
current pool, which the whole process shares.
Comment on lines +106 to +108

@coderabbitai coderabbitai Bot Nov 6, 2026

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Minor · Functional Correctness · Quick win

Supported by static analysis

Script executed:

git diff 63efb2db31f06ce4c1edc0c3fe5bcd78d8cb720a 036e4a07c8ec8d9fb32630e2d581f886f2df61a5 -- cuda_core/cuda/core/_memory/_pinned_memory_resource.pyx cuda_core/cuda/core/_memory/_pinned_memory_resource.pyi cuda_core/cuda/core/_memory/_memory_pool.pyx
rg -n -F -e 'cuMemGetMemPool' -e 'cuMemPoolSetAccess' -e 'cuMemPoolGetAccess' -e 'Kernels on other devices' -e 'current pool' cuda_core/cuda/core/_memory cuda_core/tests
nl -ba cuda_core/cuda/core/_memory/_pinned_memory_resource.pyx | sed -n '65,210p;260,355p'
rg -n -F -e '_initialize' -e 'cuMemGetMemPool' -e 'cuMemPoolCreate' cuda_core/cuda/core/_memory/_memory_pool.pyx

Repository: NVIDIA/cuda-python

Length of output: 20733

Web query:

CUDA driver API cuMemGetMemPool host location current pool CU_MEM_LOCATION_TYPE_HOST cuMemPoolSetAccess other device access flags independent

Result:

Sources consulted:
- 6.39. Stream Ordered Memory Allocator — CUDA Driver API Reference Manual 13.5 documentation — https://docs.nvidia.com/cuda/cuda-driver-api/cuda_driver_api/group__CUDA__MALLOC__ASYNC.html
- 4.3. Stream-Ordered Memory Allocator — CUDA Programming Guide — https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/stream-ordered-memory-allocation.html
- CUDA Driver API — https://docs.nvidia.com/cuda/pdf/CUDA_Driver_API.pdf
- 6.5. Data types used by CUDA driver — CUDA Driver API Reference Manual 13.5 documentation — https://docs.nvidia.com/cuda/cuda-driver-api/cuda_driver_api/group__CUDA__TYPES.html
- …and 28 more

Script executed:

nl -ba cuda_core/cuda/core/_memory/_memory_pool.pyx | sed -n '400,435p'
nl -ba cuda_core/cuda/core/_memory/_location.pyx | sed -n '1,120p'

Repository: NVIDIA/cuda-python

Length of output: 1675

Web query:

CUDA 13 driver cuMemPoolSetAccess "specified" locations "access" map allocations does not affect other locations

Result:

Sources consulted:
- CUDA Driver API :: CUDA Toolkit Documentation — https://docs.nvidia.com/cuda/archive/13.0.0/cuda-driver-api/group__CUDA__MALLOC__ASYNC.html
- 6.39. Stream Ordered Memory Allocator — CUDA Driver API Reference Manual 13.5 documentation — https://docs.nvidia.com/cuda/cuda-driver-api/cuda_driver_api/group__CUDA__MALLOC__ASYNC.html
- 4.3. Stream-Ordered Memory Allocator — CUDA Programming Guide — https://docs.nvidia.com/cuda/archive/13.2.1/cuda-programming-guide/04-special-topics/stream-ordered-memory-allocation.html
- …and 17 more

Script executed:

rg -n -F 'cumemlocation_from_id' cuda_core/cuda/core/_memory/_location.pxd cuda_core/cuda/core/_memory
nl -ba cuda_core/cuda/core/_memory/_location.pxd | sed -n '1,110p'
nl -ba cuda_core/cuda/core/_memory/_peer_access_utils.pyx | sed -n '135,175p'

Repository: NVIDIA/cuda-python

Length of output: 7737

Web query:

site:docs.nvidia.com/cuda/archive/13.0.0/cuda-driver-api/group__CUDA__MALLOC__ASYNC.html "cuMemPoolSetAccess" "The access" "locations"

Result:

Sources consulted:
- CUDA Driver API :: CUDA Toolkit Documentation — https://docs.nvidia.com/cuda/archive/13.0.0/cuda-driver-api/group__CUDA__MALLOC__ASYNC.html
- CUDA Driver API :: CUDA Toolkit Documentation — https://docs.nvidia.com/cuda/archive/13.0.0/cuda-driver-api/group__CUDA__VA.html
- CUDA Driver API — https://docs.nvidia.com/cuda/archive/13.0.0/pdf/CUDA_Driver_API.pdf
- CUDA Driver API :: CUDA Toolkit Documentation — https://docs.nvidia.com/cuda/archive/13.0.2/cuda-driver-api/group__CUDA__MALLOC__ASYNC.html
- …and 19 more

important: Qualify the other-device access claim.

Two resources created without options can share the current pool. If they are created on different supported devices, both devices can retain access to all allocations in that pool. Qualify the claim here and in _pinned_memory_resource.pyi.


To create an IPC-Enabled memory resource (MR) that is capable of sharing
allocations between processes, specify ``ipc_enabled=True`` in the initializer
option. When IPC is enabled and ``numa_id`` is not specified, the NUMA node
Expand Down Expand Up @@ -141,13 +149,7 @@ cdef class PinnedMemoryResource(_MemPool):
raise TypeError("Cannot allocate from a mapped IPC-enabled memory resource")
cdef Stream s = Stream_accept(stream)
device = s.device
cdef bint supported = (
device.properties.host_numa_memory_pools_supported
if self._numa_id >= 0
else device.properties.host_memory_pools_supported
)

if not supported:
if not _host_pool_supported(device, self._numa_id):
raise RuntimeError(
f"CUDA device {device.device_id} does not support the requested "
"host memory pool for PinnedMemoryResource. Use "
Expand Down Expand Up @@ -244,7 +246,7 @@ cdef class PinnedMemoryResource(_MemPool):

@property
def is_device_accessible(self) -> bool:
"""Return True. This memory resource provides device-accessible buffers."""
"""Return True. The device that is current when the resource is created can read and write its buffers."""

@coderabbitai coderabbitai Bot Nov 6, 2026

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Minor · Functional Correctness · Quick win

Supported by static analysis

Script executed:

git diff 63efb2db31f06ce4c1edc0c3fe5bcd78d8cb720a 036e4a07c8ec8d9fb32630e2d581f886f2df61a5 -- cuda_core/cuda/core/_memory/_pinned_memory_resource.pyx
rg -n -F -e 'from_allocation_handle' -e '_PMR_init' -e 'is_device_accessible' -e 'cuMemPoolImport' -e 'cuMemPoolSetAccess' cuda_core/cuda/core/_memory cuda_core/docs cuda_core/tests
nl -ba cuda_core/cuda/core/_memory/_pinned_memory_resource.pyx | sed -n '1,410p'

Repository: NVIDIA/cuda-python

Length of output: 33079

Web query:

CUDA driver cuMemPoolImportFromShareableHandle imported memory pool device access cuMemPoolSetAccess access permissions importing process

Result:

Sources consulted:
- CUDA Programming Guide — https://docs.nvidia.com/cuda/cuda-programming-guide/pdf/cuda-programming-guide.pdf
- 6.39. Stream Ordered Memory Allocator — CUDA Driver API Reference Manual 13.5 documentation — https://docs.nvidia.com/cuda/cuda-driver-api/cuda_driver_api/group__CUDA__MALLOC__ASYNC.html
- 4.3. Stream-Ordered Memory Allocator — CUDA Programming Guide — https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/stream-ordered-memory-allocation.html
- cuda.bindings._v2.driver.mem_pool_import_from_shareable_handle — cuda.bindings — https://nvidia.github.io/cuda-python/cuda-bindings/latest/module/generated/driver_v2/cuda.bindings._v2.driver.mem_pool_import_from_shareable_handle.html
- 4.3. Stream-Ordered Memory Allocator — CUDA Programming Guide — https://docs.nvidia.com/cuda/archive/13.1.1/cuda-programming-guide/04-special-topics/stream-ordered-memory-allocation.html
- …and 14 more

Script executed:

nl -ba cuda_core/cuda/core/_memory/_ipc.pyx | sed -n '380,485p'
nl -ba cuda_core/cuda/core/_memory/_pinned_memory_resource.pyi | sed -n '60,190p'
git diff 63efb2db31f06ce4c1edc0c3fe5bcd78d8cb720a 036e4a07c8ec8d9fb32630e2d581f886f2df61a5 -- cuda_core/docs/source/release
rg -n -F -e 'memory_pool_import' -e 'mem_pool_import' -e 'cuMemPoolImportFromShareableHandle' cuda_core/cuda/core/_cpp cuda_core/cuda/core/_memory

Repository: NVIDIA/cuda-python

Length of output: 12111

Web query:

CUDA cuMemPoolCreate HOST_NUMA pool initially only accessible host cuMemPoolImportFromShareableHandle access imported host memory pool GPU access permissions

Result:

Sources consulted:
- 6.39. Stream Ordered Memory Allocator — CUDA Driver API Reference Manual 13.5 documentation — https://docs.nvidia.com/cuda/cuda-driver-api/cuda_driver_api/group__CUDA__MALLOC__ASYNC.html
- 6.38. Stream Ordered Memory Allocator — CUDA Driver API Reference Manual 13.4 documentation — https://docs.nvidia.com/cuda/archive/13.4.2/cuda-driver-api/cuda_driver_api/group__CUDA__MALLOC__ASYNC.html
- 4.3. Stream-Ordered Memory Allocator — CUDA Programming Guide — https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/stream-ordered-memory-allocation.html
- …and 19 more

Script executed:

nl -ba cuda_core/cuda/core/_cpp/rt/memory.cpp | sed -n '108,175p'
rg -n -F -e 'create_mempool_handle_ipc' cuda_core/cuda/core/_cpp/rt cuda_core/cuda/core/_memory/_ipc.pyx
nl -ba cuda_core/docs/source/release/1.3.0-notes.rst | sed -n '138,155p'

Repository: NVIDIA/cuda-python

Length of output: 4542

important: Qualify the device-access claims for imported pools.

from_allocation_handle() does not grant device access, and imported pools do not inherit the exporter's access grants. A kernel that uses an imported allocation can fault.

Exclude imported pools from the guarantee in the property, class documentation, stub, and release notes.

return True

@property
Expand All @@ -253,6 +255,14 @@ cdef class PinnedMemoryResource(_MemPool):
return True


cdef inline bint _host_pool_supported(device, int numa_id):
return (
device.properties.host_numa_memory_pools_supported
if numa_id >= 0
else device.properties.host_memory_pools_supported
)


cdef inline _PMR_init(PinnedMemoryResource self, options):
from .._device import Device

Expand All @@ -264,6 +274,7 @@ cdef inline _PMR_init(PinnedMemoryResource self, options):
cdef size_t max_size = 0
cdef cydriver.CUmemLocationType loc_type
cdef int numa_id = -1
cdef cydriver.CUmemAccessDesc access

if opts is not None:
ipc_enabled = opts.ipc_enabled
Expand Down Expand Up @@ -309,6 +320,17 @@ cdef inline _PMR_init(PinnedMemoryResource self, options):
max_size,
)

# A pool on a host location is accessible only from the host by default
# (see cuMemPoolCreate), so grant the current device read/write access. A
# device without host pool support cannot allocate from it anyway.
dev = Device()
if _host_pool_supported(dev, numa_id):
access.location = cumemlocation_from_id(
cydriver.CUmemLocationType.CU_MEM_LOCATION_TYPE_DEVICE, dev.device_id)
access.flags = cydriver.CUmemAccess_flags.CU_MEM_ACCESS_FLAGS_PROT_READWRITE
with nogil:
HANDLE_RETURN(cydriver.cuMemPoolSetAccess(as_cu(self._h_pool), &access, 1))


def _deep_reduce_pinned_memory_resource(mr: object) -> tuple[object, ...]:
check_multiprocessing_start_method()
Expand Down
7 changes: 7 additions & 0 deletions cuda_core/docs/source/release/1.3.0-notes.rst
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -142,6 +142,13 @@ New features
Fixes and enhancements
----------------------

- Kernels can read and write :class:`PinnedMemoryResource` buffers. The driver
makes a host memory pool accessible only from the host by default, so a
kernel that touched such a buffer failed with ``CUDA_ERROR_ILLEGAL_ADDRESS``.
The resource now grants the device that is current when it is created read
and write access to the pool.
(`#3205 <https://github.com/NVIDIA/cuda-python/issues/3205>`__)

- ``register()`` on a :class:`~_memory.DeviceMemoryResource` or a
:class:`PinnedMemoryResource` now raises ``ValueError`` when the resource
already has a different UUID. The check used to be an ``assert``, so under
Expand Down
45 changes: 45 additions & 0 deletions cuda_core/tests/test_memory.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -51,17 +51,21 @@
DeviceMemoryResource,
DeviceMemoryResourceOptions,
GraphMemoryResource,
LaunchConfig,
LegacyPinnedMemoryResource,
ManagedBuffer,
ManagedMemoryResource,
ManagedMemoryResourceOptions,
MemoryResource,
PinnedMemoryResource,
PinnedMemoryResourceOptions,
Program,
ProgramOptions,
StreamOptions,
VirtualMemoryBuffer,
VirtualMemoryResource,
VirtualMemoryResourceOptions,
launch,
)
from cuda.core._dlpack import DLDeviceType
from cuda.core._memory._ipc import IPCBufferDescriptor
Expand Down Expand Up @@ -1242,6 +1246,47 @@ def test_pinned_memory_resource_initialization(init_cuda):
buffer.close()


@pytest.mark.human_reviewed
@pytest.mark.parametrize("pool", ["current", "new", "new_numa"])
def test_pinned_memory_resource_kernel_access(init_cuda, pool):
"""A kernel on the current device reads and writes PinnedMemoryResource buffers (#3205)."""
device = init_cuda
skip_if_pinned_memory_unsupported(device)
if pool == "current":
options = None
elif pool == "new":
options = PinnedMemoryResourceOptions(max_size=POOL_SIZE)
else:
numa_id = device.properties.host_numa_id
if numa_id < 0 or not device.properties.host_numa_memory_pools_supported:
pytest.skip("Device does not support host NUMA memory pools")
options = PinnedMemoryResourceOptions(max_size=POOL_SIZE, numa_id=numa_id)

code = r"""
extern "C" __global__ void double_plus_one(int* data, size_t n) {
size_t i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) data[i] = data[i] * 2 + 1;
}
"""
kernel = (
Program(code, "c++", ProgramOptions(arch=f"sm_{device.arch}")).compile("cubin").get_kernel("double_plus_one")
)

n = 1024
mr = create_pinned_memory_resource_or_xfail(options, xfail_device=device)
buffer = _allocate_pinned_buffer_or_xfail(mr, n * ctypes.sizeof(ctypes.c_int), device=device)
stream = device.default_stream
stream.sync()
data = (ctypes.c_int * n).from_address(int(buffer.handle))
data[:] = range(n)

launch(stream, LaunchConfig(grid=(n + 255) // 256, block=256), kernel, buffer, ctypes.c_size_t(n))
stream.sync()
assert data[:] == [2 * i + 1 for i in range(n)]
buffer.close()
mr.close()


@pytest.mark.agent_authored(model="cursor-grok-4.5")
def test_pinned_memory_resource_rejects_unsupported_host_pool(init_cuda):
"""allocate() must fail on devices without host memory pool support (see #2486)."""
Expand Down
Loading