AcademySoftwareFoundation / AcademySoftwareFoundation/openvdb
NanoVDB: injectable CUDA memory resources — design & roadmap
Nobody has claimed this yet.
- Dominant language
- C++
- Stars
- 3.4k
- Forks
- 777
- Avg merge
- 3d 9h
- Merged PRs (30d)
- 34
Description
Summary
NanoVDB's CUDA path hard-wires its allocations — grid storage via GridHandle<BufferT>, scratch via cuda::TempPool/cuda::DeviceResource, and raw cudaMallocAsync/cudaMallocHost in builders. This tracks a layered design to let downstream inject its own allocator wherever NanoVDB allocates, aligned with CCCL's cuda::mr / cuda::buffer model — without a hard CCCL dependency. New APIs land alongside the existing ones; the legacy dual buffer + handle surface are deprecated and removed a release or two later, so there is no hard break at introduction.
Why
- Inject your own allocator without forking headers — downstream currently forks NanoVDB headers to route through their pool (e.g. PyTorch's
c10::cuda::CUDACachingAllocator; cf. openvdb/fvdb-core#655, which vendored three headers). - Fix the pinned-host stall — the synchronous
cudaMallocHostinsidecuda::DeviceBufferserializes streams where a dual buffer is actually constructed host-side. Note this is not everyDeviceBufferuse:initallocates host or device but never both, so buffers created with a real device id never touchcudaMallocHost(see the correction under B1). - Share one pool across tensors, scratch, and grid buffers (caching; fewer fragmentation OOMs).
- Familiar + future-proof — method/type shapes match CCCL, so a
resource_ref/cuda::bufferadapter is a thin shim later, while NanoVDB keeps its low CUDA floor.
Design — three pieces, each a minimal NanoVDB analog of a CCCL standard
- Resource (allocator): two tiers matching CCCL's refinement. A synchronous resource exposes
allocate(bytes, alignment)/deallocate(...)— no stream (cuda::PinnedResourceis one). A stream-ordered resource addsallocate_async/deallocate_async(bytes, alignment, stream)and still provides the sync pair (typically thin delegates through the null stream), sois_async_resource<R>impliesis_resource<R>— exactlycuda::mr's shape.cuda::DeviceResourceprovides all four; its legacy staticallocateAsync/deallocateAsyncare[[deprecated]]. Plusdefault_resource<R>()and the two detection traits. A custom allocator is a ~25-line struct. cuda::Buffer<T,R>(storage) +cuda::BufferView<T>(view): one single-space, resource-aware, stream-ordered container modelled on CCCL'scuda::buffer<T>(the RMMdevice_buffer/device_uvectorlineage).Buffer<T>typed,Buffer<std::byte>raw. Constructor order and shape followcuda::bufferexactly —(stream, resource, count, noInit), the tag mandatory: there is no implicitly initializing count constructor (implicit fill of fresh memory is a hidden cost the allocate-then-overwrite pattern wastes).size()counts elements,size_bytes()bytes; a synchronousRyields a buffer with no stream API at all. Alongside it,BufferView<T>: a typed non-owning view with span semantics (element constness inT, trivially copyable, no resource, no stream) that satisfies the same static interface — so aGridHandlecan wrap externally owned memory (an ONNX Runtime or Torch tensor) with zero copies. Owning and viewing do not share a type: the legacymManagedowns-or-wraps flag is deleted, not re-spelled.- Single-space
GridHandle(surface): one buffer per handle;grid()returns its pointer; host↔device is an explicit stream-carryingcopyTo.
Rules that matter
- Stream ownership: a buffer retains its allocation stream and frees on it (not the null stream); the stream must outlive the buffer or be reset via a non-synchronizing
setStream(the RMMcuda_stream_viewcontract).resizeorders everything — including the free of the old block, whose last use is the prefix copy — on the passed stream. - Resource ownership (revised — decision of record):
cuda::Bufferholds the resource by value, matchingcuda::buffer— but ownership semantics are selected by what is placed in that slot, completing the CCCL model rather than adopting half of it (cuda::bufferowns anany_resource; CCCL's borrowing and sharing tiers areresource_refandshared_resource, and its constructor'sstatic_assertpoints users at them). The tiers here: a concrete resource in the slot is owned as a copy (so it must be a cheap-to-copy handle, stateless or pointer-to-state);ResourceRef<R>(#2269) borrows — a non-owning, dependency-free static analog ofcuda::mr::resource_refthat is itself a resource, withenable_if-gated async members so a ref over a synchronous resource does not misreport its tier, and pointer-identity equality; a refcountedSharedResource<R>analog is deferred until the Python-bindings need materializes. Non-owning containers (TempPool, and the builders'ResourceT*members over time) compose withBufferviaResourceRef. History: an earlier decision ("reference-at-API, pointer-as-member", pre-cuda::bufferreshape) was superseded by the by-value rule without reconciling the artifacts built on it —TempPool's pointer contract and the stateful test resources — which is what the firstTempPoolconversion attempt tripped over. Type-erased CCCLresource_refinterop remains step 4. - Synchronous arenas (e.g. ONNX Runtime's) are wrapped as synchronous resources — never as a silent
allocate_asyncfacade (a facade misrepresents its semantics and hands multi-stream callers unexpected serialization). Contract:deallocateimplies the memory is quiescent. The explicit lift isAsyncFromSync<R>(#2272), the analog of CCCL'ssynchronous_resource_adapter:allocate_asyncforwards (synchronously allocated memory is valid on every stream — stronger than stream-ordering requires) anddeallocate_asyncsynchronizes the stream first, making the quiescence contract hold; the serialization cost is documented at the type and chosen by the caller. - Naming: types are CamelCase per OpenVDB style (
cuda::Buffer,cuda::BufferView) — type names are aliasable, so the opt-inusing Buffer = ::cuda::buffer<T>at CUDA ≥ 13.2 is preserved. Member names keep the standard spelling (data,size,size_bytes,allocate_async) because member matching is structural and cannot be aliased. - CUDA graph capture: the async resource tier is capture-safe — stream-ordered allocation records as graph allocation/free nodes, and
Buffer's async path performs no hidden synchronization or initialization (verified by a capture → instantiate → relaunch test). The sync tier can never be captured (it synchronizes), visibly: a sync-Rbuffer has no stream API. - No hard CCCL dependency; an optional adapter bridges to
cuda::mr/cuda::buffer/cuda::std::span.
Roadmap
-
Step 1 — resource concept (#2231, merged):
DeviceResourceinstance methods +default_resource+ detection traits; newPinnedResource;TempPoolroutes through a resource instance and frees on its retained stream;PointsToGridinstance-injection seam. Additive. Follow-up in review: #2244 (point encoding for any resource). -
Step 2 —
cuda::Buffer<T,R>+BufferView<T>+ scratch retrofit (complete: B1 #2268, B2 #2269, B3 #2270, B4 #2272; plus #2273 consumer seams and #2286 small-builder scratch): ship the container, the view, the synchronousis_resourcetrait and the sync/async dispatch (PR A — #2251, merged); then convert the raw alloc/free onto the container — the "no raw allocation" cleanup — with a builder-coverage audit. The retrofit was originally scoped as a single PR B; a scoping pass found the three targets differ enough in risk that they land separately:-
B1 —
TopologyBuilder(#2268): 8 of its 10cuda::DeviceBuffermembers are device-only (zero host.data()uses), as are 3 function-local buffers. Convert those 11 toBuffer<std::byte,R>and add a defaultedResourceTparameter; onlyDilateGridandMergeGridsinstantiateTopologyBuilder, so existing callers are unaffected.mProcessedRootandmDataare genuinely dual-space and stay until Step 3. Carries the builder-coverage audit table. Correction: an earlier revision of this entry described B1 as fixing the pinned-host stall. It does not.DeviceBuffer::initallocates host memory or device memory, never both, andTopologyBuilderalways passed a real device id, so nocudaMallocHostwas ever on this path. Measured on dilate at 1k/20k/200k points: no difference beyond run-to-run noise. B1's value is injectability (the last builder without the resource seam) and scope-based ownership, not speed. The pinned-host stall, where it exists, is onGridHandle-side paths that construct host-side buffers, and belongs to Step 3. -
B2 — resource ergonomics + the remaining seams (#2269):
SyncFromAsync<Derived>CRTP mixin (a custom resource is two methods rather than four, with the mandatory synchronize in one audited place — its first users were the two test resources inTestMemoryResource, which turned out never to have modelled the concept); aResourceTseam forMeshToGrid, the last builder on a hard-wiredDeviceResource;ResourceRef<R>per the revised ownership rule above; andTempPool's bytes ontoBuffer<std::byte, ResourceRef<R>>— same pointer contract and stream retention, block freed by ownership (discard-on-growth viadestroy(stream)+ move-assignment; the pool keeps asize_tmirror because cub's two-pass API wants a mutablesize_t&).Buffer::swapand thedestroy/set_streamalignment landed in #2268. -
B3 —
PointsToGrid(#2270, merged): 48 sites, and the only target that is not mechanical.mData.*are device-visible raw pointers uploaded at two separate points;std::swap(d_indx, mData.d_indx)moves ownership between a local and one of those aliased fields after the first upload; agotoretry loop over voxel density makes lifetimes non-lexical; and allocations are freed across three exit paths ~500 lines apart. Wants a written ownership design plus a separategoto-to-whileprep commit before the RAII change. Sequenced after #2244, which edits the same function. -
B4 — builders on a synchronous resource (#2272):
MallocResource(synchronouscudaMalloc/cudaFree, works on pool-less devices) +AsyncFromSync<R>per the rule above. Closes the scratch half of the acceptance criterion as pure library code —PointsToGrid<BuildT, AsyncFromSync<MallocResource>>, no build flags. The grid handle's buffer still needsNANOVDB_USE_SYNC_CUDA_MALLOCuntil Step 3 deliversGridHandle<cuda::Buffer<std::byte,R>>. Tested by drivingPointsToGridend-to-end on an injected stateful synchronous resource with balanced accounting.
Splitting matters more than usual here because #2264 means no CUDA test runs on a GPU in CI, so a mistake in the intricate conversion would land unverified by review. Additive throughout. Acceptance criterion (from #2255): NanoVDB builds and runs on a pool-less vGPU (
cudaDevAttrMemoryPoolsSupported == 0, e.g. AWSg6f) via an injected synchronouscudaMalloc-backed resource; interim coverage via theNANOVDB_USE_SYNC_CUDA_MALLOCmacro (the caller's explicit opt-in). Theutil::cudawrappers deliberately do not silently fall back: on a pool-less device they fail with an actionable diagnostic (#2256) rather than making an async resource misrepresent its semantics; CI detects the capability with acudaDevAttrMemoryPoolsSupportedprobe and sets the macro. -
-
Step 3 — single-space grid storage, via deprecation (introduce + deprecate complete, shipping in OpenVDB 13.1; the remove step is tracked below) (end state:
GridHandle<cuda::Buffer<std::byte,R>>directly — one buffer per handle, no dual surface, cross-space viacopyTo;GridHandleoverBufferViewfor externally owned blobs):- Introduce: ship the single-space handle path; re-implement the legacy
cuda::DeviceBufferinternally as a composition of twocuda::Buffers (transparent — same API, gains the pinned-host fix); migrate NanoVDB's own internal uses + tests/examples off theDeviceBuffername so the later deprecation fires only externally. - Deprecate:
using DeviceBuffer [[deprecated]] = <impl>+[[deprecated]]on the dualGridHandle/NodeManagermethods (deviceUpload/deviceDownload/deviceGrid/deviceData). Both still compile; warnings external only. Soak one or two releases. - Remove: delete the dual buffer + dual surface +
hasDeviceDual; flip entry-point defaults to single-space. The break lands here, after the window, only for code that didn't migrate. - Direction of adaptation (rule of record):
cuda::Bufferis the type the legacy buffers are being replaced by, so it keeps the interface we want to live with — thecuda::buffer-aligned spelling — andGridHandle/NodeManageradapt to it. Not the reverse: legacy names are not accreted ontocuda::Buffermerely becauseHostBuffer/DeviceBuffer/UnifiedBuffercarry them, since those three are being deleted. Concretely,GridHandleshould move todestroy()andempty()rather thancuda::Buffergainingclear()andisEmpty()permanently. Where a shim is genuinely needed to survive the deprecation window it is marked transitional at the declaration, naming what replaces it and when it goes — ascuda::Buffer::clearis in #2268. This applies to every spelling the window might otherwise drag across:isEmpty,deviceData,deviceUpload,deviceDownload,create. - Known gap:
GridHandle::copy()callsmBuffer.isEmpty(), whichcuda::BufferandBufferViewdo not provide (they haveempty()). It is latent today becausecopy()is only instantiated on use, but it blocks the end state. Per the rule above the fix is inGridHandle, not incuda::Buffer— and it needs more than a rename regardless, sincecopy()also callsOtherBufferT::create(...)and performs a host-sidestd::memcpyover the buffer contents, neither of which is valid for device memory.
- Introduce: ship the single-space handle path; re-implement the legacy
-
Candidate — capacity-bounded, sync-free (graph-capturable) build path: graph capture prohibits host logic on device-computed values, and
PointsToGridreads count reductions back to size allocations and launch dimensions — so building cannot be captured regardless of allocator (cf. NVIDIA/warp GH-1606, which reimplemented grid building privately for exactly this reason). Sketch: caller-supplied capacity bounds; one up-front allocation through the injected resource or into caller-owned memory viaBufferView; launch dimensions from capacity; counts consumed on device; optional deferred readback. Requirements gathered from the Warp team: an overflow-clamped grid must remain safe to traverse from the root (orphaned leaves are acceptable — safety, not full well-formedness); point-mask support is required; a CPU counterpart is a plus (single-source maintenance); build performance must not regress; capacity-growth policy stays with applications. The memory seam is the prerequisite infrastructure; this is its own item. -
Step 4 — optional CCCL adapter (gated on availability): native
cuda::mr/cuda::bufferinterop,cuda::std::spanconversions, and runtime (resource_ref) selection.
Migration note — custom ResourceT contract
Releases v12.1.0–v13.0.0 accepted a custom resource with static allocateAsync(bytes, alignment, stream) / deallocateAsync(...). Since #2231, builders call instance methods, and the concept now matches CCCL's refinement as described above. The static methods on cuda::DeviceResource are deprecated and will be removed after a deprecation window. Migrating a v13-era custom resource is mechanical: drop static, rename to the snake_case instance forms, add the two-line sync delegates (~8 lines total — see cuda/DeviceResource.h for the reference shape).
Downstream payoff (fvdb-core)
fvdb deletes its three forked headers and writes a small TorchAllocatorResource (forwarding to c10::cuda::CUDACachingAllocator), then uses PointsToGrid<…, TorchAllocatorResource> or GridHandle<cuda::Buffer<std::byte, TorchAllocatorResource>>; ONNX Runtime kernels wrap ORT-owned grid blobs zero-copy via GridHandle<BufferView<std::byte>>. No fork, no patch fragility.
Progress — 2026-09-10
Everything through the deprecation window has merged and ships in OpenVDB 13.1:
- Introduce: #2288 (single-space
GridHandle<cuda::Buffer<std::byte,R>>), #2292 (cuda::copyTocross-space transfers + pinned/managed handles), #2293 (reset→destroy,NodeManagersingle-space overload). - Deprecate: #2301 (
DeviceBuffer→ transitionalDualDeviceBuffer+[[deprecated]]alias; every in-tree consumer migrated;copyToadopts validated metadata so it runs no kernel and is host-callable; handle construction from host bytes now validates the whole grid chain) and #2319 (DistributedPointsToGridper-device resources;UnifiedBuffer→DualUnifiedBuffer+[[deprecated]]alias; multi-GPU tests/examples onto managedcuda::Buffer). Both names deprecate together; default-using callers are warning-free until removal.
Remaining — the remove step (one release after 13.1):
- Flip the tools'
BufferTdefaults tocuda::Buffer<std::byte>— the largest API-visible event of the epic (tool return types change). - Delete the dual family:
DualDeviceBuffer,DualUnifiedBuffer,CudaDeviceBuffer,hasDeviceDual, the dualGridHandle/NodeManageraccessors, and the transitionalclear()/isEmpty()shims. - Decide the resource-direct signature question: with
create()gone, the tools' pool/buffer parameters are vestigial;createNodeManager's resource overload is the model. - Headers:
UnifiedBuffer.hholds nothing but the removed pieces — replace with a one-line#errortombstone naming the replacement for one release, then delete.DeviceBuffer.his gated: the python bindings expose aDeviceBufferclass whose API is the dual model (deviceUpload/deviceDownload), retargeted onto the implementation by #2301; deleting the header requires a python-API redesign around single-space handles andcopyTo— a deliberate decision, not a cleanup.
Queued alongside: #2312 (type the seven Mask sidecar buffers as Buffer<Mask<N>>, unlocked by #2310), step 4 (optional CCCL adapter). The multi-GPU test failures in #2245 are confirmed environmental and fixed by iommu=pt on the affected host — the full mgpu suite now passes on two GPUs.
Contributor guide
First steps
- Read the whole issue, then the project's contributing guide.
- Comment on the issue to say you are picking it up — it saves two people doing the same work.
- Fork the repository and make your change on a branch.
- Open a pull request that references the issue number.
Research direction
Review the remaining roadmap in issue #2232 and the linked follow-up issues, especially the GridHandle storage work after the completed resource and Buffer steps. Start by checking which roadmap items are still open and their referenced acceptance criteria; done means the planned injectable CUDA allocation design is implemented without breaking existing callers and the stated tests or runtime checks pass.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- hpc
- Issue type
- Feature
- Difficulty
- 5/5
- Estimated time
- Over a week
- Activity status
- Quiet
- Clarity
- Mostly clear
- Newbie friendliness
- 35/100