Skip to content

NanoVDB: resource forms of the point-set pointsToGrid/voxelsToGrid (CUDA) - #2379

Merged
swahtz merged 4 commits into
AcademySoftwareFoundation:masterfrom
harrism:nanovdb-multiset-resource-forms
Oct 21, 2026
Merged

swahtz merged 4 commits into
AcademySoftwareFoundation:masterfrom
harrism:nanovdb-multiset-resource-forms

Conversation

@harrism

@harrism harrism commented Oct 21, 2026 •

Copy link
Copy Markdown
Contributor

Adds resource forms of the point-set overloads of tools::cuda::pointsToGrid and voxelsToGrid, completing the set started in #2374. Like the single-set resource forms, they take (sets, stream, resource). Each set's grid, the merged result and all internal scratch allocate through resource on stream, and the result is a GridHandle<cuda::Buffer<std::byte, ResourceT>>, defaulting to DeviceResource:

std::vector<std::tuple<const Coord* const, size_t, double>> sets{{d_a, nA, 1.0}, {d_b, nB, 0.5}};
auto handle = nanovdb::tools::cuda::voxelsToGrid<float, const Coord*>(sets, stream);       // DeviceResource
auto viaRef = nanovdb::tools::cuda::voxelsToGrid<float, const Coord*>(sets, stream, ref);  // borrowed resource

The point-set prototype forms are constrained the same way as the single-set ones, so a cudaStream_t argument selects the resource form. So does a stream wrapper that converts to cudaStream_t, such as at::cuda::CUDAStream; before, the wrapper was deduced as the prototype buffer and failed deep inside the build, in the single-set forms too. A literal 0 in that position now fails with a message saying to write cudaStream_t(0).

Empty inputs give empty outputs, as in MeshToGrid, the topology tools and merge/split, so a pipeline doesn't have to read a device count back before building:

  • Builders: a zero point count builds a valid empty grid in every form, including the PointsToGrid builder and DistributedPointsToGrid. An empty list of sets gives an empty handle, and a set without points contributes an empty grid. Before, a zero count launched a zero-block kernel, which exits the process in debug builds and builds on garbage in release builds.
  • Checksum: updateChecksum with CheckMode::Full handles an empty grid the same way the host does.
  • Merge and copy: mergeGridHandles of no grids and copyTo of an empty handle return an empty handle over the prototype's resource for single-space buffers, instead of throwing for a resource that can't be default-constructed.
  • Validator: tools::checkGrid no longer asserts in debug builds on a grid with no nodes. That assert also fired for host-built empty grids.

Invalid arguments still throw std::invalid_argument: a negative count, a bad voxel size or map, and now PointType::Disable for a point grid, which used to throw only after the grid was allocated.

This also fixes pointsToGrid with PointType::Default for pointer-to-const coordinates (e.g. const Vec3f*). It failed to compile in every form, because it assigned into a const element type; the other element-type uses in the file already strip the const.

This targets master. Whether it goes into 13.2 is still open: it's additive and nothing in 14.0 depends on it.

Validation on one RTX PRO 6000 Blackwell GPU with nvcc 12.6:

  • New test: TestBuffer.ResourceFormsOfThePointSetOverloads covers the voxel and point set forms with a counting resource, the default resource, byte equality with the prototype forms, the per-set point counts, balanced allocations, empty and mixed set lists, and the invalid-voxel-size and PointType::Disable errors. Its point sets use const pointers, so it also covers the PointType::Default fix.
  • Empty inputs: ZeroPointsBuildAnEmptyGrid, EmptyHandlesMergeAndCopyForEveryBufferType and StreamWrapperSelectsTheResourceForm (all in TestBuffer), plus ZeroPoints_DistributedCudaPointsToGrid run single-GPU. The empty grids pass isValid(CheckMode::Full), and a -G debug build of the device and host paths also validates them.
  • Existing suites: all single-GPU ctest targets pass. Unit tests, tools, examples and the Python module build with -Werror, and the 38 Python tests pass.

🤖 Generated with Claude Code

…UDA)

The single-set pointsToGrid and voxelsToGrid have resource forms that
allocate the output and all scratch through a given resource on a given
stream; the point-set overloads, which build one grid per set and merge
them, did not. Add (sets, stream, resource) forms that route every
per-set build, the merge and the scratch through the resource, and
constrain the point-set prototype forms like the single-set ones: a
cudaStream_t argument selects the resource form, and a literal 0 there
fails with a message saying to write cudaStream_t(0). Every point-set
overload now rejects an empty set with std::invalid_argument instead of
returning an empty handle or failing inside the merge.

Also fix pointsToGrid with PointType::Default for pointer-to-const
coordinates, which assigned into a const element type and failed to
compile.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
@harrism
harrism requested a review from kmuseth as a code owner October 21, 2026 01:39

