NanoVDB: batch-capable TopologyBuilder (CUDA)

[swahtz]

Summary

tools::cuda::TopologyBuilder can now build a batch of B independent grids into one buffer in a single pass of kernel launches, laid out back to back with mGridIndex = g and mGridCount = B. Single-grid builds are unchanged: every consumer (RefineGrid, CoarsenGrid, DilateGrid, PruneGrid, MergeGrids, MeshToGrid) compiles without edits and produces byte-identical output.

Motivation: downstream users that build many small grids per step (fVDB's generated convolution topologies, openvdb/fvdb-core#785) currently loop over members, paying several stream synchronizations per grid plus one more per grid to merge the handles. With this change they run the builder once over the batch and pay the synchronizations once.

What changed

TopologyBuilderData gains the grid's batch position (gridIndex, gridCount) and the first index of its tiles and nodes in the batch-wide scratch arrays (tileCount, tileBase, upperBase, lowerBase, leafBase, processedRootOffset). The builder holds one Data per grid (pinned host staging plus a device array); data() still means grid 0.

Processed roots are staged per grid in one pinned buffer. allocateProcessedRoots(tileCounts) lays out B roots; allocateProcessedRoot(bytes) is the single-grid form and resets the batch to one grid. uploadProcessedRoot also uploads a tileToGrid table for batches. hostProcessedRoot(g) / deviceProcessedRoot(g) default to grid 0.

countNodes keeps its batch-wide EnumerateNodesFunctor launch and inclusive scans; the three scalar total readbacks become one B-thread gather of per-grid counts and bases at each grid's tile boundaries and one D2H of the Data array. getBuffer lays the grids back to back in one allocation. d_upperOffsets is now set for every grid, including empty ones (it was left uninitialized when nodeCount[2] == 0).

Build functors locate their grid and rebase batch-wide node offsets into that grid's node arrays: BuildGridTreeRootFunctor / InitGridTreeRootFunctor run over B threads and write the grid's own index and count; BuildUpperNodesFunctor and ProcessLowerNodesFunctor resolve the grid through tileToGrid; the leaf-offset, bbox and post-process functors resolve it by binary search over the per-grid bases (topology::detail::gridOfIndex). Leaf mOffset values and tree.mVoxelCount are rebased per grid. New processGridTreeRoot(stream) launches the header functor over the batch.

ProcessLowerNodesFunctor (util/cuda/Morphology.cuh) gets a batched overload taking the Data array and tileToGrid; the single-grid overload stays and both share one body.

Speculative-root helpers in topology::detail, for batch callers that must stage one root per grid: tileSortKey (the PointsToGrid offset-shifted key), ProcessedTileMap, insertProcessedTile, insertProcessedTiles (every 4096-aligned tile overlapping a bbox, with floor semantics for negative coordinates) and packProcessedRoot. The existing consumers keep their own root enumeration, so the batch test checks these helpers against that code rather than against themselves.

The mask-fill and leaf-fill functors (*InternalNodesFunctor, *LeafMasksFunctor) need no changes: a batch driver offsets their processed-root and mask-row pointers per grid, as the new test does.

Memory-resource conventions (#2232, #2322). All new scratch is typed cuda::Buffer<T> sized in elements with the mandatory noInit, memsets use size_bytes(), device buffers borrow the injected resource through ResourceRef, and only the variable-length root image stays a byte buffer. The builder's host-side staging is pinned throughout: the per-grid Data array (Buffer<Data, PinnedResource>, replacing the pageable Data member, so the countNodes readback is now genuinely asynchronous), the processed roots, and the tile-to-grid table. No pageable memory is handed to cudaMemcpyAsync anywhere in the builder.

How the batched build works

The example is the new unit test: three grids refined in one pass, with g1 empty. The top panel is master, where batching means B independent builds (each with its root readback, count sync, final sync and the two syncs in the GridHandle constructor) followed by mergeGridHandles and one more sync per grid. The bottom panel is this PR: scratch is indexed by processed tile across the batch, countNodes gathers each grid's node counts and bases at its tile boundaries with one readback, getBuffer lays the grids back to back, and every build kernel resolves its grid (tileToGrid[t] for tile-indexed kernels, a binary search over the per-grid bases for node-indexed ones) and subtracts that grid's bases. The consumer's mask-fill and leaf-fill functors are unchanged; a batch driver only offsets their processed-root and mask-row pointers per grid.

Compatibility

Not a breaking change for existing consumers, whose files are unchanged in this PR. Every method, accessor and Data field they use keeps its signature and meaning for one grid; their own lambdaKernel<<<1,1>>>(1, BuildGridTreeRootFunctor, deviceData()) launches keep working (index 0 of a one-element array). The only behavioral change is inside countNodes, where one D2H of the Data array replaces three scalar copies under the same sync-before-getBuffer contract.

Testing

• New TestNanoVDBCUDA.TopologyBuilderBatch_ValueOnIndex: three members (the refine/coarsen fixture, an empty grid, voxels across the +-4096 tile boundaries and in negative octants) refined in one builder pass on single-space GridHandle<cuda::Buffer<std::byte>> handles; the device handle constructor validates the three-grid chain, and the result is brought to a pinned handle with cuda::copyTo and byte-compared against RefineGrid on each member alone after setting the grid index and count. The test uses no dual-buffer accessors.
• Full nanovdb_test_cuda run: 65 of 66 pass; the one failure, UnifiedBuffer_IO, cannot open data/3_spheres.nvdb from the build directory and fails identically without this change. nanovdb_test_cuda_memory_resource: 19 of 19 pass.
• compute-sanitizer --tool memcheck (2025.3 and 2025.4) reports cudaErrorIllegalInstruction inside the ProcessLowerNodesFunctor launch on both this branch and untouched master with the same toolchain, so it is a pre-existing sanitizer/toolchain interaction on this setup, not something this change introduces. The tests themselves pass outside the sanitizer.

Built with NANOVDB_CUDA_WERROR=ON, CUDA 13.2, sm_120 (RTX PRO 6000 Blackwell). Rebased on current master; pendingchanges/nanovdb.txt has an entry under Improvements.

🤖 Generated with Claude Code