Skip to content

NanoVDB: SyncFromAsync, ResourceRef, and resource seams for MeshToGrid and TempPool (CUDA) - #2269

Merged
kmuseth merged 18 commits into
AcademySoftwareFoundation:masterfrom
harrism:nanovdb-temppool-buffer
Aug 14, 2026
Merged

NanoVDB: SyncFromAsync, ResourceRef, and resource seams for MeshToGrid and TempPool (CUDA)#2269
kmuseth merged 18 commits into
AcademySoftwareFoundation:masterfrom
harrism:nanovdb-temppool-buffer

Conversation

@harrism

@harrism harrism commented Aug 5, 2026

Copy link
Copy Markdown
Contributor

Follow-up to #2231 and #2251, part of #2232. This is B2 of step 2's retrofit.

Stacked on #2268 — it branches from that PR, so the diff shown here includes its commits until it merges. Review only the commit titled "add SyncFromAsync, and give MeshToGrid a resource seam".

SyncFromAsync<Derived>

A stream-ordered resource must also model the synchronous concept (is_async_resource<R> implies is_resource<R>, matching cuda::mr's refinement), so a custom resource has to write four methods where two would do. This CRTP base supplies the synchronous pair in terms of the stream-ordered one:

struct MyResource : nanovdb::cuda::SyncFromAsync<MyResource> {
    static constexpr size_t DEFAULT_ALIGNMENT = 256;
    void* allocate_async(size_t bytes, size_t alignment, cudaStream_t stream);
    void  deallocate_async(void* p, size_t bytes, size_t alignment, cudaStream_t stream);
};

The synchronous pair is deliberately not a bare delegate. Memory returned by allocate must be usable immediately on any stream, so the null-stream allocation is synchronized before it is returned — which is what DeviceResource::allocate already did, in a place nobody else could reuse. Omitting that synchronization produces memory that satisfies the concept but is not actually synchronous: a race, not a compile error. deallocate does not synchronize, because the synchronous concept's contract is that the memory is already quiescent.

The two test resources in TestMemoryResource.cu are the first users, and were already wrong in exactly this way — they provide only the async pair, so they never modelled is_async_resource. Nothing caught it because TempPool duck-types rather than asserting. They now derive from the mixin and conform.

MeshToGrid resource seam

MeshToGrid was the last builder allocating from a hard-wired DeviceResource, via the TempDevicePool alias. It now takes a ResourceT parameter (defaulted, so existing callers are unaffected) and threads it into both its TopologyBuilder and its pool.

As with Data in #2268, BoxTrianglePair is hoisted out of the class as MeshToGridBoxTrianglePair, with an in-class alias. It carries no dependence on the resource, and leaving it nested would give every ResourceT instantiation its own incompatible type — which the device functors, templated on BuildT alone, could not name.

ResourceRef, and the TempPool conversion it unblocks

The first attempt at converting TempPool failed its own unit tests, and the failure exposed a real gap rather than a bug in the conversion:

  • TempPool holds Resource* — non-owning, documented "must outlive this pool" — and TestMemoryResource's resources carry their counters inline, so they are stateful and lossy to copy.
  • cuda::Buffer holds R by value, so a pool holding a Buffer had to copy the resource, and the pool then recorded traffic into its copy while the tests asserted on the original.

Checking the model we borrowed showed we had adopted half of a two-part design. cuda::buffer genuinely owns its resource — its constructor says so in a static_assert: "Buffer owns a copy of the memory resource…" — but in CCCL the ownership semantics are selected by what is placed in the by-value slot: an owning any_resource, a borrowing resource_ref, or a refcounted shared_resource. We shipped the slot without the borrowing type, so a non-owning container had no legitimate way to express itself.

nanovdb::cuda::ResourceRef<R> is the missing piece — a non-owning reference that is itself a resource:

  • copying the ref shares the underlying instance, so a by-value Buffer over a ref borrows;
  • its async methods exist only when R models AsyncResource (enable_if-gated, the same pattern as Buffer's stream API), so a ref over a synchronous resource does not misreport its tier — covered by static_asserts over ResourceRef<PinnedResource>;
  • two refs compare equal exactly when they reference the same resource, which is the equality_comparable semantics CCCL's concepts require ("memory allocated by x may be deallocated by y");
  • it is the static, dependency-free analog of cuda::mr::resource_ref, and the same shape as std::pmr::polymorphic_allocator over memory_resource*.

TempPool now keeps its bytes in a Buffer<std::byte, ResourceRef<Resource>>: same Resource* contract, same stream retention, same discard-on-growth reallocation (via destroy(stream) + move-assignment — scratch is never prefix-copied), but the block is freed by ownership rather than by hand. The TempPool unit tests — the ones the by-value attempt failed — pass unchanged, which is the regression guard for the borrow semantics.

The ownership rule of record is written up on #2232: Buffer stays by-value; a concrete resource in the slot is owned, a ResourceRef borrows; a refcounted SharedResource analog is deferred until the Python-bindings need materializes.

Testing

New tests in TestBuffer.cu cover the mixin (supplies a working synchronous pair; a Buffer over a mixin-based resource round-trips; a mixin user models both concepts) and ResourceRef (traffic reaches the original stateful resource, not a copy; tier gating via PinnedResource; equality is identity). A pendingchanges note covers the builder seams, the mixin, and the ref.

Verified locally on an RTX 6000 Ada, CUDA 12.6:

  • nanovdb_cuda_buffer_unit_test27/27 (two new: a stateful resource observes all traffic through a by-value Buffer over a ref; ref equality is identity)
  • nanovdb_cuda_memory_resource_unit_test9/9, including the two TempPool tests the by-value conversion failed, unchanged
  • nanovdb_test_cuda52/53; the failure is UnifiedBuffer_IO, a missing test-data file (data/3_spheres.nvdb) in an unrelated component, which fails identically on an unmodified baseline.

Per #2264 the CUDA tests are excluded from CI, so CI will build but not run these.

harrism and others added 3 commits August 5, 2026 01:18
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>
@harrism
harrism requested a review from kmuseth as a code owner August 5, 2026 01:52
…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>
@harrism harrism changed the title NanoVDB: add SyncFromAsync, and give MeshToGrid a resource seam (CUDA) NanoVDB: SyncFromAsync, ResourceRef, and resource seams for MeshToGrid and TempPool (CUDA) Aug 5, 2026
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
@harrism

harrism commented Aug 5, 2026

Copy link
Copy Markdown
Contributor Author

@kmuseth ready for review after #2268 (stacked on it — the two commits after #2268's set are the payload). Adds SyncFromAsync, ResourceRef (the resource-ownership reconciliation recorded on #2232), MeshToGrid's resource seam, and TempPool onto cuda::Buffer. The TempPool unit tests pass unchanged, which is the regression guard for the borrow semantics.

harrism and others added 3 commits August 5, 2026 05:02
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>
@harrism
harrism force-pushed the nanovdb-temppool-buffer branch from e3c8930 to f619056 Compare August 5, 2026 05:06
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>
@harrism

harrism commented Aug 5, 2026

Copy link
Copy Markdown
Contributor Author

Copilot cycles (fork mirror: harrism#2). Net changes: TopologyBuilder's scratch now borrows its resource through ResourceRef instead of copying it — a rule-of-record alignment found while triaging (f619056); the builders assert the stream-ordered-resource requirement directly with a diagnostic naming the builder; a DEFAULT_ALIGNMENT >= alignof(uint64_t) assert closes the custom-resource alignment corner (2608ea8); and SyncFromAsync documents its synchronize cost.

harrism and others added 2 commits August 5, 2026 07:14
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>

@kmuseth kmuseth left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

looks good but I have two questions

Comment thread nanovdb/nanovdb/cuda/Buffer.h
Comment thread nanovdb/nanovdb/cuda/Buffer.h
harrism and others added 4 commits August 8, 2026 00:35
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>
…trofit

Signed-off-by: Mark Harris <mharris@nvidia.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
@harrism
harrism marked this pull request as draft August 8, 2026 00:59
harrism and others added 3 commits August 10, 2026 22:08
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>
…buffer

Signed-off-by: Mark Harris <mharris@nvidia.com>

# Conflicts:
#	nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh
#	nanovdb/nanovdb/unittest/TestBuffer.cu
@harrism
harrism marked this pull request as ready for review August 11, 2026 23:02
@harrism

harrism commented Aug 11, 2026

Copy link
Copy Markdown
Contributor Author

@kmuseth #2268 merged — thanks! I've merged latest master into this branch (no rebase, so the commits you looked at are unchanged) and the diff now shows only this PR's payload: SyncFromAsync, ResourceRef, TempPool/TopologyBuilder/MeshToGrid resource injection, and their tests.

Both of your comment threads were on Buffer.h lines that belonged to #2268's diff; that code merged there with the doxygen @deprecated tag on clear() included. Buffer.h is no longer part of this diff. Marking ready for review.

@kmuseth kmuseth left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

Approved!

Thanks for you changes/comments

I cloned your repo, build the branch locally and ran all nanovdb_test_cuda* unit-test on a Blackwell 6000. nanovdb_test_mgpu (still) fails, but it's clearly unrelated to your changes. The other tests passed, so I'll approve and merge

@kmuseth
kmuseth merged commit ba9a2c0 into AcademySoftwareFoundation:master Aug 14, 2026
19 checks passed
swahtz added a commit that referenced this pull request Aug 18, 2026
#2273)

Part of #2232 (B5), following #2268 which gave
tools::cuda::TopologyBuilder an injectable-resource seam and #2269 which
gave one to MeshToGrid.

The five classes that instantiate TopologyBuilder -- DilateGrid,
MergeGrids, PruneGrid, RefineGrid and CoarsenGrid, all in
nanovdb/tools/cuda/ -- were left behind as unaffected callers: each held
a TopologyBuilder<BuildT> constructed as mBuilder(stream), binding
default_resource<DeviceResource>() with no injection path. Downstream,
fvdb-core needs to route these operators' scratch through PyTorch's
c10::cuda::CUDACachingAllocator, and cannot delete its forked headers
until the seam exists.

Each of the five headers receives the conversion established by
MeshToGrid:

  - the class template gains typename ResourceT =
    nanovdb::cuda::DeviceResource, with an is_async_resource guard
    static_assert first in the class body;
  - the builder member becomes TopologyBuilder<BuildT, ResourceT>
    mBuilder;
  - the constructor(s) gain a trailing defaulted ResourceT& resource
    argument, forwarded as mBuilder(stream, resource). Trailing position
    keeps every existing call source-compatible. MergeGrids takes the
    argument on both its N-ary and its binary-convenience constructor,
    the delegating one forwarding;
  - every out-of-class member definition moves to the two-parameter
    form, including the dual-template getHandle.

No nested-type hoisting is needed: TopologyBuilder::Data was hoisted to
TopologyBuilderData in #2268, and none of the five classes declares its
own nested struct. The change is additive in behavior -- ResourceT
defaults everywhere, so existing callers compile unchanged with
identical allocation behavior.

Left out of scope: each operator's mBuilder.mProcessedRoot and
TopologyBuilder's internal mData are still host+device dual-space
nanovdb::cuda::DeviceBuffer, and the grid handle's output buffer still
goes through BufferT::create. None of these is observed by the injected
resource, which is expected and is noted in the new tests' comments; the
dual-space buffers are Step 3's problem.

Adds one injected-resource test per operator to TestMemoryResource.cu
(DilateGrid_InjectedResourceSeam, MergeGrids_InjectedResourceSeam,
PruneGrid_InjectedResourceSeam, RefineGrid_InjectedResourceSeam,
CoarsenGrid_InjectedResourceSeam), mirroring the existing PointsToGrid
seam tests: build a small ValueOnIndex source grid, run the operator
with a CountingResource passed as the new trailing constructor argument,
synchronize, and assert allocs > 0 && allocs == deallocs. MergeGrids
merges two disjoint grids, PruneGrid supplies a single all-on Mask<3>
sidecar over a one-leaf grid, and CoarsenGrid spreads voxels over a
2x2x2 leaf block so coarsening is non-degenerate.

---------

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
swahtz added a commit to openvdb/fvdb-core that referenced this pull request Aug 20, 2026
… 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>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants