Replace coordinate-list grid-topology construction with leaf-mask morphology by swahtz · Pull Request #712 · openvdb/fvdb-core

swahtz · 2026-07-31T10:27:48Z

Most of fvdb's sparse-grid topology builders construct their output grid by materializing an expanded int32 candidate-coordinate list — torch.empty({V·K, 3}) plus a per-candidate batch-index array — and then deduplicating it with a radix sort inside NanoVDB PointsToGrid. For a builder with a K-voxel stencil (K = 8 for dual/subdivide, 27 for a 3³ conv, V for coarsen) this costs hundreds to ~1000 transient bytes per output voxel, allocates raw cudaMalloc scratch that bypasses the torch caching allocator, and silently corrupts results once the candidate count crosses 2³¹ (an (int)pointCount cast in PointsToGrid.cuh).

This PR rewrites the device path of those builders to operate directly on the 512-bit leaf activity masks via nanovdb::tools::cuda::TopologyBuilder (DilateGrid / PruneGrid / CoarsenGrid / RefineGrid / MergeGrids) and the PadGrid driver introduced in PR #710. The candidate list — and its sort scratch — is never materialized. Result: 1–2 orders of magnitude less time and up to ~1500× less transient memory on the affected ops, and the >2³¹ overflow is removed.

It also hardens these ops (and the pre-existing dilate/prune/clone helpers they use) for sliced/indexed GridBatch views — a latent wrong-grid bug surfaced in review — via view-aware GridBatchData accessors and a compacting-copy helper (see Sliced / non-contiguous batch correctness below).

Stacked on PR #710 (dual-grid-leaf-mask-morphology), which fixed the flagship dual_grid case and added src/fvdb/detail/utils/nanovdb/PadGrid.cuh. Filed from issue #711.

Motivation

A user OOM'd a ~500M-voxel dual_grid on an 80 GB GPU. Root cause: the 8× candidate list alone is ~500M·8·16 B ≈ 64 GB of transient allocation before the grid is even built. #710 fixed dual_grid; #711 (this PR) applies the same leaf-mask approach to the rest of the builders that share the anti-pattern.

What changed

The mask-morphology rewrite is on the CUDA specializations; the CPU proxy-grid paths and the PrivateUse1/DistributedPointsToGrid legs keep their existing algorithm (the sliced-batch fix below does update the CPU per-item loops to be view-aware). For the hot cases the mask path is taken; non-power-of-two coarsen/subdivide factors and general strides keep the coordinate-list path as a correctness fallback.

Step	Op(s)	Change
1	shared infra	Reject `>2³¹` candidates in `_createNanoGridFromIJK` + dense volume guard; short-circuit `jidx_from_joffsets` for a single list (empty jidx, matching the existing single-tensor / `JIdxForGrid` convention); drop a redundant int64→int32 copy
2	`clipped_grid`	in-bounds bool mask → `PruneGrid` (was: active-in-bounds coord list → rebuild)
3	`from_nearest_voxels_to_points`	floored-points grid + one positive `PadGrid` pass (was: 8× candidate list); deletes `NearestIjkForPoints.{cu,h}`. Also fixes an `eidx*8` int32 overflow above ~268M points
4	`from_dense`	build the grid once and `GridHandle::copy` across the batch (was: re-sort identical coords `B` times)
5	`coarsened_grid`, `conv_grid`	coarsen 2ⁿ → repeated `CoarsenGrid`; conv 3³ s1 → `DilateGrid`, 2³ s1 → positive `PadGrid`; conv re-points its coarsen short-circuit at the new builder
6	`refined_grid`, `conv_transpose_grid`	subdivide 2ⁿ → `RefineGrid` (masked → `PruneGrid` then `RefineGrid`); conv-transpose s2/k3³ → `RefineGrid` + negative `PadGrid`
7	`from_mesh`	fixes int32 truncation of the surface-sample index (`tid`/`totalSamples`) and removes a dead `1 << 31` guard (signed-overflow UB) in `ijkForMesh`. Same parametric sampler — no topology change

Performance

Measured on an NVIDIA RTX PRO 6000 Blackwell (96 GB) with the same script run on main (before) and this branch (after). Inputs are solid S³ cubes built via from_ijk (worst-case dense topology) and random point clouds for nearest. Time is the median of 7 iters (warmup + cuda.synchronize(), input built outside the timed loop). Memory is peak torch high-water allocated during the op — i.e. the candidate tensor being removed — cross-checked against driver-level peak (polling mem_get_info() every 0.3 ms, which also captures the raw cudaMalloc sort scratch); on the old path the two agree to within ~1%, confirming the candidate list dominates.

"fit" scale — both paths succeed (apples-to-apples)

op	out voxels	time before → after	speedup	mem before → after	reduction
`dual_grid`	57,066,625	184.3 → 1.81 ms	102×	9.07 GB → 26.9 MB	337×
`from_nearest_voxels_to_points`	40,326,494	135.1 → 22.19 ms	6×	6.88 GB → 978 MB	7×
`conv_grid` k3 s1	11,543,176	246.0 → 1.85 ms	133×	11.23 GB → 8.6 MB	1302×
`conv_transpose_grid` k3 s1	11,543,176	237.1 → 1.86 ms	127×	9.72 GB → 8.6 MB	1126×
`conv_grid` k2 s1	11,390,625	66.5 → 1.84 ms	36×	3.33 GB → 5.8 MB	573×
`coarsened_grid` ×2	11,239,424	39.9 → 1.01 ms	39×	1.82 GB → 22.8 MB	80×
`refined_grid` ×2	191,102,976	88.9 → 1.54 ms	58×	3.83 GB → 49.2 MB	78×
`clipped_grid`	11,390,625	11.2 → 3.08 ms	4×	2.28 GB → 829 MB	3×

