NanoVDB: own PointsToGrid's device arrays with cuda::Buffer (CUDA) - #2270
Conversation
Name the free operation destroy, as cuda::buffer does, and add the stream-taking overload it also provides. clear stays but only as a transitional delegate, marked as such: it exists because GridHandle::reset still calls it, and goes away with the legacy dual buffers that cuda::Buffer replaces. Rename setStream to set_stream. Member names cannot be aliased, so matching the standard spelling is the whole reason the resource concept kept allocate_async; the same argument applies here. Note in passing that ours deliberately does not synchronize, matching cuda::buffer's set_stream_unsynchronized rather than its set_stream, whose own documentation and implementation disagree (NVIDIA/cccl#10649). Add swap. The generic std::swap already does the right thing through the move operations, but both std::vector and rmm::device_buffer provide one. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
Eleven of the builder's buffers are device-only -- nothing reads them on the host -- yet they were cuda::DeviceBuffer, whose host pointer and per-device array they never use. Move them to the single-space cuda::Buffer and give the builder a resource parameter, so its scratch is injectable like PointsToGrid's already is, and freed by scope rather than by hand. mProcessedRoot and mData are read on the host and stay dual. This is not a speedup: DeviceBuffer::init allocates host or device memory but never both, and these buffers always passed a real device id, so no cudaMallocHost was on this path to begin with. Measured dilate on 1k/20k/ 200k points, and the difference is within run-to-run noise. Hoist the Data struct out of the class. It does not depend on the resource, and leaving it nested would give every ResourceT its own incompatible type for the device functors to name. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
A stream-ordered resource must also model the synchronous concept, which means writing four methods where two would do. The synchronous pair is not a bare delegate -- memory from allocate must be usable on any stream when it returns, so the null-stream allocation has to be synchronized first -- and omitting that yields memory which satisfies the concept but is not actually synchronous. Put it in one place rather than leaving each author to rediscover it. The two resources in TestMemoryResource are the first users, and were already wrong in exactly that way: they provide only the async pair, so they never modelled is_async_resource. TempPool duck-typed and never checked, so nothing caught it. MeshToGrid was the last builder allocating from a hard-wired DeviceResource, through TempDevicePool. Give it a ResourceT parameter and thread it into both its TopologyBuilder and its pool. As with Data in the builder, BoxTrianglePair is hoisted out of the class: it does not depend on the resource, and leaving it nested would give every ResourceT its own incompatible type for the device functors to name. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
…uffer Buffer holds its resource by value, matching cuda::buffer -- whose model this completes: in CCCL the ownership semantics are selected by what is placed in the by-value slot, an owning any_resource or a borrowing resource_ref. We adopted the slot without the borrowing type, so a container like TempPool, whose contract is a non-owning pointer to a possibly stateful resource, had no way to hold a Buffer without copying that resource and stranding its state. ResourceRef is the missing piece: a non-owning reference that is itself a resource, so copying the ref shares the underlying instance. Its async methods exist only when R models AsyncResource, so a ref over a synchronous resource does not misreport its tier, and two refs compare equal exactly when they reference the same resource. TempPool now keeps its bytes in a Buffer<std::byte, ResourceRef<R>>: same resource contract, same stream retention, same discard-on-growth reallocation, but the block is freed by ownership rather than by hand. The TempPool unit tests, which assert traffic against the caller's own resource instance, pass unchanged. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
The bisection search over voxel size was written as a backward goto, which made the lifetimes of the buffers it retries over non-lexical. Rewrite it as while(true) with continue on retry and break on convergence; the six hand-written frees before the jump are unchanged. The change is easiest to review with whitespace ignored, since the loop body re-indents: git diff -w shows 27 changed lines. d_keys and d_node_count carry results past the loop, so their declarations move above it, as does the copy event, which was created inside the retried region on every iteration but destroyed only once at the end -- each retry leaked the previous handle. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
Every device array PointsToGrid allocates is now owned by a Buffer<T, ResourceRef<ResourceT>> borrowing the injected resource -- members where the pipeline frees them in a later member function than the one that allocated them (countNodes allocates; processUpperNodes, processLeafNodes, processPoints and processBBox release), locals where the lifetime is contained. The raw pointers survive only as views: the device-visible fields inside mData, and the working pointers the cub and kernel calls take. Every owner event -- assignment, swap, destroy -- immediately refreshes its view, and released views are nulled so a stale use faults instead of reading a freed block. The index ping-pong becomes a swap of owners across the member/local boundary, replacing the bare pointer swap whose safety depended on nothing reading the device copy of d_indx between the two uploads. The density-search retry keeps its free-before-reallocate order via explicit destroy calls, so peak device memory is unchanged. The hand-matched byte sizes at every free site disappear, and the arrays released one line before scope exit now just leave scope. One behavior change worth naming: the too-many-points-per-leaf throw previously leaked the reduction scratch; ownership now releases it during unwind. Verified against the previous commit with a counting resource over three shapes (bulk segmented-sort branch, serial per-tile branch, and the bisection retry engaged): allocation count, free count, and total bytes are identical. Full CUDA and memory-resource suites unchanged. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
Each release site sits in a function that receives the stream, so pass it to destroy rather than relying on the retained stream matching -- they are the same on every current path, but the explicit form does not depend on that staying true. Also drop the cudaGetDevice calls whose result the Buffer conversion left unused. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
The scratch buffers held their resource by value, so each of the eight carried its own copy -- fine for the stateless default, wrong for a stateful resource, whose accounting would be split across copies while the caller's instance saw nothing. Borrow through ResourceRef instead, the same reconciliation TempPool uses. Assert the stream-ordered requirement directly in TopologyBuilder and MeshToGrid so a synchronous-only resource fails with a diagnostic that names the builder, not just the pool inside it. Note SyncFromAsync's synchronize cost on its allocate. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
The too-many-points-per-leaf throw unwinds past the event's manual destroy, leaking the handle. Own it with a small guard so unwinding releases it, consistent with the buffer ownership in this function. Also assert the stream-ordered resource requirement on the class, so a synchronous-only resource fails naming PointsToGrid rather than the pool inside it. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
The byte scratch is reinterpreted as word-sized types, which is valid for every resource whose DEFAULT_ALIGNMENT is at least word alignment -- all CUDA allocation paths give 256 -- but nothing said so. Assert it, so a custom resource with a weaker guarantee fails at compile time instead of misaligning on the device. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
Legal without one -- the macro expands to a braced block -- but every other use in the file spells it as a statement. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
|
Copilot cycles (fork mirror: harrism#3). Net changes: the copy event is owned by a scope guard so the too-many-points-per-leaf throw cannot leak the handle (4957582 — pre-existing leak), |
cuda::buffer's set_stream is deliberately non-synchronizing -- its documented synchronization is a stale note left behind when the synchronization was removed -- so our non-synchronizing set_stream matches it in both name and contract, and the comment claiming a divergence was wrong. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
Documentation-level only: the [[deprecated]] attribute would warn from GridHandle::reset and NodeManager::reset, template members in our own headers that must keep calling clear() until every buffer type provides destroy() -- an unactionable diagnostic for callers of reset(), and a build break under -Werror. The attribute lands when those callers migrate. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
…trofit Signed-off-by: Mark Harris <mharris@nvidia.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
The trait checks detect allocate/deallocate through unevaluated contexts, which never odr-use them, so nvcc warned that every synchronous pair and stub member was declared but never referenced (AcademySoftwareFoundation#177-D). The attribute route is closed -- nvcc's front end ignores [[maybe_unused]] for this diagnostic -- so reference them the honest way: a test that verifies the synchronous halves of the counting and stream-recording doubles behave like their stream-ordered halves, and one that pins the trait probes' stub behavior. The TU now builds with no warnings at all. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
…buffer Signed-off-by: Mark Harris <mharris@nvidia.com> # Conflicts: # nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh # nanovdb/nanovdb/unittest/TestBuffer.cu
…rid-buffer Signed-off-by: Mark Harris <mharris@nvidia.com> # Conflicts: # pendingchanges/nanovdb.txt
|
@kmuseth #2269 merged — thank you! Latest master is merged in (no rebase; existing commits unchanged) and the diff now shows only this PR's payload: PointsToGrid's ~47 hand-paired device allocations replaced with RAII |
kmuseth
left a comment
There was a problem hiding this comment.
looks good - a clear progression from 2269 (I'm starting to get the bigger picture here)
| ~EventGuard() { cudaCheck(cudaEventDestroy(event)); } | ||
| EventGuard(const EventGuard&) = delete; | ||
| EventGuard& operator=(const EventGuard&) = delete; | ||
| } eventGuard; |
… via upstream memory-resource seams (#732) ## Summary fvdb's grid builders allocate their device scratch from nanoVDB's default `DeviceResource` — a second `cudaMallocAsync` pool that partitions VRAM against PyTorch's. Large workloads (e.g. multi-frame TSDF integration) then hit a clean OOM even when the GPU has free memory in aggregate. This routes that scratch — O(N-points) sort keys, CUB temp storage, topology mask buffers — through PyTorch's CUDA allocator instead, so it shares one pool with fvdb / PyTorch tensors. **22 sites across 13 `.cu` files plus `PadGrid.cuh`.** Note this is not hardcoded to Torch's *native* caching allocator: `c10::cuda::CUDACachingAllocator` is a namespace, and its `raw_alloc_with_stream` / `raw_delete` free functions dispatch through `CUDACachingAllocator::get()` — the runtime-swappable allocator Torch itself allocates tensors from. fvdb's scratch therefore follows whatever allocator the user has installed: the native caching allocator (including `PYTORCH_CUDA_ALLOC_CONF` knobs), the `cudaMallocAsync` backend (`PYTORCH_CUDA_ALLOC_CONF=backend:cudaMallocAsync`), or a custom allocator installed via `torch.cuda.memory.change_current_allocator(CUDAPluggableAllocator(...))`. **Supersedes #655**, which vendored modified nanoVDB headers into the tree. This instead uses the injectable-memory-resource seams we developed upstream (AcademySoftwareFoundation/openvdb#2232; PRs [#2268](AcademySoftwareFoundation/openvdb#2268), [#2269](AcademySoftwareFoundation/openvdb#2269), [#2270](AcademySoftwareFoundation/openvdb#2270), [#2272](AcademySoftwareFoundation/openvdb#2272), [#2273](AcademySoftwareFoundation/openvdb#2273)) — now merged, so the pin is plain upstream master. No fork, no include-path shadowing, no resync procedure. ## What's in this PR 1. **Pin nanovdb to upstream master.** `src/cmake/get_nanovdb.cmake` → `AcademySoftwareFoundation/openvdb @ 7946f17e`, which includes the small-builder `ResourceT` seams ([#2286](AcademySoftwareFoundation/openvdb#2286)), the synchronous resource adapters ([#2272](AcademySoftwareFoundation/openvdb#2272)), and the `MeshToGrid` `CALL_CUBS` `#undef` fix ([#2284](AcademySoftwareFoundation/openvdb#2284)). 2. **`fvdb::TorchResource`.** A ~40-line stateless resource ([`src/fvdb/TorchResource.h`](src/fvdb/TorchResource.h)) modeling nanoVDB's stream-ordered `AsyncResource` concept over `c10::cuda::CUDACachingAllocator::raw_alloc_with_stream` / `raw_delete` — the dispatchers to Torch's currently active CUDA allocator (see Summary). Passed as the `ResourceT` template parameter at all 13 upstream builder call sites — `voxelsToGrid`, `DilateGrid`, `MergeGrids`, `PruneGrid`, `RefineGrid`, `CoarsenGrid` — always via the `fvdb::BuilderResource` alias ([`src/fvdb/BuilderResource.h`](src/fvdb/BuilderResource.h)), never named directly, so the allocator policy lives in a single line (a non-torch build, e.g. the ONNX Runtime EP planned in #579, retargets the alias there instead of touching every op). Being stateless, it binds through each builder's defaulted constructor argument, so no instance is plumbed through. Retains #655's `FVDB_NANOVDB_TRACE_ALLOCS` tracing (`=1` traces ≥ 256 KiB, a value starting with `2` traces everything). 3. **`PadGrid` gains a `ResourceT` seam.** The conv builders used `DilateGrid<..., TorchResource>` for odd kernels and `PadGrid` — on the rival pool — for even ones: same loop, same grid, a different allocator depending on kernel parity. fvdb's own `morphology::PadGrid` drives nanoVDB's `TopologyBuilder` (internal mask buffers, `countNodes` CUB scratch, `TempPool`) but hardcoded `DeviceResource`. It now takes a `ResourceT` parameter mirroring the upstream `DilateGrid` signature and forwards it, with `BuilderResource` passed at all 7 call sites. The default keeps it source-compatible. 4. **CUB scratch in `BuildFineGridFromCoarse`.** `cub::DeviceSegmentedReduce` temp storage used a bare `cudaMallocAsync`; it now routes through `BuilderResource`. Both `cub` calls are also now `C10_CUDA_CHECK`-wrapped — previously unchecked, as was the allocation. 5. **The `SaveNanoVDB` CUDA path.** The save path allocated its largest device buffers from nanoVDB's default pool: the per-batch `(N+1)`-element value staging buffer, the `indexToGrid` output grid handle, and the defensive host-upload buffer. All three now use `TorchDeviceBuffer`, and `indexToGrid`'s internal scratch routes through `TorchResource` via the #2286 seam. Stream-ordering is preserved: the replaced stream-ordered `DeviceBuffer` constructors become `raw_alloc` on the same current stream the copies and kernels are queued on. The host path (`indexToGridHost`) and the `HostBuffer` file-staging buffers are unchanged. ## Not routed (no upstream seam yet; all off the hot paths) - `DistributedPointsToGrid` multi-GPU scratch (deferred upstream behind AcademySoftwareFoundation/openvdb#2248) — the most valuable remaining seam - `VoxelBlockManager` / `buildVoxelBlockManager` scratch in `ReinitializeSdf.cu` - the builders' small dual-space `mProcessedRoot` / `mData` buffers (upstream roadmap Step 3) `MeshToGrid` is the one merged seam fvdb does not use: `BuildGridFromMesh.cu` does its own parametric surface sampling and goes through `_createNanoGridFromIJK`, so there is nothing to route. ## Test plan - [x] `./build.sh install` succeeds on a clean tree against the new pin (full CUDA build, `-Werror`). - [x] Injection verified live via `FVDB_NANOVDB_TRACE_ALLOCS`: a 500k-point `Grid.from_points` + `dilated_grid(2)` prints 42 `TorchResource` traces with correct results (484,631 → 8,461,871 voxels); `from_nearest_voxels_to_points` at 2M points shows `PadGrid` scratch routed. - [x] 584 tests passed — conv semantics + integration (203), conv/conv-transpose default + prune + empty grids (103), basic ops (276, 1 skipped), sliced batch (2, covering the `BuildFineGridFromCoarse` CUB path). - [x] `test_io.py` — 622 passed against the `7946f17e` pin; a traced `save_nanovdb` (`FVDB_NANOVDB_TRACE_ALLOCS=2`) shows the `indexToGrid` scratch flowing through `TorchResource`. ## Followups - Rebase the TSDF / ESDF / Occupancy stack (#656) onto this branch in place of #655, threading `TorchResource` through the new ops it adds. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Signed-off-by: Jonathan Swartz <jonathan@jswartz.info> Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
Follow-up to #2231/#2251, part of #2232. This is B3, the last of step 2's scratch retrofit:
PointsToGridwas the remaining builder with hand-pairedallocate_async/deallocate_asynccalls — 47 of them.Commit 1 — express the density search as a loop
The bisection search over voxel size was a backward
goto, which made the lifetimes of the buffers it retries over non-lexical. It becomeswhile (true)withcontinueon retry andbreakon convergence. Behavior-preserving; the loop body re-indents, sogit diff -wshows the real change: 27 lines. Two locals (d_keys,d_node_count) carry results past the loop and hoist above it, as does the copy event — which was created inside the retried region on every iteration but destroyed once at the end, so each retry leaked an event handle; it is now created once.Commit 2 — ownership
Design rule, refined from the note on the roadmap: owners own; raw pointers are views, refreshed immediately after every owner event (assignment, swap, destroy). The device-visible fields inside
mData— which ismemcpy'd to the device twice — are therefore never stale, replacing the previous unwritten invariant that nothing reads the device copy ofd_indxin the ~185 lines between the two uploads. Released views are nulled, so a stale use faults instead of silently reading a freed block.countNodesand freed in a different member function (processUpperNodes,processLeafNodes,processPoints,processBBox), so their owners are members; contained lifetimes are locals. All owners areBuffer<T, ResourceRef<ResourceT>>, borrowing the injected resource per the ownership rule of record (NanoVDB: injectable CUDA memory resources — design & roadmap #2232).mIndxBuf.swap(indxScratch)) instead of a bare pointer swap between a local and an uploaded struct field.destroy()calls, so peak device memory at retry is unchanged (move-assignment alone would briefly hold old + new).throwpreviously leaked its reduction scratch; ownership releases it during unwind.Verification
Alloc/free parity gate (the roadmap's acceptance test for this conversion): a counting resource over three shapes — bulk segmented-sort branch (200k points, many tiles), serial per-tile branch (2k points, one tile), and the bisection retry engaged (
PointsToGrid<Point>(4, 0, 8), 50k points):Full CUDA suite 52/53 (the failure is the pre-existing
UnifiedBuffer_IOmissing-data-file issue, identical on baseline); memory-resource suite 9/9; all instantiations (Point+custom,float+custom,Point+default) compile clean.Per NanoVDB: please run the CUDA unit tests on a GPU in CI #2264, CI builds but does not run these — the numbers above are from a local RTX 6000 Ada, CUDA 12.6.
DistributedPointsToGrid(49 sites) stays deferred behind #2248. The remaining small builders (IndexToGrid4,AddBlindData2,GridStats2,SignedFloodFill2,VoxelBlockManager1) and the sync-resource dispatch (step 2's vGPU acceptance criterion) are follow-ups.