@swahtz swahtz 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.

Thanks, this fills in the last of the point-set overloads nicely, and the const PtrT fix for PointType::Default is a good catch. I have three things I'd like addressed before merge. A set with zero points still gets through the new empty check. A stream wrapper type can select the prototype form. And one of the new throws has no test. I also left two smaller suggestions, on the duplicated empty check and on taking the set vector by const&.

Comment thread nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh Outdated
Comment thread nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh
Comment thread nanovdb/nanovdb/unittest/TestBuffer.cu Outdated
Comment thread nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh Outdated
Comment thread nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh Outdated
harrism and others added 3 commits October 21, 2026 06:47
A zero point count, an empty set list, or an empty handle now yields a
valid empty grid or an empty handle, instead of a zero-block launch that
exited the process or a throw:

- PointsToGrid, DistributedPointsToGrid, pointsToGrid and voxelsToGrid
  build an empty grid for no points; the set overloads merge no sets
  into an empty handle.
- updateChecksum and blockedCRC32 handle empty grids and zero bytes as
  the host does.
- mergeGridHandles and copyTo return an empty handle over the
  prototype's resource for single-space buffer types.
- tools::checkGrid bounds leaf nodes by the end of the grid, so it no
  longer asserts on a grid without nodes.

Also:
- A stream wrapper that converts to cudaStream_t selects the resource
  form instead of being deduced as the prototype buffer.
- The resource set overloads take the set list by const reference.
- PointsToGrid<Point>::setPointType rejects PointType::Disable up
  front, and DistributedPointsToGrid validates its map.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
… code

The empty-source branch built its result with the single-space GridHandle
constructor, which is defined only in cuda/GridHandle.cuh, so host-compiled
translation units calling copyTo with a single-space destination failed to
link. Build the empty handle through HandleFactory, as the non-empty path
does.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
@swahtz swahtz added the vdb13.2 label Oct 21, 2026

@swahtz swahtz 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, thanks for addressing my notes.

@swahtz
swahtz merged commit 8ea6110 into AcademySoftwareFoundation:master Oct 21, 2026
16 checks passed
swahtz added a commit to swahtz/openvdb that referenced this pull request Oct 21, 2026
The fix for cuda::updateChecksum on grids shorter than one checksum block
landed on master with AcademySoftwareFoundation#2379, which notes it.

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
swahtz added a commit that referenced this pull request Oct 22, 2026
* NanoVDB: batch-capable TopologyBuilder (CUDA)

TopologyBuilder can now build 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: the existing consumers
are untouched and produce byte-identical output.

TopologyBuilderData carries the grid's batch position and the first index
of its tiles and nodes in the batch-wide scratch arrays; the builder holds
one per grid. Processed roots are staged per grid in one pinned buffer
(allocateProcessedRoots), countNodes gathers per-grid counts and bases at
each grid's tile boundaries with one readback, getBuffer lays the grids
back to back, and the build functors rebase batch-wide node offsets into
each grid's own node arrays. d_upperOffsets is now set for every grid,
including empty ones. ProcessLowerNodesFunctor gains a batched overload;
the single-grid one stays and both share one body.

topology::detail gains the helpers a batch caller needs to stage one
speculative root per grid (tileSortKey, ProcessedTileMap,
insertProcessedTile(s), packProcessedRoot). The existing consumers keep
their own root code; moving them onto the helpers is a separate change.

Adds TestNanoVDBCUDA.TopologyBuilderBatch_ValueOnIndex: three members
(the refine/coarsen fixture, an empty grid, voxels across the +-4096 tile
boundaries) refined in one pass through the new helpers and byte-compared
against RefineGrid, on its original root enumeration, for each member alone.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

* NanoVDB: align the batched TopologyBuilder with the memory-resource conventions and document it

Follows the rules of record from #2232 and #2322 in the new code: the
tile-to-grid table is staged in a pinned cuda::Buffer<uint32_t> so every
transfer the builder issues is asynchronous, and the batch test builds on
single-space GridHandle<cuda::Buffer<std::byte>> handles and reads results
through cuda::copyTo instead of the deprecated dual accessors. Adds a
class-level usage example with the synchronization contract and
declaration comments on each stage, keeps the private type aliases in one
block, makes updateTileLayout private, has allocateProcessedRoots return
void, drops the this-> prefixes the header does not otherwise use, gives
the test a typed device-upload helper, and records the capability in
pendingchanges/nanovdb.txt.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>

* NanoVDB: make TopologyBuilder::getBuffer wait for the node counts and drop the tile-to-grid table on reset