"big" scale — larger inputs; the trend holds and steepens

op	out voxels	time before → after	speedup	mem before → after	reduction
`dual_grid`	135,005,697	428.1 → 2.17 ms	197×	21.50 GB → 61.7 MB	348×
`from_nearest_voxels_to_points`	120,948,271	382.1 → 56.46 ms	7×	20.64 GB → 2.93 GB	7×
`conv_grid` k3 s1	24,389,000	519.7 → 1.68 ms	309×	23.87 GB → 15.9 MB	1502×
`conv_transpose_grid` k3 s1	24,389,000	489.7 → 1.69 ms	291×	20.65 GB → 15.9 MB	1299×
`conv_grid` k2 s1	24,137,569	143.3 → 1.64 ms	87×	7.08 GB → 12.3 MB	575×
`coarsened_grid` ×2	21,952,000	78.4 → 0.99 ms	79×	3.55 GB → 44.5 MB	80×
`refined_grid` ×2	452,984,832	203.4 → 1.88 ms	108×	9.07 GB → 111.8 MB	81×
`clipped_grid`	22,188,041	21.7 → 5.27 ms	4×	4.44 GB → 1.62 GB	3×

Reading the numbers

Transient memory scales linearly with input — that is the reported OOM. Old conv_grid k3 needs ~980 B of scratch per output voxel (11.5M → 11.2 GB; 24.4M → 23.9 GB); the new path is < 1 B/voxel. An ~80M-voxel conv OOMs an 80 GB card on main and takes < 20 MB here. dual_grid costs ~160 B/input-voxel old → ~500M voxels ≈ 80 GB → OOM, matching the original bug report.
Kernel-cubed builders win most (3³ conv materializes 27 candidates/voxel; dual/subdivide 8) — the mask path never forms that list.
clip (3–4×) and nearest (6–7×) are the smaller wins by design: their fast paths still touch per-voxel data (clip builds an in-bounds mask and prunes; nearest floors points to a grid then does one PadGrid pass), so they cut the constant factor rather than eliminate the list — while also fixing the int32 overflow in nearest.

Correctness & safety

Reordering is safe. All 8 builders are pure-forward topology construction (no autograd). Every differentiable consumer (pool / refine / gather-scatter conv / inject) maps voxels by IJK coordinate, not linear index, and TopologyBuilder enumerates in NanoVDB canonical order — the same order PointsToGrid produces — so the "feature row i ↔ voxel with getValue()==i+1" invariant is preserved. Tests assert ordered-ijk CUDA↔CPU parity plus a consumer round-trip per op.
clipGridFeaturesWithMask row alignment (features rmasked in input-active order vs the pruned grid) is explicitly tested — input-active order is canonical, so they align.
Non-power-of-two factors / general strides fall back to the coordinate-list path, so those cases are byte-for-byte unchanged.

Sliced / non-contiguous batch correctness

Reviewing the mask-morphology fast paths surfaced a latent bug shared with several pre-existing helpers. A sliced or indexed GridBatch (grid[idx], grid[a:b]) is a view: it shares the underlying NanoVDB handle (which keeps every grid) and only shrinks batchSize() and the per-grid metadata, so batchSize() < nanoGridHandle().gridCount() and each item's grid lives at its mCumBytes byte offset — not at physical index i. The fast paths resolved grids with mGridHdl->deviceGrid(i) (physical) and returned whole-handle copies, so on an indexed batch they silently produced the wrong grids or a voxelSizes/gridCount mismatch.

View-aware accessors on GridBatchData: deviceGridPtrAt(i) / hostGridPtrAt(i) return the i-th logical grid resolved by byte offset (the mapping Accessor::grid(i) already used), so op authors reach for the correct pointer by default instead of the deviceGrid(i) footgun. Every site — conv, conv_transpose, coarsen, subdivide, dilate, prune, across CUDA fast paths and CPU proxy loops — was retrofitted.
Compacting copy for identity / whole-copy paths: MakeContiguous now exposes contiguousGridHandle() (compact the selected grids: a per-grid byte copy + mGridIndex/mGridCount header fixup, no radix sort) and cloneGridHandleAt(); makeContiguous and cloneGrid build on them. cloneGrid had been copying every physical grid of a shared handle — the source of a dilated_grid(0) mismatch on a slice.
Zero-copy identity: coarsened_grid(1) / refined_grid(1) (unmasked) / dilated_grid(0) now return the input grid unchanged in Python, preserving its view and metadata with no copy (the C++ ops keep a compacting-copy fallback for direct callers).
Test: tests/unit/test_sliced_batch.py runs every op on tail/gap/reversed/single-item views (CPU+CUDA) against an independently-built contiguous reference.

Testing

Python: full pytest unit/ — 2775 passed, 10 skipped. Adds per-op CUDA↔CPU parity, consumer round-trip, edge-case (empty / empty-middle-item / single-voxel / leaf & root-tile boundary), peak-memory, and sliced/non-contiguous-batch (test_sliced_batch.py) tests.
C++: ctest — 40/41 pass. The one failure, PredGatherIGemmTest, is a pre-existing "Support for cp.async instructions has not been enabled" CUTLASS igemm error on this sm_120/Blackwell setup; this branch touches zero igemm/gather/CUTLASS files (git diff --name-only main...HEAD confirms), so it is unrelated.

Notes for the reviewer

