Repository navigation
Conversation
…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>
swahtz
reviewed
Oct 21, 2026
swahtz
left a comment
Contributor
There was a problem hiding this comment.
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&.
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
approved these changes
Oct 21, 2026
swahtz
left a comment
Contributor
There was a problem hiding this comment.
Looks good, thanks for addressing my notes.
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>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Adds resource forms of the point-set overloads of
tools::cuda::pointsToGridandvoxelsToGrid, 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 throughresourceonstream, and the result is aGridHandle<cuda::Buffer<std::byte, ResourceT>>, defaulting toDeviceResource:The point-set prototype forms are constrained the same way as the single-set ones, so a
cudaStream_targument selects the resource form. So does a stream wrapper that converts tocudaStream_t, such asat::cuda::CUDAStream; before, the wrapper was deduced as the prototype buffer and failed deep inside the build, in the single-set forms too. A literal0in that position now fails with a message saying to writecudaStream_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:PointsToGridbuilder andDistributedPointsToGrid. 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.updateChecksumwithCheckMode::Fullhandles an empty grid the same way the host does.mergeGridHandlesof no grids andcopyToof 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.tools::checkGridno 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 nowPointType::Disablefor a point grid, which used to throw only after the grid was allocated.This also fixes
pointsToGridwithPointType::Defaultfor pointer-to-const coordinates (e.g.const Vec3f*). It failed to compile in every form, because it assigned into aconstelement type; the other element-type uses in the file already strip theconst.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:
TestBuffer.ResourceFormsOfThePointSetOverloadscovers 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 andPointType::Disableerrors. Its point sets useconstpointers, so it also covers thePointType::Defaultfix.ZeroPointsBuildAnEmptyGrid,EmptyHandlesMergeAndCopyForEveryBufferTypeandStreamWrapperSelectsTheResourceForm(all inTestBuffer), plusZeroPoints_DistributedCudaPointsToGridrun single-GPU. The empty grids passisValid(CheckMode::Full), and a-Gdebug build of the device and host paths also validates them.-Werror, and the 38 Python tests pass.🤖 Generated with Claude Code