You signed in with another tab or window. Reload to refresh your session.You signed out in another tab or window. Reload to refresh your session.You switched accounts on another tab or window. Reload to refresh your session.Dismiss alert
{{ message }}
Repository navigation
feat(cuda.core): share VirtualMemoryResource buffers with fabric handles - #3230
Buffers of a VirtualMemoryResource configured with handle_type="fabric" can now be shared between processes on one node. This adds the fabric transport that #2980 tracks, beside the POSIX fd transport from #3028.
Changes
New VirtualMemoryFabricBufferDescriptor (listed in api_private.rst): one opaque 64-byte CUmemFabricHandle per physical allocation, the allocation sizes, the buffer size, and the exporting device's UUID. It is plain data: pickle and multiprocessing carry it without the file-descriptor machinery, it compares and hashes by value, and from_handles() rebuilds one received through a transport of your own, with validation. The export zeroes the handle struct first, so two exports of one allocation carry the same bytes and no stack bytes travel in a descriptor.
Both VMM descriptors derive from a private _VirtualMemoryDescriptor, so Buffer.from_ipc_descriptor and the resource's import run one loop for both transports. The fd descriptor keeps its ownership and identity semantics and its from_fds().
VirtualMemoryResource construction checks Device.properties.handle_type_fabric_supported for handle_type="fabric" and raises RuntimeError that names the NVLink fabric and IMEX channel requirements, before an allocation can fail opaquely. A device can report support inside a container that has the IMEX channel device node while the GPU has no fabric; a refused allocation then carries a note that says so. is_ipc_enabled is True for "fabric" on Linux with device-located memory.
Fabric descriptors pin nothing. The exporter keeps the buffer alive until every importer has imported it; after import, the imported allocation holds its own driver references. A stale or forged handle fails at import with a driver error and a note about stale handles, forgery, and IMEX channels. The docstrings, the release note, and VMM_DESIGN.md say so wherever descriptor lifetime is described.
A location mismatch at import names the owning device by UUID and says which local device it is, or that it is not visible in this process.
_ipc.pyi and _virtual_memory_resource.pyi are regenerated by the stub hook.
Tests: tests/memory_ipc/test_vmm_fabric.py covers export shape, descriptor validation and pickling, Queue, Pipe, Process-argument, and Pool round trips for descriptors and buffers, a second-GPU import with peer access, handle-type mismatch, garbage and stale handles, re-export of imports, a queued buffer that does not pin, and the unsupported-device error. The fabric fixtures skip on devices without fabric handle support, and on devices that report it but refuse the allocation (one probe allocation per device).
Validation
Fabric handles need an NVLink fabric and an accessible IMEX channel. The project's CI has no such runner, so the fabric cases skip there; the release note says that this release's fabric support is not covered by CI. On the CI runners the attribute is true (the container has the channel device node) while every fabric allocation is refused with a driver error, which is why the fixtures probe instead of trusting the attribute.
Validation on NVIDIA internal hardware: a two-GPU B300 system ran the memory_ipc suite and test_memory.py twice: on the first commit (478 passed, 21 skipped; the one failure found the uninitialized padding in the exported handle struct) and on the fix and on the PR head with main merged (479 passed, 21 skipped each), and a Hopper system without fabric support ran the fabric and fd suites (30 passed, 17 skipped, with the unsupported-device error test among the passes). Multi-node sharing is out of scope and untested.
Related Work
Part of #2980. Follows #3028 and #3229; this PR includes the #3229 commit until that PR merges.
…alMemoryFdBufferDescriptor
The class carries POSIX file descriptors for VirtualMemoryResource
buffers. A fabric-handle descriptor will sit beside it, so its name now
says which transport it is. No alias: the class first appears in the
unreleased 1.3.0. No behavior change.
Written by Codex (gpt-6.1-sol); reviewed by Claude.
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Add VirtualMemoryFabricBufferDescriptor, a plain-data descriptor that
carries one opaque 64-byte CUmemFabricHandle per physical allocation,
the allocation sizes, the buffer size, and the exporting device's UUID.
VirtualMemoryResource exports it when config.handle_type is "fabric",
and Buffer.from_ipc_descriptor imports it through the shared
_VirtualMemoryDescriptor interface that the POSIX fd descriptor also
implements, so both transports run one import loop.
Fabric descriptors travel through plain pickle and multiprocessing
without the file-descriptor machinery and pin nothing: the exporter
keeps the buffer alive until every importer has imported it. A stale or
forged handle fails at import with a driver error and a note. Resource
construction checks Device.properties.handle_type_fabric_supported and
raises RuntimeError when the device lacks an NVLink fabric or an
accessible IMEX channel. from_handles() rebuilds a descriptor received
through a transport of your own. A location mismatch at import names
the owning device by UUID.
Sharing is supported on one node; multi-node sharing is untested. The
project's CI has no fabric-capable runner; the release note says so.
Part of NVIDIA#2980.
Written by Codex (gpt-6.1-sol); reviewed by Claude.
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Linux virtual-memory sharing now supports fabric handles for device memory alongside POSIX file descriptors.
Fabric descriptors can be pickled and transferred between processes on the same node. They do not pin allocations, so keep the exporting buffer alive until import completes. Imported buffers cannot be re-exported.
Fabric sharing requires a supported device, NVLink fabric, and accessible IMEX channel. Unsupported devices and invalid or incompatible imports raise errors.
API Updates
POSIX file-descriptor descriptors are now named VirtualMemoryFdBufferDescriptor; fabric descriptors use VirtualMemoryFabricBufferDescriptor.
Linux virtual-memory IPC now supports fabric handles alongside POSIX file descriptors. The change adds fabric descriptors and resource export and import paths, defines fabric-handle validation and lifetime behavior, and adds tests for transfer, errors, and unsupported devices.
Changes: Linux VMM Fabric IPC
Layer / File(s)
Summary
Shared and fabric descriptor model cuda_core/cuda/core/_memory/_ipc.pxd, cuda_core/cuda/core/_memory/_ipc.pyi, cuda_core/cuda/core/_memory/_ipc.pyx
A shared descriptor base supports handle types and allocation sizes. VirtualMemoryFdBufferDescriptor replaces the previous POSIX descriptor name. VirtualMemoryFabricBufferDescriptor validates and serializes fabric handles, sizes, and an optional owner UUID.
Resource export and import cuda_core/cuda/core/_memory/_virtual_memory_resource.pyx, cuda_core/cuda/core/_memory/_virtual_memory_resource.pyi, cuda_core/cuda/core/_cpp/rt/VMM_DESIGN.md, cuda_core/docs/source/api_private.rst, cuda_core/docs/source/release/1.3.0-notes.rst
Linux resource handling recognizes both handle types. Fabric export returns a fabric descriptor; import uses the shared descriptor interface and reports fabric-specific errors. Fabric descriptors do not pin allocations, and the exporter must keep the buffer alive until import completes. Resource construction raises RuntimeError when the device lacks fabric-handle support.
New tests cover descriptor construction, serialization, cross-process sharing, ownership, lifetime, import errors, and unsupported devices. Existing POSIX-FD tests use the renamed descriptor.
Possibly related issues
NVIDIA/cuda-python#2980: The change adds fabric descriptor creation, import paths, and resource checks that address the fabric-support gap described by the issue.
NVIDIA/cuda-python#2077: The change adds fabric tests that skip when required hardware is unavailable; it does not add a fabric-capable CI runner.
Suggested reviewers:rwgk, mdboom
Priority: Normal
Change: Feature
Merge Risk: Moderate · up to 7dc24
A real fabric allocation-path failure can be misreported as unavailable, leaving the new fabric IPC behavior untested on an affected device. Restrict skips to physical-allocation refusals before merging.
- [ ] Autofix · Keep fixing CodeRabbit findings and required CI, and resolving merge conflicts
Comment @coderabbitai help to get the list of available commands.
Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 10 remain after this review.
Zero the 64-byte handle struct before a fabric export. The driver fills
its fields and leaves the padding, so two exports of one buffer carried
different stack bytes (seen on B300); the export now returns the same
bytes for the same allocation, and no stack bytes travel in a descriptor.
The fabric test fixtures allocate once per device after the attribute
check. Inside a container that has the IMEX channel device node the
attribute is true while the GPU has no fabric, and the allocation is
refused; that is what the CI runners do. A refused physical allocation
on a fabric resource now carries a note about the fabric and IMEX
requirements, and the docs say that the attribute can be true where the
allocation still fails.
The import-error tests check cuda.core's note, not the driver's status,
which varies with the IMEX setup. Garbage handles use fixed patterns
instead of os.urandom. The release note names the POSIX fd descriptor
as the one that pins. One note per import failure. _handles is readonly
on both descriptors.
Written by Claude after the de novo review, CodeRabbit, CI, and the
Hopper and B300 runs.
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 10 remain after this review.
The reason will be displayed to describe this comment to others. Learn more.
Major · Functional Correctness · Quick win
important: Restrict which allocation errors cause a skip.
If the probe raises CUDA_ERROR_OUT_OF_MEMORY, this catch caches the error and skips every later fabric test for that device. Let unrelated CUDAError failures propagate. Skip only statuses that establish fabric is unavailable. CUDA lists out-of-memory as a possible allocation error. (docs.nvidia.com)
As per coding guidelines, “Capability probes must not swallow real bugs.”
The reason will be displayed to describe this comment to others. Learn more.
Minor · Functional Correctness · Quick win
Supported by static analysis
Web query:
cuMemExportToShareableHandle CU_MEM_HANDLE_TYPE_FABRIC repeated export same handle bytes deterministic documentation
Result:
<source_evidence>
<source kind="open_page">
<title>6.47. Virtual Memory Management — CUDA Driver API Reference Manual 13.4 documentation</title>
<location>https://docs.nvidia.com/cuda/archive/13.4.2/cuda-driver-api/cuda_driver_api/group__CUDA__VA.html</location>
<excerpt>citeturn2view0 [wordlim: 200] Crawled: today; Content type: text/html; Source: open({"ref_id":"turn1search6","lineno":null}); Total lines: 519
L0: [Input: Search docs] [Input] [Input]
L1:
L2: 【0†CUDA Driver API Reference Manual】
L3: # 6.47. Virtual Memory Management【1†】
L4:
L5: This section describes the virtual memory management functions of the low-level CUDA driver application programming interface.
L6:
L7: Functions
L8:
L9: CUresult 【2†cuMemAddressFree】(CUdeviceptr ptr, size_t size)
L10:
L11:
L12: Free an address range reservation.
L13:
L14: CUresult 【3†cuMemAddressReserve】(CUdeviceptr *ptr, size_t size, size_t alignment, CUdeviceptr addr, unsigned long long flags)
L15:
L16:
L17: Allocate an address range reservation.
L18: CUresult 【4†cuMemCreate】(CUmemGenericAllocationHandle *handle, size_t size, const CUmemAllocationProp *prop, unsigned long long flags)
L19:
L20:
L21: Create a CUDA memory handle representing a memory allocation of a given size described by the given properties.
L22:
L23: CUresult 【5†cuMemExportToShareableHandle】(void *shareableHandle, CUmemGenericAllocationHandle handle, CUmemAllocationHandleType handleType, unsigned long long flags)
L24:
L25:
L26: Exports an allocation to a requested shareable handle type.
L27: CUresult 【6†cuMemGetAccess】(unsigned long long *flags, const CUmemLocation *location, CUdeviceptr ptr)
L28:
L29:
L30: Get the access `flags` set for the given `location` and `ptr` .
L31:
L32: CUresult 【7†cuMemGetAllocationGranularity】(size_t *granularity, const CUmemAllocationProp *prop, CUmemAllocationGranularity_flags option)
L33:
L34:
L35: Calculates either the minimal or recommended granularity.
L36:
L37: CUresult 【8†cuMemGetAllocationPropertiesFromHandle】(CUmemAllocationProp *prop, CUmemGenericAllocationHandle handle)
L38: Retrieve the contents of the property structure defining properties for this handle.
L39:
L40: CUresult 【9†cuMemImportFromShareableHandle】(CUmemGenericAllocationHandle *handle, void *osHandle, CUmemAllocationHandleType shHandleType)
L41:
L42:
L43: Imports an allocation from a requested shareable handle type.
L44:
L45: CUresult 【10†cuMemMap】(CUdeviceptr ptr, size_t size, size_t offset, CUmemGenericAllocationHandle handle, unsigned long long flags)
L46:
L47:
L48: Maps an allocation handle to a reserved virtual address range.
L49: CUresult 【11†cuMemMapArrayAsync】(CUarrayMapInfo *mapInfoList, unsigned int count, CUstream hStream)
L50:
L51:
L52: Maps or unmaps subregions of sparse CUDA arrays and sparse CUDA mipmapped arrays.
L53:
L54: CUresult 【12†cuMemRelease】(CUmemGenericAllocationHandle handle)
L55:
L56:
L57: Release a memory handle representing a memory allocation which was previously allocated through cuMemCreate.
L58:
L59: CUresult 【13†cuMemRetainAllocationHandle】(CUmemGenericAllocationHandle *handle, void *addr)
L60: Given an address `addr` , returns the allocation handle of the backing memory allocation.
L61:
L62: CUresult 【14†cuMemSetAccess】(CUdeviceptr ptr, size_t size, const CUmemAccessDesc *desc, size_t count)
L63:
L64:
L65: Set the access flags for each location specified in `desc` for the given virtual address range.
L66:
L67: CUresult 【15†cuMemUnmap】(CUdeviceptr ptr, size_t size)
L68:
L69:
L70: Unmap the backing memory of a given address range.
L71: ## 6.47.1. Functions【16†】
L72:
L73: 【17†CUresult】 cuMemAddressFree(【18†CUdeviceptr】 ptr, size_t size)【19†】
L74:
L75:
L76: Free an address range reservation.
L77:
L78: Frees a virtual address range reserved by cuMemAddressReserve. The size must match what was given to memAddressReserve and the ptr given must match what was returned from memAddressReserve.
L79:
L80: See also
L81:
L82: 【3†cuMemAddressReserve】
L83:
L84: Parameters
L85:
L86:
L87: * ptr – [in] - Starting address of the virtual address range to free
L88: * size – [in] - Size of the virtual address region to free
L89:
L90: Returns
L91:
L92:
L93: 【20†CUDA_SUC…[truncated]</excerpt>
</source>
<source kind="search">
<title>6.47. Virtual Memory Management — CUDA Driver API Reference Manual 13.4 documentation</title>
<location>https://docs.nvidia.com/cuda/archive/13.4.2/cuda-driver-api/cuda_driver_api/group__CUDA__VA.html</location>
<excerpt>citeturn1search6 [wordlim: 200] Crawled: last week; CUDA Driver API Reference Manual ... CUresult cuMemExportToShareableHandle(void *shareableHandle, CUmemGenericAllocationHandle handle, CUmemAllocationHandleType handleType, unsigned long long flags)
</excerpt>
</source>
<source kind="search">
<title>4.17. Virtual Memory Management — CUDA Programming Guide</title>
<location>https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/virtual-memory-management.html</location>
<excerpt>citeturn0search0 [wordlim: 200] Crawled: last month; Exporting Memory Handle The CUDA virtual memory management API expose a new mechanism for interprocess communication using handles to exchange necessary information about the allocation and physical address space. ... cuMemExportToShareableHandle(&fd, handle, handleType, 0); ... CUmemAllocationHandleType handleType = CU_MEM_HANDLE_TYPE_FABRIC;
Exporting Memory Handle The CUDA virtual memory management API expose a new mechanism for interprocess communication using handles to exchange necessary information about the allocation and physical address space. One can export handles for OS-specific IPC or fabric-specific IPC. OS-specific IPC handles can only be used on a single-node setup. Fabric-specific handles can be used on a single or multi node setups.
[Input] OS-Specific Handle (Linux)
CUmemAllocationHandleType handleType = CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR;
CUmemGenericAllocationHandle handle = allocatePhysicalMemory(0, 1<<21);
int fd;
cuMemExportToShareableHandle(&fd, handle, handleType, 0);
[Input] Fabric Handle
CUmemAllocationHandleType handleType = CU_MEM_HANDLE_TYPE_FABRIC;
CUmemGenericAllocationHandle handle = allocatePhysicalMemory(0, 1<<21);
CUmemFabricHandle fh;
cuMemExportToShareableHandle(&fh, handle, handleType, 0);
Note
OS-specific handles require all processes to be part of the same OS.
Note
Fabric-specific handles require IMEX channels to be enabled by sysadmin.
The memMapIpcDrv sample can be used as an example for using IPC with VMM allocations.
### 4.17.3.2. Share and Import#
Sharing Memory Handle Once the handle is exported, it must be shared with the receiving process or processes using an inter-process communication protocol. The developer is free to use any method for sharing the handle. The specific IPC method used depends on the application’s design and environment. Common methods include OS-specific inter-process sockets and distributed message passing. Using OS-specific IPC offers high-performance transfer, but is limited to processes on the same machine and not portable. Fabric-specific IPC is simpler and more portable. However, fabric-specific IPC requires system-level support. The chosen method must securely and reliably transfer the handle data to the target process so it can be used to import the memory and establish a valid mapping. = NULL) &&
(cmptr->cmsg_len == CMSG_LEN(sizeof(int)))) {
if ((cmptr->cmsg_level != SOL_SOCKET) || (cmptr->cmsg_type != SCM_RIGHTS)) {
free(control_un.control);
return -1;
}
memmove(&receivedfd, CMSG_DATA(cmptr), sizeof(receivedfd));
*fd = receivedfd;
} else {
free(control_un.control);
return -1;
}
free(control_un.control);
return 0;
}
[Input] Receive: OS-Specific IPC (WIN)
int ipcRecvShareableHandle(HANDLE &handle, HANDLE *shareableHandle) {
DWORD cbRead;
ReadFile(handle, shareableHandle, (DWORD)sizeof(*shareableHandles), &cbRead, NULL);
return 0;
}
[Input] Receive: Fabric IPC
MPI_Recv(&fh, sizeof(CUmemFabricHandle), MPI_BYTE, 1, 0, MPI_COMM_WORLD);
Importing Memory Handle Again, the user can import handles for OS-specific IPC or fabric-specific IPC. OS-specific IPC handles can only be used on a single-node. Fabric-specific handles can be used for single- or multi-node.
[Input] OS-Specific Handle (Linux)
CUmemAllocationHandleType handleType = CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR;
cuMemImportFromShareableHandle(handle, (void*) &fd, handleType);
[Input] Fabric Handle
CUmemAllocationHandleType handleType = CU_MEM_HANDLE_TYPE_FABRIC;
cuMemImportFromShareableHandle(handle, (void*) &fh, handleType);
</excerpt>
</source>
<source kind="search">
<title>4.19. Extended GPU Memory — CUDA Programming Guide</title>
<location>https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/extended-gpu-memory.html</location>
<excerpt>citeturn0search1 [wordlim: 200] Crawled: last month; In addition `CU_MEM_HANDLE_TYPE_FABRIC` should be defined as the requested handle type. ... After creating allocation handle using `cuMemCreate` the user can export that handle to the other node, Node B, calling `cuMemExportToShareableHandle`:
### 4.19.2.3. Multi-Node, Multi-GPU#
Beyond memory allocation, remote peer access does not have EGM-specific modification and it follows CUDA inter process (IPC) protocol. See CUDA Programming Guide for more details in IPC.
The user should allocate memory using `cuMemCreate` and again the user has to explicitly provide `CU_MEM_LOCATION_TYPE_HOST_NUMA` as the location type and numaID as the location identifier. In addition `CU_MEM_HANDLE_TYPE_FABRIC` should be defined as the requested handle type. The following code snippet shows allocating physical memory on Node A:
CUmemAllocationProp prop{};
prop.type = CU_MEM_ALLOCATION_TYPE_PINNED;
prop.requestedHandleTypes = CU_MEM_HANDLE_TYPE_FABRIC;
prop.location.type = CU_MEM_LOCATION_TYPE_HOST_NUMA;
prop.location.id = numaId;
size_t granularity = 0;
cuMemGetAllocationGranularity(&granularity, &prop,
MEM_ALLOC_GRANULARITY_MINIMUM);
size_t padded_size = ROUND_UP(size, granularity);
size_t page_size = ...;
assert(padded_size % page_size == 0);
CUmemGenericAllocationHandle allocHandle;
cuMemCreate(&allocHandle, padded_size, &prop, 0);
After creating allocation handle using `cuMemCreate` the user can export that handle to the other node, Node B, calling `cuMemExportToShareableHandle`:
cuMemExportToShareableHandle(&fabricHandle, allocHandle,
CU_MEM_HANDLE_TYPE_FABRIC, 0);
// At this point, fabricHandle should be sent to Node B via TCP/IP.
On Node B, the handle can be imported using `cuMemImportFromShareableHandle` and treated as any other fabric handle
// At this point, fabricHandle should be received from Node A via TCP/IP.
CUmemGenericAllocationHandle allocHandle;
cuMemImportFromShareableHandle(&allocHandle, &fabricHandle,
CU_MEM_HANDLE_TYPE_FABRIC);
When handle is imported at Node B, then the user can reserve an address space and map it locally in a regular fashion:
</excerpt>
</source>
<source kind="search">
<title>CUDA Driver API</title>
<location>https://docs.nvidia.cn/cuda/pdf/CUDA_Driver_API.pdf</location>
<excerpt>citeturn0search12 [wordlim: 200] Published: 1.1 years ago; If shHandleType is CU_MEM_HANDLE_TYPE_FABRIC and the importer process has not been granted access to the same IMEX channel as the exporter process, this API will error as CUDA_ERROR_NOT_PERMITTED. ... cuMemExportToShareableHandle, cuMemMap, cuMemRelease
Modules
CUresult cuMemImportFromShareableHandle (CUmemGenericAllocationHandle *handle, void *osHandle, CUmemAllocationHandleType shHandleType)
Imports an allocation from a requested shareable handle type.
Parameters
handle
- CUDA Memory handle for the memory allocation.
osHandle
- Shareable Handle representing the memory allocation that is to be imported.
shHandleType
- handle type of the exported handle CUmemAllocationHandleType.
Returns
CUDA_SUCCESS, CUDA_ERROR_INVALID_VALUE, CUDA_ERROR_NOT_INITIALIZED, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_PERMITTED, CUDA_ERROR_NOT_SUPPORTED
Description
If the current process cannot support the memory described by this shareable handle, this API will error as CUDA_ERROR_NOT_SUPPORTED.
If shHandleType is CU_MEM_HANDLE_TYPE_FABRIC and the importer process has not been granted access to the same IMEX channel as the exporter process, this API will error as CUDA_ERROR_NOT_PERMITTED.
Note: Importing shareable handles exported from some graphics APIs(Vulkan, OpenGL, etc) created on devices under an SLI group may not be supported, and thus this API will return CUDA_ERROR_NOT_SUPPORTED. There is no guarantee that the contents of handle will be the same CUDA memory handle for the same given OS shareable handle, or the same underlying allocation.
See also:
cuMemExportToShareableHandle, cuMemMap, cuMemRelease
CUDA Driver API
TRM-06703-001_v13.0 | 276</excerpt>
</source>
<omitted_sources count="33">Additional source records were omitted from this bounded result.</omitted_sources>
</source_evidence>
Citations:
- 1: https://docs.nvidia.com/cuda/archive/13.4.2/cuda-driver-api/cuda_driver_api/group__CUDA__VA.html
- 2: https://docs.nvidia.com/cuda/archive/13.4.2/cuda-driver-api/cuda_driver_api/group__CUDA__VA.html
suggestion: Do not require identical handle bytes across exports.
The assertion at lines 133-135 requires two exports of one allocation to produce equal ipc_descriptor values. The CUDA Driver API does not document that exported fabric handle bytes are stable across exports. The test therefore pins an unspecified driver detail. The cuda_core/tests guidance asks tests to pin behavior the contract guarantees, not internal representation. A driver can return a different valid handle for a later export. This test would then fail even though both descriptors import the same memory.
Replace the byte-equality check with behavior-level checks:
Import each successive export and confirm the imported buffer exposes the same data.
Keep descriptor value-equality coverage on a copied or pickled descriptor, which tests the descriptor's own equality contract.
…r helper
Cython reported the stray line as unreachable code, and the wheel builds
treat Cython warnings as errors. The return goes back to the note helper
it belonged to.
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
The fixture probe skipped every fabric test on any CUDAError from the
probe allocation, which would hide a real failure (an out-of-memory
condition, for one) as 'no fabric here'. It now skips on the statuses
that mean the device has no usable fabric (NOT_PERMITTED on the CI
runners, INVALID_VALUE on an H200 without an IMEX channel, NOT_SUPPORTED)
and lets anything else propagate, per the test guide's rule on
capability probes.
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Reviewing files that changed from the base of the PR and between bd69003 and 7dc241c.
Files selected for processing: 2
cuda_core/tests/conftest.py
cuda_core/tests/memory_ipc/test_vmm_ipc.py
Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 11 remain after this review.
important: Classify only physical-allocation refusals as unavailable.
The probe's try block calls mr.allocate(1), and VirtualMemoryResource.allocate() also reserves address space, maps the allocation, and sets up the deallocation stream. A CUDA_ERROR_INVALID_VALUE from any of those stages matches _FABRIC_REFUSED. The status is then cached for that device, and the fabric fixtures skip for the rest of the session. A real failure on a fabric-capable device can therefore appear as a skip.
Limit the skip to the physical-allocation stage, and let other CUDAErrors propagate. VirtualMemoryResource adds a fabric-specific note only when a physical allocation fails. Check that note through exc.__notes__. str(exc) does not include notes. Test for the note's presence, not its wording, because the text is not a stable contract.
This branch has not been deployed
No deployments
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
cuda.coreEverything related to the cuda.core modulefeatureNew feature or requestP0High priority - Must do!
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Andy-Jost
added
P0
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
Summary
Buffers of a
VirtualMemoryResourceconfigured withhandle_type="fabric"can now be shared between processes on one node. This adds the fabric transport that #2980 tracks, beside the POSIX fd transport from #3028.Changes
VirtualMemoryFabricBufferDescriptor(listed inapi_private.rst): one opaque 64-byteCUmemFabricHandleper physical allocation, the allocation sizes, the buffer size, and the exporting device's UUID. It is plain data:pickleandmultiprocessingcarry it without the file-descriptor machinery, it compares and hashes by value, andfrom_handles()rebuilds one received through a transport of your own, with validation. The export zeroes the handle struct first, so two exports of one allocation carry the same bytes and no stack bytes travel in a descriptor._VirtualMemoryDescriptor, soBuffer.from_ipc_descriptorand the resource's import run one loop for both transports. The fd descriptor keeps its ownership and identity semantics and itsfrom_fds().VirtualMemoryResourceconstruction checksDevice.properties.handle_type_fabric_supportedforhandle_type="fabric"and raisesRuntimeErrorthat names the NVLink fabric and IMEX channel requirements, before an allocation can fail opaquely. A device can report support inside a container that has the IMEX channel device node while the GPU has no fabric; a refused allocation then carries a note that says so.is_ipc_enabledisTruefor"fabric"on Linux with device-located memory.VMM_DESIGN.mdsay so wherever descriptor lifetime is described._ipc.pyiand_virtual_memory_resource.pyiare regenerated by the stub hook.tests/memory_ipc/test_vmm_fabric.pycovers export shape, descriptor validation and pickling, Queue, Pipe, Process-argument, and Pool round trips for descriptors and buffers, a second-GPU import with peer access, handle-type mismatch, garbage and stale handles, re-export of imports, a queued buffer that does not pin, and the unsupported-device error. The fabric fixtures skip on devices without fabric handle support, and on devices that report it but refuse the allocation (one probe allocation per device).Validation
Fabric handles need an NVLink fabric and an accessible IMEX channel. The project's CI has no such runner, so the fabric cases skip there; the release note says that this release's fabric support is not covered by CI. On the CI runners the attribute is true (the container has the channel device node) while every fabric allocation is refused with a driver error, which is why the fixtures probe instead of trusting the attribute.
Validation on NVIDIA internal hardware: a two-GPU B300 system ran the
memory_ipcsuite andtest_memory.pytwice: on the first commit (478 passed, 21 skipped; the one failure found the uninitialized padding in the exported handle struct) and on the fix and on the PR head withmainmerged (479 passed, 21 skipped each), and a Hopper system without fabric support ran the fabric and fd suites (30 passed, 17 skipped, with the unsupported-device error test among the passes). Multi-node sharing is out of scope and untested.Related Work
Part of #2980. Follows #3028 and #3229; this PR includes the #3229 commit until that PR merges.
🤖 Generated with Claude Code