Stacked on PR Build dual/padded grids via leaf-mask morphology to cut memory #710 (dual-grid-leaf-mask-morphology) — review after Build dual/padded grids via leaf-mask morphology to cut memory #710 merges; the Build dual/padded grids via leaf-mask morphology to cut memory #710 commits form the base of this branch.
The view-aware GridBatchData accessors and the compacting-copy helpers also fix the same sliced-batch bug in the pre-existing dilate/prune/clone helpers those ops route through.

dual_grid materialized an 8x-expanded candidate coordinate list before deduplicating it, costing ~420-450 bytes of transient memory per input voxel (and overflowing an int32 cast in NanoVDB's segmented radix sort above ~2^31 candidates). A ~500M-voxel grid OOMed on an 80GB GPU. Build the padded topology directly from the 512-bit leaf activity masks via nanovdb::tools::cuda::TopologyBuilder instead, so scratch is O(node count) rather than O(voxel count). - New src/fvdb/detail/utils/nanovdb/PadGrid.cuh: a one-sided (octant) analogue of nanovdb's DilateGrid. The internal-node functor reuses DilateInternalNodesFunctor's scatter verbatim with only the neighbor stencil swapped for a one-sided one; the leaf and erosion functors are the one-sided shift-OR / shift-AND derivations. - BuildPaddedGrid.cu CUDA path composes bmax positive + -bmin negative unit passes (pad, or erode + PruneGrid for exclude_border), preserving the primal/dual transform swap. dual_grid is one positive pass. - Guard the empty-erosion case (PruneGrid cannot build an empty grid); use GridHandle::copy for the bmin==bmax==0 identity. - Reject exclude_border on PrivateUse1 explicitly (was a latent crash) and require bmin <= 0 <= bmax. - Add a low-level build_padded_grid binding for the generic box, plus tests covering CUDA/CPU parity, leaf/root boundary crossings, empty batch items, erode-to-empty, and peak memory. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

buildPaddedGrid unconditionally swapped the source's primal/dual transforms in its metadata tail. That is correct for dual_grid, whose result voxels sit on the corner (dual) lattice, but wrong for the general build_padded_grid primitive exposed in this PR: a plain padded grid stays on the *same* lattice as its source, so build_padded_grid(0, 0) was not an identity (it shifted the origin by half a voxel) and every box pad silently moved to the dual lattice. Thread a dualTransform flag through ops::buildPaddedGrid: dual_grid passes true (keep the swap), build_padded_grid passes false (carry the source transforms verbatim). populateGridMetadata already recomputes the correct primal/dual transforms from the source's (voxelSize, origin), so the non-dual path just copies them over unchanged. Adds test_dual_grid_transform_is_dual_lattice (origin shifts by -0.5 voxel) and test_build_padded_grid_preserves_transform (origin/voxel size unchanged for every box), both on CPU and CUDA. Addresses the review on openvdb#710. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

The valueMask reference bound in PadInternalNodesFunctor was never read -- the padding stencil reads each child leaf's valueMask() directly -- so it only produced an unused-variable warning. The const_cast reference binding had no side effect, so drop it. Addresses the Copilot review note on openvdb#710. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

A sliced or indexed GridBatchData is a view: it shares the underlying NanoVDB handle (which keeps every grid) and only shrinks batchSize() and the metadata array, so batchSize() < nanoGridHandle().gridCount() and each item's grid is located by its mCumBytes byte offset, not by physical index. The CUDA path indexed the handle physically (mGridHdl->deviceGrid(i)) and the identity fast-path copied the whole handle, so on a sliced batch dual_grid returned the wrong grids and build_padded_grid pulled in the excluded siblings (with an out-of-bounds metadata read in the transform fix-up). The CPU path looped gridCount() with the same effect, corrupting the heap. main's CUDA path was correct here (it went through the Accessor); this restores that. - CUDA padding loop: resolve each source grid as bufBase + cumBytesAt(i) (the mapping Accessor::grid() uses), iterating batchSize(). No input copy. - CUDA identity (bmin==bmax==0): keep the one-shot whole-handle copy for a contiguous batch; for a sliced view rebuild from Accessor-mapped coordinates (only build_padded_grid(0,0) reaches this; dual_grid, being (0,1), never does). - CPU paths: iterate batchSize() and resolve grids by byte offset. Adds test_sliced_batch_padding_matches_contiguous (CPU+CUDA) over tail, gap, reversed, and single-item selections. Addresses the review on openvdb#710. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

…dexing PadLeafNodesFunctor and ErodeKeepMaskFunctor built a uint64_t(&)[10][3][3] reference by reinterpret_cast'ing the address of the middle element originalWordsShifted[1][1][1], then indexed it with negative and out-of-range subscripts (originalWords[-1][..], originalWords[8][..], originalWords[i-1][..][-1]). The notional [10][3][3] array does not exist at that address (it also extends past the real buffer), so the negative/out-of-range subscripting is undefined behavior an optimizer may miscompile. Replace the reference with a small centered-accessor lambda that indexes the real array with an explicit +1 offset -- originalWords(i, dBj, dBk) maps to originalWordsShifted[i + 1][dBj + 1][dBk + 1] -- so every access is in bounds. The computed addresses (and generated code) are identical; the CUDA<->CPU parity tests are unchanged. Addresses the review on openvdb#710. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

…grids Two defensive fixes from the openvdb#710 review, both on the exclude_border erosion path: - Allocate the per-leaf keep-mask sidecar as uint64 rather than uint8 so its data pointer is guaranteed 8-byte aligned for Mask<3> (8 uint64_t words), which the kernels dereference directly. torch's allocators over-align in practice, but the tensor API doesn't guarantee it. The emptiness check runs on a zero-copy uint8 view, since torch's any() (an 'or' reduction) isn't implemented for uint64. - In the erosion pass loop, if the first pass sees leafCount == 0 (a grid with voxels but no leaf nodes, e.g. a tile-only grid), push an explicit empty grid handle instead of the default-constructed one, which would crash downstream. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

The root-speculation step copies the source root+upper nodes to a host buffer with cudaMemcpyAsync(..., cudaMemcpyDeviceToHost, mStream) and then dereferences it on the host to enumerate tiles. nanovdb::HostBuffer::create() returns pageable memory today, which makes that D2H copy implicitly host-synchronous, so it works -- but a pinned host buffer would leave the copy in flight and the host read would race on uninitialized data. Add an explicit cudaStreamSynchronize(mStream) before the dereference so correctness no longer depends on the host-buffer allocation strategy. Cost is negligible: one small copy per grid. Addresses the review on openvdb#710. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

Copilot

Pull request overview

This PR removes the high-transient-memory “expanded candidate coordinate list + radix sort” approach from several CUDA sparse-grid topology builders by switching to NanoVDB leaf-mask morphology via nanovdb::tools::cuda::TopologyBuilder (including a new one-sided PadGrid primitive). It also hardens topology ops for sliced/indexed GridBatch views by introducing view-aware grid-pointer accessors and compacting copy helpers, and adds guards against NanoVDB’s int32 candidate-count overflow.

Changes:

Rework CUDA topology builders (conv_grid, conv_transpose_grid, coarsened_grid, refined_grid, from_nearest_voxels_to_points, clip_grid, and dense-grid batching) to use morphology/TopologyBuilder fast paths where applicable and fall back to coordinate lists otherwise.
Fix sliced/non-contiguous GridBatch correctness via GridBatchData::{deviceGridPtrAt,hostGridPtrAt} and handle-compacting helpers (contiguousGridHandle, cloneGridHandleAt).
Add overflow guards and extensive tests (CUDA↔CPU parity, sliced views, and peak-memory regressions).

Reviewed changes

Copilot reviewed 32 out of 32 changed files in this pull request and generated 1 comment.

Show a summary per file

File	Description
tests/unit/test_sliced_batch.py	New regression coverage for sliced/indexed `GridBatch` topology ops (CPU/CUDA).
tests/unit/test_dual.py	Expanded dual/padded-grid correctness, transform semantics, sliced-view parity, and CUDA/CPU parity tests.
tests/unit/test_basic_ops.py	Adds peak-memory and fast-path parity tests for nearest/clip/conv ops on CUDA.
src/tests/JaggedTensorTest.cpp	Updates expectations for single-list `jidx_from_joffsets` convention (empty `jidx`).
src/python/GridBatchOps.cpp	Exposes `build_padded_grid` binding; makes `dual_grid` call `buildPaddedGrid(..., dualTransform=true)`.
src/fvdb/JaggedTensor.cpp	Short-circuits `jidx_from_joffsets` to empty `jidx` for 0/1-list JaggedTensors.
src/fvdb/GridBatchData.h	Adds view-aware logical-grid pointer APIs (`deviceGridPtrAt`, `hostGridPtrAt`).
src/fvdb/GridBatchData.cu	Implements logical-grid pointer APIs via byte-offset resolution.
src/fvdb/detail/utils/nanovdb/PadGrid.cuh	New one-sided octant padding morphology implementation for NanoVDB device grids.
src/fvdb/detail/ops/NearestIjkForPoints.h	Deleted (no longer needed by nearest-voxel builder).
src/fvdb/detail/ops/NearestIjkForPoints.cu	Deleted (replaced by floored-point grid + `PadGrid`).
src/fvdb/detail/ops/MakeContiguous.h	Adds `contiguousGridHandle` and `cloneGridHandleAt` APIs for view-safe compacting copies.
src/fvdb/detail/ops/MakeContiguous.cu	Implements compacting-copy handle creation and refactors `makeContiguous` to use it.
src/fvdb/detail/ops/IjkForMesh.cu	Fixes indexing/truncation for large surface-sample counts using int64 tids + `packed_accessor64`.
src/fvdb/detail/ops/CloneGrid.cu	Fixes clone-on-view bug by compacting selected grids rather than copying full underlying handle.
src/fvdb/detail/ops/ClipGrid.cu	Switches clip topology construction to prune-by-mask instead of coord-list rebuild.
src/fvdb/detail/ops/BuildPrunedGrid.cu	Uses byte-offset logical-grid pointer accessors (CPU/CUDA).
src/fvdb/detail/ops/BuildPaddedGrid.h	Extends API with `dualTransform` toggle and documents semantics.
src/fvdb/detail/ops/BuildPaddedGrid.cu	Rewrites CUDA padded-grid builder to TopologyBuilder morphology (`PadGrid`/`PruneGrid`) and fixes view handling.
src/fvdb/detail/ops/BuildGridFromNearestVoxelsToPoints.cu	Replaces 8-candidate emission with floored points + one positive `PadGrid` pass.
src/fvdb/detail/ops/BuildGridFromIjk.cu	Adds >2³¹ per-grid candidate-count rejection to avoid NanoVDB silent corruption.
src/fvdb/detail/ops/BuildGridForConvTranspose.cu	Adds CUDA morphology/subdivide fast paths and uses logical-grid pointer accessors.
src/fvdb/detail/ops/BuildGridForConv.cu	Adds CUDA morphology/coarsen fast paths and uses logical-grid pointer accessors.
src/fvdb/detail/ops/BuildFineGridFromCoarse.h	Exposes `fineGridHandleFromCoarseCUDA` for reuse.
src/fvdb/detail/ops/BuildFineGridFromCoarse.cu	Implements RefineGrid-based power-of-two subdivision fast path (+ mask prune) and identity compact-copy.
src/fvdb/detail/ops/BuildDilatedGrid.cu	Fixes mixed-batch zero-dilation cloning for views using `cloneGridHandleAt`; uses logical-grid accessors.
src/fvdb/detail/ops/BuildDenseGrid.cu	Avoids repeated identical radix sorts by building once and copying across the batch; adds dense-volume guard.
src/fvdb/detail/ops/BuildCoarseGridFromFine.h	Exposes `coarseGridHandleFromFineCUDA` for reuse.
src/fvdb/detail/ops/BuildCoarseGridFromFine.cu	Implements CoarsenGrid-based power-of-two coarsening fast path + identity compact-copy.
src/CMakeLists.txt	Removes deleted `NearestIjkForPoints.cu` from build.
fvdb/grid_batch.py	Adds Python-side identity short-circuits (return `self`) for coarsen/refine/dilate where safe.
fvdb/_fvdb_cpp.pyi	Adds `build_padded_grid` to typing stubs.

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

The erode keep-mask emptiness check reduced over a uint8 view of the uint64 sidecar because any() (an 'or' reduction) is not implemented for uint64 on CUDA. (mask != 0).any() is equivalent and clearer -- the != yields a bool tensor whose any() reduction is implemented, so no dtype-reinterpreting view is needed. (.view(dtype) did compile and work; this just avoids the byte-reinterpretation.) Addresses the review on openvdb#712. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

…penvdb#711) Step 1 of the leaf-mask topology refactor (issue openvdb#711): cross-cutting hardening of the coordinate-list grid-construction path, independent of the per-op rewrites that follow. - Add a per-grid TORCH_CHECK(count <= INT32_MAX) in _createNanoGridFromIJK (BuildGridFromIjk.cu) and an unmasked-volume guard in BuildDenseGrid's checkInputs. NanoVDB's PointsToGrid radix sort casts the coordinate count to int32 (PointsToGrid.cuh:645), silently corrupting grids above ~2^31 candidates; these guards turn that into a clear error. - Short-circuit JaggedTensor::jidx_from_joffsets to an empty jidx for a single list (joffsets.size(0) <= 2), mirroring the single-tensor constructors and JIdxForGrid.cu. Every batchSize==1 op output was materializing (and often immediately discarding) a full int32 array of zeros plus a binary-search kernel; the empty-jidx convention is already honored by all consumers. Deferred: dropping the int64->int32 coordinate copy in BuildGridFromIjk.cu (needs validation of PointsToGrid's fancy-pointer/coalesced-load path; efficiency-only, low value) -- tracked as a follow-up. Full `pytest unit` passes (2763 tests) with no regressions from the jidx change. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

Step 2 of issue openvdb#711. clipGridWithMask computed a per-voxel in-bounds boolean mask, then materialized every active coordinate (activeGridCoords), masked it, and rebuilt the grid through createNanoGridFromIJK (a full radix sort with ~40 B/kept-voxel of raw cudaMalloc scratch). Clipping only ever keeps a subset of the source, so prune the source grid to the mask directly with pruneGrid (leaf-mask morphology, PruneGrid + a 64 B/leaf sidecar) -- no coordinate list, no sort. pruneGrid preserves the source transform and canonical voxel order, so the features rmask'd in clipGridFeaturesWithMask remain row-aligned with the clipped grid. Both clipGrid's helpers and pruneGrid are CPU+CUDA only; clip never supported PrivateUse1 (activeVoxelsInBoundsMask already TORCH_CHECKs it off), so no device coverage is lost. Existing test_clip_grid (single + batch, cpu+cuda) already pins counts and the exact gradient row-alignment; added test_clip_grid_mask_based_parity covering a non-dense grid with CUDA/CPU ordered-parity and a direct feature-alignment check. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

…b#711) Step 3 of issue openvdb#711. buildGridFromNearestVoxelsToPoints emitted the 8 corner voxels of each point's cell as an explicit coordinate list (int32[8N,3] + two 8N int32 batch-index arrays, ~160 B/point) and rebuilt the grid through a radix sort over 8N candidates -- and NearestIjkForPoints computed the write offset as `eidx * 8` in int32, silently wrapping above ~268M points, before the sort's own int32 count overflow. The 8 nearest voxels of a point are floor(p) + {0,1}^3, and Minkowski sum distributes over the point union, so the nearest-voxel grid equals the floored-point grid padded by one positive octant. Build the base grid from one voxel per point (floor), then apply a single positive PadGrid pass (leaf-mask morphology). Points are unstructured so one sort remains, but on N candidates instead of 8N -- which also lifts the overflow threshold 8x and removes the eidx*8 wrap entirely. Delete NearestIjkForPoints.{cu,h} (only used here). CPU path unchanged; the op has no PrivateUse1 leg. Existing set-equality tests (cpu+cuda, single+batch) already pin the exact floor(p)+{0,1}^3 topology; added a peak-memory test showing the torch-visible drop. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

…b#711) Step 4a of issue openvdb#711. The CUDA and PrivateUse1 dense-grid builders looped over the batch calling voxelsToGrid / DistributedPointsToGrid on the *same* coordinate list every iteration, re-running the radix sort (and its raw cudaMalloc scratch) batchSize times to produce batchSize identical grids. Every dense item is the same box (a mask, if provided, is shared across the batch), so build the grid once and GridHandle::copy it for the remaining items -- exactly what the kCPU path already does. This is the low-risk half of the dense rewrite; eliminating the coordinate list itself (an analytic-mask TopologyBuilder driver for the unmasked box) is a larger standalone change tracked separately. The Step 1 volume guard already turned the >2^31-cell overflow into a clear error. test_dense_interface (124) passes. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

Step 5a of issue openvdb#711 (Cluster A, coarsen half). buildCoarseGridFromFine materialized one coarse coordinate per fine voxel (floor(f/factor)) and rebuilt the grid through a radix sort. fvdb's mapping is floor(f/factor), and NanoVDB's CoarsenGrid maps f to floor(f/2) per pass (its coarsenComponent is floor(n/2) for all n, unioning each 2^3 fine block), so a uniform power-of-two factor is exactly that many CoarsenGrid passes -- leaf-mask morphology, no coordinate list, no sort. Non-power-of-two or non-uniform factors keep the coordinate path. Rewrites the kCUDA specialization only; kCPU (proxy grid) and kPrivateUse1 (coordinate list) are unchanged. The shared coarseIJKForFineGrid helper is retained (still used by the fallback and by buildGridForConv's short circuit -- conv is handled in a follow-up commit). Pooling relies on fine<->coarse correspondence by IJK coordinate, which is preserved; nn_modules pooling (18) and basic_ops pool/coarsen incl. forward+backward grad (36) pass. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

…nvdb#711) Step 5b of issue openvdb#711 (Cluster A, conv half). buildGridForConv emitted kernelVolume candidate coordinates per source voxel (int32[V*K,3] + jidx + a full N*K bool mask, then a compacting gather while the expanded tensors were still live) and rebuilt the grid through a radix sort -- ~1.9 KB of transient memory per input voxel for a 3x3x3 stride-1 conv, and an int32 overflow above ~79.5M voxels. For stride-1 convs the output is exactly S (+) window (the kernel offset box), so add leaf-mask fast paths: - kernel_size == 1 || stride == kernel_size: pure coarsening by stride -- reuse coarseGridHandleFromFineCUDA (CoarsenGrid for uniform power-of-two stride), extracted from BuildCoarseGridFromFine. - stride 1, uniform odd kernel k: (k-1)/2 symmetric DilateGrid passes. - stride 1, uniform even kernel k: (k-1) positive PadGrid passes. - everything else (non-uniform kernels, general strides > 1, which involve decimation rather than a clean morphology op): coordinate-list fallback. Conv topology feeds gather-scatter convolution by IJK coordinate, so the (canonical) enumeration order is preserved. kCPU proxy path unchanged; the op has no PrivateUse1 leg. conv_default/ground_truth/igemm/simple_unet (75) and nn_modules (68) pass forward+backward; added an explicit conv_grid cpu/cuda fast-path topology parity test. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

Step 6a of issue openvdb#711 (Cluster B, subdivide half). buildFineGridFromCoarse materialized factor^3 fine coordinates per coarse voxel (int32[V*K,3] + jidx, and for the masked path an int64 cumsum + a blocking .item() sync) and rebuilt the grid through a radix sort. fvdb subdivision maps coarse c to the block c*factor + [0, factor-1]^3, and NanoVDB's RefineGrid maps c to 2c + {0,1}^3 per pass, so a uniform power-of-two factor is that many RefineGrid passes -- leaf-mask morphology, no coordinate list, no sort. Masked subdivision prunes the coarse grid to the per-voxel mask first (PruneGrid) and refines the result. Non-power-of-two / non-uniform factors keep the coordinate path. The kCUDA body is exposed as fineGridHandleFromCoarseCUDA so buildGridForConvTranspose can reuse it for its (kernel_size == 1 || stride == kernel_size) subdivision short circuit (conv-transpose handled in a follow-up commit). Refine consumes the fine grid by IJK, so canonical order is preserved; refined_grid (incl. masked, factor-1, empty) and refine forward+backward pass (53 subdivide/refine + 5 consumer tests). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

…vdb#711) Step 6b of issue openvdb#711 (Cluster B, conv-transpose half). buildGridForConv- Transpose emitted kernelVolume candidate coordinates per source voxel (dstIjk = srcIjk*stride + offset) and rebuilt the grid through a radix sort. Add leaf-mask fast paths: - kernel_size == 1 || stride == kernel_size: pure subdivision by stride -- reuse fineGridHandleFromCoarseCUDA (RefineGrid for power-of-two stride). - stride 1, uniform kernel: S (+) window, same as forward conv -- odd kernel -> symmetric dilations, even kernel -> positive pad passes. - stride 2, kernel 3 (the classic upsampling conv-transpose): the output is 2S (+) [-1,1]^3; RefineGrid gives 2S (+) {0,1}^3 and one negative pad pass adds (+) {-1,0}^3, composing to (+) [-1,1]^3. - everything else (general strides, larger stride-2 kernels, non-uniform): coordinate-list fallback. Conv-transpose topology feeds gather-scatter conv by IJK, so canonical order is preserved. kCPU proxy path unchanged; no PrivateUse1 leg. conv_transpose default + ground_truth (62) pass forward+backward; added an explicit conv_transpose_grid cpu/cuda fast-path topology parity test. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

The CUDA surface-sample generator for from_mesh truncated its sample indices to int32: the per-thread `tid` and the `totalSamples` count (read from an int64 cumulative-sum) were both int32, silently wrapping once a mesh generates more than 2^31 surface samples. The dead `if (outIJK.numel() >= 1 << 31)` guard picked between two byte-identical launch branches and relied on `1 << 31` signed-overflow UB. Widen `tid`/`numTris`/`totalSamples` to int64 (the accessors already use 64-bit indexing) and collapse the duplicated launch into a single path. No behavior change below 2^31 samples; correct above it. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

Step 1 made jidx_from_joffsets short-circuit a single-list joffsets to an empty jidx, matching the established convention: the single-tensor JaggedTensor constructors and JIdxForGrid already return an empty jidx for one list, and consumers treat an empty jidx as every element mapping to batch 0. StaticUtilityFunctions still asserted the naive full-length-zeros result for that edge case; update it to expect the empty tensor. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

A sliced/indexed GridBatch (grid[idx], grid[a:b]) is a view: it shares the underlying handle (which keeps every grid) and only shrinks batchSize() and the per-grid metadata, so batchSize() < nanoGridHandle().gridCount() and each item's grid is located by its mCumBytes byte offset, not by physical index. The mask-morphology fast paths resolved grids with mGridHdl->deviceGrid(i) (physical) and returned whole-handle copies, so on an indexed batch they silently produced the wrong grids or a voxelSizes/gridCount mismatch -- the same bug class fixed in buildPaddedGrid (openvdb#710), across the rest of the openvdb#711 ops and the pre-existing dilate/prune/clone helpers they use. - GridBatchData gains view-aware accessors deviceGridPtrAt(i)/hostGridPtrAt(i): the i-th *logical* grid resolved by byte offset (like Accessor::grid(i)). Every physical-index site (conv, conv_transpose, coarsen, subdivide, dilate, prune -- CUDA fast paths and CPU proxy loops) now uses them. - MakeContiguous exposes contiguousGridHandle() (compact the selected grids: a per-grid byte copy + mGridIndex/mGridCount header fixup, no radix sort) and cloneGridHandleAt(); makeContiguous and cloneGrid build on them. cloneGrid had copied *every physical grid* of a shared handle -- wrong for a sliced batch. - Zero-copy identity: coarsened_grid(1) / refined_grid(1) / dilated_grid(0) now return the input grid unchanged in Python, preserving its view and metadata with no copy (the C++ ops keep a compacting-copy fallback for direct callers). Adds test_sliced_batch.py: every op on tail/gap/reversed/single-item views (CPU+CUDA) matches an independently-built contiguous reference. Full unit suite green (2775 passed). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

Copilot

Pull request overview

Copilot reviewed 32 out of 32 changed files in this pull request and generated no new comments.

Suppressed comments (3)

src/fvdb/detail/ops/MakeContiguous.cu:92

cloneGridHandleAt() has the same assumption as contiguousGridHandle(): it calls getCurrentCUDAStream(input.device().index()) for CUDA without checking has_index(). Add the same guard to avoid index == -1 cases.

    const bool isCpu = input.device().is_cpu();
    cudaStream_t stream =
        isCpu ? cudaStream_t(0) : at::cuda::getCurrentCUDAStream(input.device().index()).stream();
    uint8_t *dst       = isCpu ? buffer.data() : buffer.deviceData();

src/fvdb/detail/ops/MakeContiguous.cu:64

contiguousGridHandle() uses at::cuda::getCurrentCUDAStream(input.device().index()) when the device is CUDA. If the device is constructed without an explicit index (index == -1 / has_index() == false), this can fail or pick an unintended stream. Add an explicit guard (consistent with other CUDA code paths in this repo) before calling getCurrentCUDAStream.

This issue also appears on line 89 of the same file.

    const bool isCpu = input.device().is_cpu();
    cudaStream_t stream =
        isCpu ? cudaStream_t(0) : at::cuda::getCurrentCUDAStream(input.device().index()).stream();
    uint8_t *dstBase       = isCpu ? buffer.data() : buffer.deviceData();

src/fvdb/detail/ops/BuildPaddedGrid.cu:347

In the CUDA identity case (bmin==bmax==0) for a non-contiguous/sliced batch, this falls back to rebuilding via paddedIJKForGrid + _createNanoGridFromIJK, which reintroduces the expensive coordinate-list + radix-sort path even though the result is just a compacted copy of selected grids. Since this PR introduces contiguousGridHandle() for exactly this use-case, prefer using it here to avoid the sort and the joffsets().cpu() sync.

        if (baseBatchHdl.isContiguous()) {
            return baseBatchHdl.nanoGridHandle().copy<TorchDeviceBuffer>(guide);
        }
        return ops::_createNanoGridFromIJK(
            paddedIJKForGrid<torch::kCUDA>(baseBatchHdl, nanovdb::CoordBBox::createCube(0, 0)));

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

buildPaddedGrid pioneered the byte-offset grid resolution (bufBase + cumBytesAt(i)) before the view-aware GridBatchData accessors existed. Now that deviceGridPtrAt(i)/hostGridPtrAt(i) are available, use them here as well -- the three sites (both CPU proxy loops and the CUDA dispatch) computed exactly what the accessors return, so this is a behavior-preserving DRY cleanup that routes every op through the one correct implementation. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

The CUDA identity case (bmin==bmax==0) on a sliced/non-contiguous batch rebuilt topology via paddedIJKForGrid + _createNanoGridFromIJK, which reintroduces the coordinate-list + radix-sort path and a joffsets().cpu() sync -- even though the result is just a compacted copy of the selected grids. Use contiguousGridHandle() (byte copy + header fix-up, no sort, no sync) instead, matching the coarsen/subdivide identity paths. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

Copilot

Pull request overview

Copilot reviewed 32 out of 32 changed files in this pull request and generated no new comments.

Suppressed comments (1)

src/fvdb/detail/ops/BuildGridFromIjk.cu:248

The new >2^31-candidate guard copies ijk.joffsets() to CPU here, but the CUDA and PrivateUse1 dispatch paths already do their own ijk.joffsets().cpu() copy (see earlier in this same file). This adds a second host sync/copy per call on the hot device paths. Consider moving the guard into the device-specific dispatch functions (reusing their existing ijkBOffsetTensor) or otherwise reusing a single host copy of the offsets.

Also, this block now uses std::numeric_limits—please ensure <limits> is included explicitly in this translation unit rather than relying on transitive includes.

    {
        const torch::Tensor joffsetsCpu = ijk.joffsets().cpu();
        const auto joffsetsAcc          = joffsetsCpu.accessor<fvdb::JOffsetsType, 1>();
        for (int64_t gi = 0; gi + 1 < joffsetsCpu.size(0); gi += 1) {
            const int64_t nCoords = joffsetsAcc[gi + 1] - joffsetsAcc[gi];

The >2^31-candidate guard in _createNanoGridFromIJK did its own ijk.joffsets().cpu() copy, but the CUDA and PrivateUse1 dispatch paths already copy joffsets to host -- so the guard added a second host sync per call on the hot device paths. Extract the check into a file-local helper (checkCandidateCountsFitInt32) and call it inside each dispatch, reusing the host offsets copy each already makes: one sync per path, no logic duplication, same universal guarantee. Also include <limits> explicitly rather than relying on a transitive include for std::numeric_limits. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

swahtz and others added 4 commits July 31, 2026 16:10

Formatting fixes

49912e5

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

Address review: fix 1<<31 UB and bitwise & in PadGrid root speculation

a496cf2

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

format fix

f152909

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

swahtz requested a review from a team as a code owner July 31, 2026 10:27

swahtz requested review from matthewdcong and sifakis July 31, 2026 10:27

swahtz changed the title ~~cccccbkrnntffrnvnffkgggfrgbnnldijjrdufufthtfIssue 711 leaf mask topology~~ Replace coordinate-list grid-topology construction with leaf-mask morphology Jul 31, 2026

swahtz marked this pull request as draft July 31, 2026 10:34

swahtz added optimization Performance or memory optimization Topology Operations Issues related to topology operations (prune, merge, dilate, etc. labels Jul 31, 2026

swahtz force-pushed the issue-711-leaf-mask-topology branch from 903ae2c to 35fded2 Compare August 1, 2026 06:23

swahtz self-assigned this Aug 1, 2026

swahtz and others added 5 commits August 1, 2026 18:50

swahtz added this to the v0.6 milestone Aug 3, 2026

swahtz force-pushed the issue-711-leaf-mask-topology branch from 35fded2 to 5878cec Compare August 3, 2026 01:24

swahtz requested a review from Copilot August 3, 2026 01:39

Copilot started reviewing on behalf of swahtz August 3, 2026 01:40 View session

Copilot AI reviewed Aug 3, 2026

View reviewed changes

Comment thread src/fvdb/detail/ops/BuildPaddedGrid.cu Outdated

swahtz and others added 6 commits August 3, 2026 13:53

swahtz force-pushed the issue-711-leaf-mask-topology branch from 5878cec to 11fd2eb Compare August 3, 2026 01:55

swahtz requested a review from Copilot August 3, 2026 01:57

Copilot started reviewing on behalf of swahtz August 3, 2026 01:57 View session

Copilot AI reviewed Aug 3, 2026

View reviewed changes

swahtz and others added 5 commits August 3, 2026 14:05

Remove verbose gridPtrAt comment

4ca6625

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

Comment cleanups

1f884ef

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

format fix

24d06a4

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

swahtz requested a review from Copilot August 3, 2026 02:35

Copilot started reviewing on behalf of swahtz August 3, 2026 02:36 View session

Copilot AI reviewed Aug 3, 2026

View reviewed changes

swahtz linked an issue Aug 3, 2026 that may be closed by this pull request

Grid topology ops build coordinate lists where leaf-mask operations would be far cheaper #711

Open

swahtz mentioned this pull request Aug 3, 2026

Generalize leaf-mask morphology to arbitrary factors: non-power-of-two coarsen/subdivide and general conv strides still pay the coordinate-list + radix-sort path #716

Open

swahtz mentioned this pull request Aug 3, 2026

GridEdgeNetwork emits 8 vertices + 12 edges per voxel with no deduplication (~368 B/voxel) #720

Open

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

Replace coordinate-list grid-topology construction with leaf-mask morphology - #712

Replace coordinate-list grid-topology construction with leaf-mask morphology#712
swahtz wants to merge 28 commits into
openvdb:mainfrom
swahtz:issue-711-leaf-mask-topology

swahtz commented Jul 31, 2026 •

edited

Loading

Uh oh!

Copilot AI left a comment

Uh oh!

Uh oh!

Copilot AI left a comment

Uh oh!

Copilot AI left a comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Uh oh!

Conversation

swahtz commented Jul 31, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Motivation

What changed

Performance

"fit" scale — both paths succeed (apples-to-apples)

"big" scale — larger inputs; the trend holds and steepens

Reading the numbers

Correctness & safety

Sliced / non-contiguous batch correctness

Testing

Notes for the reviewer

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Reviewed changes

Uh oh!

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

swahtz commented Jul 31, 2026 •

edited

Loading