Repository navigation
Support explicit CUDA managed-memory allocation #1523
Description
Description
Add a first-class way for users to explicitly allocate CUDA managed-memory arrays in Warp without changing the default CUDA allocation behavior.
The feature should:
- Provide an opt-in allocator path for CUDA managed memory through Warp's existing allocator configuration mechanisms.
- Preserve the CUDA device/context associated with each allocation so arrays still have normal Warp CUDA device metadata.
- Record allocation provenance so Warp can distinguish managed-memory arrays from default CUDA, CUDA memory-pool, CPU, custom, and externally wrapped allocations.
- Update access queries and checked launch validation so managed arrays are validated through CUDA managed-memory capabilities instead of peer-access or CUDA memory-pool access state.
- Reject unsupported CUDA IPC export for managed-memory arrays with a clear error.
- Allow managed arrays allocated before CUDA graph capture to be used by captured kernels.
- Reject managed allocation during CUDA graph capture with a clear error. The initial implementation uses
cudaMallocManaged()outside capture on all supported CUDA toolkit builds, including CUDA 13.x builds.
Deferred follow-up
Graph-capturable managed allocation through CUDA 13 managed memory pools is intentionally out of scope for this issue and branch. Track it separately if/when needed, since it requires additional design around managed memory pool creation, graph-capture allocation lifetime, free ordering, and CUDA 12.x compatibility.
Context
Warp currently supports default CUDA allocation and CUDA memory-pool allocation, but it does not provide a first-class opt-in path for CUDA managed-memory arrays. Users working on systems with HMM, ATS, concurrent managed access, or limited managed-memory support need a clear way to request Unified Memory while preserving Warp's normal device metadata and diagnostics.
This feature would let users choose managed memory explicitly while keeping standard CUDA arrays unchanged. It would also let Warp make allocation-aware access decisions: CUDA devices that support managed memory can accept managed arrays without relying on peer-access or memory-pool-access state, CPU access can be gated on the owning CUDA device's managed host-access capabilities, and unknown/custom allocations can remain conservative.
CUDA toolkit compatibility is important. Warp's PyPI build currently uses CUDA 12.9, where CUDA managed memory pools are not available through the runtime enum needed for pool creation. This issue therefore scopes initial support to cudaMallocManaged() outside CUDA graph capture on all supported builds, with clear capture-time rejection. CUDA 13+ managed memory pools can be evaluated separately for graph-capturable managed allocation.
Activity
Add ManagedAllocator support [NVIDIAGH-1523]
shi-eric commentedon Jun 26, 2026
This is addressed by e6a8ed4.
The change adds an explicit CUDA managed-memory allocator path while preserving default CUDA allocation behavior. Managed arrays keep normal Warp CUDA device metadata, report
CUDA_MANAGEDthrougharray.memory_kind, and are handled bywp.can_access()and checked launch validation using CUDA managed-memory capabilities. Unsupported CUDA IPC export for managed arrays now raises a clear error. Preallocated managed arrays can be used in CUDA graph captures; allocation during capture is rejected through thecudaMallocManaged()path with CUDA's native diagnostic.Public API behavior:
Python scope:
wp.CudaManagedAllocator()is a new allocator that works with the existing allocator APIs.wp.MemoryKindandarray.memory_kindexpose host, pinned host, CUDA device, CUDA mempool, CUDA managed, and unknown memory kinds.wp.can_access()and checked launch validation use memory kind instead of allocator identity where Warp can classify the allocation.array.ipc_handle()rejects managed-memory arrays, array views, and externally wrapped CUDA pointers with explicit errors.Kernel scope:
wp.can_access()or checked validation says the target device can access the allocation.Test coverage:
warp/tests/test_allocator.pycoverMemoryKindvalues, CPU array memory kinds, managed allocator protocol behavior, managed allocation on the selected CUDA device, deallocation context handling, and managed allocator failure reporting.warp/tests/cuda/test_unified_memory.pycover CUDA memory-kind queries, wrapped managed pointer classification, managedwp.can_access()behavior, managed IPC rejection, checked CPU launches with managed arrays, preallocated managed arrays in CUDA graph capture, and memory-kind diagnostics for unknown or unowned allocations.This issue can be closed.
Add ManagedAllocator support [NVIDIAGH-1523]
Add ManagedAllocator support [NVIDIAGH-1523]
Add ManagedAllocator support [NVIDIAGH-1523]