Build dual/padded grids via leaf-mask morphology to cut memory

[swahtz]

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 python binding for the generic box, plus tests covering CUDA/CPU parity, leaf/root boundary crossings, empty batch items, erode-to-empty, and peak memory.

[Copilot]

Pull request overview

This PR reworks the CUDA implementation of dual_grid / padded-grid construction to avoid materializing an 8× expanded coordinate list by instead building padded topology directly from NanoVDB leaf activity masks via TopologyBuilder, significantly reducing transient memory usage and avoiding large-candidate overflows.

Changes:

• Adds a new device-side one-sided morphology implementation (PadGrid.cuh) to build padded/dual topologies from leaf masks.
• Updates the CUDA padded-grid path (BuildPaddedGrid.cu) to use mask-morphology passes (and erosion+prune for exclude_border), with explicit guards for identity and empty-erosion cases.
• Exposes a low-level Python binding build_padded_grid(...) and adds CUDA/CPU parity + boundary + memory tests.

Reviewed changes

Copilot reviewed 5 out of 5 changed files in this pull request and generated 3 comments.

Show a summary per file

File	Description
tests/unit/test_dual.py	Adds CUDA/CPU parity tests for dual_grid and generic padded boxes, plus boundary/empty/memory coverage.
src/python/GridBatchOps.cpp	Exposes build_padded_grid binding and routes dual_grid through buildPaddedGrid.
src/fvdb/detail/utils/nanovdb/PadGrid.cuh	New one-sided (octant) mask-morphology topology builder for padded/dual grids on CUDA.
src/fvdb/detail/ops/BuildPaddedGrid.cu	Switches CUDA implementation to TopologyBuilder-based mask morphology; adds erosion+prune path and guards.
fvdb/_fvdb_cpp.pyi	Adds the build_padded_grid type stub.

---

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

[Copilot]

Pull request overview

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

Suppressed comments (1)

src/python/GridBatchOps.cpp:675