countNodes reads the per-grid node counts back into pinned host memory,
so the copy is now genuinely asynchronous; before, the pageable
destination blocked the host and a caller who skipped the documented
synchronize still saw valid counts. getBuffer sizes the output from
those counts, so it now waits on an event countNodes records instead of
relying on the caller. The event covers only work queued up to the
readback, and a public waitForNodeCounts() exposes the same wait for
callers that read data(g)->nodeCount directly. The class-level contract
and the batch test follow suit; the consumers' own stream syncs are left
in place as harmless.

The tile-to-grid table was only released by processLowerNodes, so a
builder reset to a single grid (or an empty batch) after an abandoned
multi-grid upload kept the old table and the tile-indexed functors
indexed the device data array out of bounds with it. resetHostData now
drops both the host staging and the device table, and uploadProcessedRoot
drops them whenever the batch does not need one.

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

* NanoVDB: limit TopologyBuilder::countNodes to the node-count fields of data(g)

countNodes uploaded the whole per-grid parameter array so the gather
kernel could read each grid's tile range, then copied the whole array
back to land six counts per grid. A host write to any other field
between countNodes and the wait was raced by the upload or overwritten
by the copy back. The gather now reads the tile ranges from a small
builder-owned pinned array and writes the counts to device scratch, and
one strided copy lands them in their own fields of data(g); the counts
and bases are adjacent in the struct and a static_assert keeps them so.
getBuffer is again the only upload of data(g), as before the batch work.

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

* NanoVDB: resolve TopologyBuilder tiles to grids by tile base and fix the tile layout once

The tile-indexed functors looked their grid up in a per-tile table that
the builder staged in pinned memory, uploaded, special-cased for a single
grid and had to keep in step with every reset, while the node-indexed
functors already found their grid with a binary search over the per-grid
node bases. The search helper now lives in Morphology.cuh as a generic
template and both tile-indexed functors use it over tileBase, so the
table, its staging, its upload, its null branch and its lifetimes are
gone and a single grid is a batch of one. The single-grid overload of
ProcessLowerNodesFunctor had no caller left and is removed.

The tile layout was re-derived from the host roots on five entry points,
with a fallback for roots already released. uploadProcessedRoot now fixes
it once and the later stages read the stored total.

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

* NanoVDB: stage TopologyBuilder's host data in one pinned block

The batch work left the builder with three pinned buffers (the per-grid
parameters, allocated in every constructor, the processed roots and the
tile bases), each a synchronous cudaMallocHost/cudaFreeHost pair, where
the single-grid builder had one for the roots. One staging block now
holds the roots, the per-grid parameters and the tile bases at fixed
offsets; it is sized when the roots are allocated, grown only when too
small, and the constructor allocates nothing pinned. processLowerNodes
releases the roots by zeroing their byte count rather than freeing, so
the root accessors return null afterwards as before. A fresh builder
reports zero grids until a root is allocated.

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

* NanoVDB: address review of the batched TopologyBuilder

- Hold the builder's event in a movable RAII wrapper and default the
  builder's moves, so it and its consumers stay movable.
- Record that event after every copy into or out of the pinned staging
  block, and wait on it before allocateProcessedRoot(s) reuses the block.
- Restore ProcessLowerNodesFunctor's single-grid operator() as a
  forwarder.
- Assert in data(g) that grid g has been allocated.
- Simplify tileSortKey; document that checksums run per grid; reword
  the MeshToGrid sync comment.
- Extend the batch test: empty grids at both ends, a grid with tiles but
  no nodes, two builds on one builder on a non-blocking stream, and
  compile-time movability checks.

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

* NanoVDB: fix cuda::updateChecksum on grids shorter than one checksum block

util::cuda::blockedCRC32 launched its per-block kernel with zero blocks
when the data held no full 4 KB block, which failed the launch (and the
blocksPerGrid assert in debug builds). The tail of an empty grid is that
short, so tools::cuda::updateChecksum with CheckMode::Full failed on it,
and so did every CUDA topology tool with setChecksum(CheckMode::Full) and
an empty output. Such data now gets the checksum util::blockedCrc32
returns on the host.

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

* NanoVDB: checksum a TopologyBuilder batch without per-grid host round trips

postProcessGridTree called tools::cuda::updateChecksum once per grid,
which builds its own lookup table and reads the grid header back into
pageable memory to learn the grid size. The builder already knows each
grid's size, so it now creates one lookup table for the batch and calls
crc32Head and blockedCRC32 directly, producing the same checksums with
no host round trip. The batch test's second build computes full
checksums and validates them on the host.

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

* NanoVDB: drop the checksum note that master now carries

The fix for cuda::updateChecksum on grids shorter than one checksum block
landed on master with #2379, which notes it.

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

---------

Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
Co-authored-by: Claude Fable 5.1 <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.

2 participants