Skip to content

NanoVDB: run the builders on a synchronous memory resource (CUDA) - #2272

Merged
kmuseth merged 42 commits into
AcademySoftwareFoundation:masterfrom
harrism:nanovdb-sync-resource-builders
Aug 18, 2026
Merged

NanoVDB: run the builders on a synchronous memory resource (CUDA)#2272
kmuseth merged 42 commits into
AcademySoftwareFoundation:masterfrom
harrism:nanovdb-sync-resource-builders

Conversation

@harrism

@harrism harrism commented Aug 5, 2026

Copy link
Copy Markdown
Contributor

Follow-up to #2231/#2251, part of #2232. This closes the scratch half of step 2's acceptance criterion: "NanoVDB builds and runs on a pool-less vGPU via an injected synchronous cudaMalloc-backed resource."

Stacked on #2270 — review the last commit.

What

The builders require a stream-ordered resource, which a device without memory-pool support (cudaDevAttrMemoryPoolsSupported == 0, e.g. AWS g6f, cf. #2255) cannot provide through cudaMallocAsync — the wrappers fail loudly there by design (#2256). Two additions close the gap as pure library code:

  • cuda::MallocResource — a synchronous cudaMalloc/cudaFree resource that works on any device. Models is_resource only.
  • cuda::AsyncFromSync<R> — presents a synchronous resource as a stream-ordered one; the mirror of SyncFromAsync, and the analog of cuda::mr's synchronous_resource_adapter. allocate_async forwards — synchronously allocated memory is already valid on every stream, a stronger guarantee than stream-ordering requires. deallocate_async synchronizes the stream first, which makes the synchronous contract's quiescence requirement hold.

So the pool-less path is PointsToGrid<BuildT, AsyncFromSync<MallocResource>> (or an AsyncFromSync<ResourceRef<YourArena>> for a stateful arena), with no build flags.

Why a wrapper rather than dispatch inside the builders

The roadmap's rule is that a synchronous arena is never a silent allocate_async facade — a facade misrepresents its semantics and hands multi-stream callers unexpected serialization. AsyncFromSync is the explicit, honest version of the same lift: every deallocation synchronizes its stream, that cost is documented at the type, and the caller chooses it. This is also exactly the design CCCL settled on (synchronous_resource_adapter synchronizes before deallocate_sync when wrapping a synchronous resource). The alternative — if constexpr dispatch at every builder release site — spreads the same synchronization across ~30 sites instead of one audited method.

What this does not close

The grid handle's buffer still goes through GridHandle<BufferT>'s buffer type (default DeviceBuffercudaMallocAsync), so a fully flag-free pool-less deployment still needs NANOVDB_USE_SYNC_CUDA_MALLOC for that one allocation until step 3 delivers GridHandle<cuda::Buffer<std::byte,R>>. Scratch — the overwhelming majority of allocations — is closed here.

Testing

New TestMemoryResource.PointsToGrid_RunsOnSynchronousResource: drives PointsToGrid end-to-end with AsyncFromSync<ResourceRef<SyncCountingResource>> — a stateful, synchronous-only resource. Asserts the grid builds, that scratch really routed through the synchronous resource (allocs > 0), and that every allocation was freed through the caller's instance (allocs == deallocs). On pooled hardware this path bypasses pools entirely, which is precisely what a pool-less device requires, so the dispatch is exercised without vGPU hardware; final validation on a real g6f would be a good CI follow-up once #2264 lands. Plus static_asserts pinning the tiers (MallocResource is sync-only; the wrapper models async).

Suites: memory-resource 10/10, TestBuffer 27/27, full CUDA 52/53 (pre-existing UnifiedBuffer_IO data-file failure). Per #2264, CI builds but does not run these — numbers are from a local RTX 6000 Ada, CUDA 12.6.

harrism and others added 13 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>
…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 builders require a stream-ordered resource, which a device without
memory-pool support cannot provide through cudaMallocAsync. Two additions
close the gap. MallocResource is a synchronous cudaMalloc/cudaFree
resource that works on any device. AsyncFromSync presents a synchronous
resource as a stream-ordered one -- the mirror of SyncFromAsync, and the
analog of cuda::mr's synchronous_resource_adapter: allocate_async
forwards, since synchronously allocated memory is already valid on every
stream, and deallocate_async synchronizes the stream first, which makes
the synchronous contract's quiescence requirement hold. The serialization
cost of that synchronize is documented at the type, so a synchronous
backend is an explicit caller choice rather than a silent substitution.

The new test drives PointsToGrid end-to-end with an injected
AsyncFromSync over a stateful synchronous resource: every scratch
allocation routes through cudaMalloc and is freed through the caller's
instance, never touching stream-ordered allocation -- which is exactly
what a pool-less device requires of the scratch path.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
harrism and others added 4 commits August 5, 2026 05:29
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>
…-builders

Signed-off-by: Mark Harris <mharris@nvidia.com>
Null-free is a no-op everywhere else, so there is no quiescence to
establish and no reason to synchronize the stream for it. Also correct
the sync-resource test's comment: the grid handle's output buffer is
allocated through BufferT, not the injected resource, so "every builder
allocation" overstated what the test demonstrates -- it is the scratch
allocations that never touch stream-ordered allocation.

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#4). Net changes (aac7f00): the sync-resource test's comment now says scratch allocations and names the grid handle's output buffer as the exception, and AsyncFromSync::deallocate_async treats null as a no-op rather than synchronizing for it. Second cycle returned no findings.

harrism and others added 4 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>
Signed-off-by: Mark Harris <mharris@nvidia.com>
…-builders

Signed-off-by: Mark Harris <mharris@nvidia.com>
harrism and others added 8 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>
Signed-off-by: Mark Harris <mharris@nvidia.com>
…-builders

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>
…-builders

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 10 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>
Signed-off-by: Mark Harris <mharris@nvidia.com>
…-builders

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
…urce-builders

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

# Conflicts:
#	pendingchanges/nanovdb.txt
@harrism
harrism marked this pull request as ready for review August 17, 2026 20:10
@harrism

harrism commented Aug 17, 2026

Copy link
Copy Markdown
Contributor Author

@kmuseth #2270 merged — thanks! Latest master merged in (no rebase) and the diff is now just this PR's payload, the last layer of the stack: nanovdb::cuda::MallocResource (a synchronous cudaMalloc-backed resource for devices without memory-pool support) and nanovdb::cuda::AsyncFromSync (presents any synchronous resource as stream-ordered by synchronizing before each deallocation), plus a test running PointsToGrid end-to-end on an injected synchronous resource. Marking ready for review.

…urce-builders

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

# Conflicts:
#	pendingchanges/nanovdb.txt

@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 - I assume this wraps up the major changes to the cuda memory allocation

@kmuseth
kmuseth merged commit 3f35101 into AcademySoftwareFoundation:master Aug 18, 2026
15 checks passed
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