• build_padded_grid is exposed as a general padding primitive, but it delegates to ops::buildPaddedGrid, which always swaps primal/dual transforms (see BuildPaddedGrid.cu where mDualTransform/mPrimalTransform are swapped). For callers, build_padded_grid(bmin=0,bmax=0) is therefore not a pure identity (it still swaps transforms), which is surprising given the name/signature. Please either document this behavior on the binding or rename the binding to make the dual-grid semantics explicit.

        [](const GBI &grid, int bmin, int bmax, bool excludeBorder) {
            if (grid.batchSize() == 0) {
                return fvdb::detail::makeEmptyGridBatchData(grid.device());
            }
            return ops::buildPaddedGrid(grid, bmin, bmax, excludeBorder);

[swahtz]

Re: the suppressed Copilot note on GridBatchOps.cpp — that build_padded_grid inherits buildPaddedGrid's unconditional primal/dual transform swap, so build_padded_grid(bmin=0, bmax=0) isn't a pure identity.

Good catch, and it's a real correctness bug rather than just a naming/doc nit. The swap is only correct for dual_grid, whose result voxels sit on the corner (dual) lattice; a plain padded grid stays on the same lattice as its source and must keep the source's transforms.

Fixed: ops::buildPaddedGrid now takes a dualTransform flag.

• dual_grid → true: keeps the primal↔dual swap (origin shifts by −½ voxel). Behavior unchanged.
• build_padded_grid → false: carries the source transforms verbatim, so build_padded_grid(0, 0) is a true identity and every box pad stays on the source lattice.

populateGridMetadata already recomputes the correct primal/dual transforms from the source's (voxelSize, origin), so the non-dual path just copies them across unchanged.

Added regression tests on CPU + CUDA: test_dual_grid_transform_is_dual_lattice (dual origin == base − ½·voxel_size) and test_build_padded_grid_preserves_transform (origin and voxel size unchanged for every box, including asymmetric [0, 2] / [-2, 1]). Full unit suite green.

[Copilot]

Pull request overview

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

Suppressed comments (1)

src/fvdb/detail/utils/nanovdb/PadGrid.cuh:179

• valueMask is declared but never used in PadInternalNodesFunctor. This will trigger an unused-variable warning (and can fail builds if warnings are treated as errors). Remove it or explicitly mark it unused.

        const auto &srcTree = srcGrid->tree();
        const auto &lower   = srcTree.template getFirstNode<1>()[lowerID];
        auto &valueMask     = const_cast<LowerMaskT &>(lower.valueMask());

[swahtz]

Re: the suppressed Copilot note on PadGrid.cuh — valueMask declared but never used in PadInternalNodesFunctor.

Fixed — removed it. The stencil reads each child leaf's mask directly (padNeighborMaskStencil<Positive>(leaf.valueMask())), so the outer const_cast reference to the lower node's valueMask() was dead and had no side effect. Clean build, unit suite green.

[Copilot]

Pull request overview

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

Suppressed comments (1)

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

• This retrieves the per-item grid via mGridHdl->deviceGrid(i), which assumes batch item i corresponds to the i-th grid in the underlying handle. That is not true for non-contiguous/sliced GridBatchData (which uses per-item mCumBytes offsets; see GridBatchData::Accessor::grid). Use the batch metadata’s byte offsets to compute the grid pointer so CUDA padding works correctly on indexed batches too.

        nanovdb::OnIndexGrid *grid = baseBatchHdl.mGridHdl->deviceGrid<nanovdb::ValueOnIndex>(i);
        TORCH_CHECK(grid, "Grid is null");

[Copilot]

Pull request overview

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

Suppressed comments (1)

src/fvdb/detail/utils/nanovdb/PadGrid.cuh:534

• Same UB pattern as in PadLeafNodesFunctor: originalWords is defined via reinterpret_cast<WordStencilT>(originalWordsShifted[1][1][1]) and later indexed with negative indices (originalWords[-1][…], originalWords[i][…][-1], etc.). Because the reference does not denote an actual [10][3][3] array object, this relies on undefined behavior and could be miscompiled.

Recommend switching to direct indexing into originalWordsShifted with explicit +1 offsets (or an accessor helper) so all indices stay within the declared array object.

        uint64_t originalWordsShifted[10][3][3] = {};
        using WordStencilT                      = uint64_t(&)[10][3][3];
        auto &originalWords = reinterpret_cast<WordStencilT>(originalWordsShifted[1][1][1]);

[Copilot]

Pull request overview

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

Suppressed comments (2)

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

• When excludeBorder==true, a grid with numVoxelsAt(i) > 0 but leafCount == 0 (e.g. a grid represented by value tiles, or otherwise malformed metadata) will hit break before handle is assigned. The code then pushes a default-constructed GridHandle, which can lead to crashes later.

Even if this situation “shouldn’t happen”, it’s safer to make the leafCount==0 branch produce an explicit empty handle (or throw) so behavior is well-defined.

                              .mNodeCount[0]
                        : baseBatchHdl.numLeavesAt(i);
                if (leafCount == 0) {
                    break; // already eroded to empty; further erosion is a no-op
                }

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

• keepTensor is allocated as a uint8 tensor and then reinterpreted as nanovdb::Mask<3>*. This relies on the returned pointer being suitably aligned for Mask<3> (likely 8-byte alignment); if it isn't, dereferencing in the CUDA kernels is undefined behavior and can cause misaligned global memory accesses.

Allocate the sidecar as a torch::kUInt64 tensor (or otherwise ensure alignment) before casting to Mask<3>*.

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

    const int64_t maskBytes = static_cast<int64_t>(sizeof(nanovdb::Mask<3>)) * leafCount;
    torch::Tensor keepTensor =
        torch::empty({maskBytes}, torch::TensorOptions().dtype(torch::kUInt8).device(device));
    auto *keepMasks = reinterpret_cast<nanovdb::Mask<3> *>(keepTensor.data_ptr());

[swahtz]

Addressed the two suppressed Copilot notes on BuildPaddedGrid.cu (in 39847f2):

• Erosion leafCount == 0: a grid with voxels but no leaf nodes (a tile-only grid) would break before handle was assigned and push a default-constructed GridHandle. It now pushes an explicit createEmptyGridHandle in that case, so the result is always well-defined.
• kUInt8 keep-mask sidecar → Mask<3>*: now allocated as kUInt64 so the data pointer is guaranteed 8-byte aligned for Mask<3> (8 uint64 words). The emptiness check runs on a zero-copy uint8 view, since torch's any() (an or reduction) isn't implemented for uint64. (There's a single such sidecar; the "also on line 372" reference didn't correspond to a second allocation.)

test_dual (including the exclude_border erosion paths) and the full unit suite are green.

[Copilot]

Pull request overview

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

Suppressed comments (2)

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

• numNegative is computed as -bmin using int arithmetic. If a caller passes bmin == INT_MIN, this negation overflows (undefined behavior) before any range checks can catch it, and totalPasses becomes incorrect. Compute counts in a wider signed type and validate that the pass count is non-negative and fits in an int loop bound.

    const int numPositive = bmax;
    const int numNegative = -bmin;
    const int totalPasses = numPositive + numNegative;

src/fvdb/detail/utils/nanovdb/PadGrid.cuh:720

• cudaMemcpyAsync(..., cudaMemcpyDeviceToHost, mStream) is followed by immediate dereference of srcRootAndUpperBuffer on the host without synchronizing the stream. This is only safe if HostBuffer::create() always returns pageable memory (making the copy implicitly synchronous); if it ever returns pinned memory, this becomes a real race and can read uninitialized data. Prefer making the copy explicitly synchronous (either cudaMemcpy or cudaStreamSynchronize(mStream) after the async copy) to keep correctness independent of the host-buffer allocation strategy.

        auto srcRootAndUpperBuffer = nanovdb::HostBuffer::create(rootAndUpperSize);
        cudaCheck(cudaMemcpyAsync(srcRootAndUpperBuffer.data(),
                                  deviceSrcRoot,
                                  rootAndUpperSize,
                                  cudaMemcpyDeviceToHost,
                                  mStream));
        auto srcRootAndUpper = static_cast<RootT *>(srcRootAndUpperBuffer.data());

        for (uint32_t t = 0; t < srcRootAndUpper->tileCount(); t++) {