Repository navigation
NanoVDB: add tools::cuda::MeshToSDF for index grids - #2417
Draft
JaeHyunLee94 wants to merge 35 commits into
Draft
JaeHyunLee94 wants to merge 35 commits into
JaeHyunLee94 wants to merge 35 commits into
Conversation
tools::cuda::signedFloodFill writes +/-background into the inactive values of a float grid. An index grid has no inactive values to write to: every inactive voxel shares the background index. PropagateSign instead takes the signs of the active voxels from a sidecar and returns, for every tree level, a mask of the inactive regions that are inside: per-leaf and per-node masks and the origins of inside root tiles. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: JaeHyun Lee <jaehlee@nvidia.com>
Records the model, the caller-contract preconditions (no empty leaves; two-sided shielding of the narrow band), the leaf-locality and leaf sweep-bound results, and the proposed unified node-local propagator. Draft-only; to be removed before the PR leaves draft. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
…tees Generalize precondition 1 to "no empty nodes" and record which builders guarantee it (PruneGrid output does; TopologyBuilder alone does not). Add the "unbounded region is outside" precondition, the proof that upward-only seeding labels every inside childless slot, and the root value-tile seeding bug. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
…-only root) Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
…oarser levels in the same kernel The leaf fill (one thread and one shared byte per voxel) and the seeding from leaf faces took most of propagate(). One warp per node now holds the node's mask in registers, one bit per slot (lane x holds x-slice x), and updates every slot in each sweep with word operations: shifts for y and z neighbors, warp shuffles for x neighbors, and a warp vote to stop. The same warp then seeds the coarser levels across the node's faces, so the leaf, lower and upper levels take one kernel each and the seven stages become four. On MeshToSDF output for 10 meshes (resolution 2048, hairball 1024), propagate() is 5-9x faster, with identical masks at every level. The sweeps are Jacobi at every level, so the lower and upper fills no longer use atomics or race. Setup, the root stage and the face-seeding helpers are unchanged. Functions are ordered with helpers before their callers, as in ConnectedComponents.cuh. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: JaeHyun Lee <jaehlee@nvidia.com>
…ed fills Describe the node fills as of 700ebd6 (one warp per node, WarpSlicedMask, seeding fused into the fill kernels) and propose a thread-agnostic Mask2D slice type with a 2D MaskShift<dj, dk> overload and detail::zShiftBits, with the warp-level coordination kept separate. Includes the validated prototype. Update the hierarchy-completeness proof and the code-level findings to the current stage structure and functor names. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
…s the design review Adopt the Mask2D proposal from the design notes. Mask2D, the 2D MaskShift<dj, dk> and detail::zShiftBits go into util/MorphologyHelpers.h, and util/cuda/MorphologyHelpers.cuh gains a 3D MaskShift<di, dj, dk> on Mask2D slices held by one warp per node (lane x holds x-slice x). These replace repeatBits and WarpSlicedMask in PropagateSign.cuh: a sweep is six unit MaskShifts, and the masks and the speed are unchanged. Also address the review findings: - probeChildlessSlot becomes probeTerminalSlot, which returns the node's LEVEL together with the node and slot instead of the literals 0-3. Root value tiles, which the preconditions exclude, now assert, also in RootWallFunctor. - Each fill asserts the sweep bound that holds under the preconditions (22, 45 and 93 for leaf, lower and upper nodes, and the Manhattan size of the array at the root), while still looping until a sweep changes nothing, so release builds converge on any input. - The root fill uses Jacobi sweeps over two arrays, so it no longer reads and writes the same bytes in one sweep, and propagate() releases the root-tile array before returning, as ConnectedComponents does. - Document preconditions 1-4 and the sign convention, size launches with util::cuda::blocksPerGrid, and use one term for a cell of the root array. The masks are identical on 29 MeshToSDF cases in release and debug builds, and the regression test passes 14/14. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: JaeHyun Lee <jaehlee@nvidia.com>
…nd, and what was not done Add an update for 2026-10-18: the code changes, the argument for the 22/45/93 sweep bound (which closes the open point on coarser levels), the parts of the Mask2D proposal that were not done and why (packing several nodes per warp measured slower in the leaf stage, so one node per warp stays), and what is still open. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: JaeHyun Lee <jaehlee@nvidia.com>
…sign Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
The warp MaskShift in util/cuda/MorphologyHelpers.cuh takes Mask2D slices, which util/MorphologyHelpers.h defines, so the CUDA header did not compile as a first include, which nanovdb_cuda_header_check (AcademySoftwareFoundation#2377) now requires. Include it: util/MorphologyHelpers.h includes this header back only at its end, where the include guard skips it, so either header can come first. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
…VDB_ include-guard prefix DeviceGridTraits now reads on a given stream (AcademySoftwareFoundation#2383), so initializeMasks passes mStream to getTreeData and getIndexBBox instead of synchronizing mStream before reading on the default stream. Also rename the include guard to NANOVDB_TOOLS_CUDA_PROPAGATESIGN_CUH_HAS_BEEN_INCLUDED, like the other tools/cuda headers (AcademySoftwareFoundation#2395). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
…y construction Review of util/MorphologyHelpers.h: - Mask2D's template parameter is LOG2DIM, its constants are DIM, SIZE (new, the slots in a slice) and WORD_COUNT, and its words are mWords, as in Mask. The words stay public, so a slice remains a plain aggregate. - The static_assert says that only LOG2DIM 3, 4 and 5 are supported, the sizes of NanoVDB's node masks, and the struct comment no longer mentions threads. - detail::yBits becomes detail::firstWordOfY0: the first word of the Mask2D whose on slots are those with y == 0. zShiftBits<LOG2DIM, dk> is documented the same way, as MaskShift<0, dk, false> of the all-on Mask2D, restricted to its first word. The warp MaskShift and PropagateSign follow the renames; facesWithOnSlots takes FirstY from firstWordOfY0 directly. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
Fold the loose lane-layout comment into the function's doc comment, and trim the include comment for util/MorphologyHelpers.h. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
Add the PartitionGrid design notes, copied unchanged from the mesh-to-sdf branch, under examples/ex_partition_grid_cuda. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: JaeHyun Lee <jaehlee@nvidia.com>
PartitionGrid takes the same input as PruneGrid, a source grid and one Mask<3> per leaf, and returns two grids: [0] holds the active voxels whose retain bit is on and is identical to PruneGrid with that mask, and [1] holds the rest, i.e. PruneGrid with the complemented mask. Callers that need both sides no longer prune twice. It runs PruneGrid's stages with two TopologyBuilders. The source root table is read and sorted once and given to both builders, and the two mask-dependent kernels each read a source leaf once and write both outputs (PartitionInternalNodesFunctor and PartitionLeafMasksFunctor, next to their PruneGrid counterparts in Morphology.cuh). An output that keeps no voxel is a valid empty grid. The handles default to the single-space cuda::Buffer. Add TestNanoVDBCUDA.PartitionGrid_ValueOnIndex, which compares both outputs byte for byte with PruneGrid for a mixed mask and for all-on and all-off masks, and install the header. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: JaeHyun Lee <jaehlee@nvidia.com>
…ry implementation and its payoff Add an update for 2026-10-19: why an empty output crashed TopologyBuilder and that AcademySoftwareFoundation#2336 fixed it upstream (which settles open item 8.1: an empty output is a valid empty grid), the binary PartitionGrid, and its measured cost against two PruneGrid calls. The n-ary entry point and the example are still to do. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: JaeHyun Lee <jaehlee@nvidia.com>
Record the new tools::cuda::PartitionGrid. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: JaeHyun Lee <jaehlee@nvidia.com>
Replace the dense root-tile array with one face-seed word per root-table entry, decide the root stage from those seeds after stage 3, and expand interior root tiles in rounds over sorted root keys. Includes the facts it relies on, a 2D example, the cost argument and the alternatives considered. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
…y one PartitionGridBinary Add tools::cuda::PartitionGrid, which splits the active voxels of a ValueOnIndex grid into one grid per label, given a label per value index (the format ConnectedComponents writes). It counts the voxels of each label with one histogram, then splits the label range at the label that halves the voxel count and recurses depth first. Each split builds a retain mask per leaf, one warp per leaf, and runs PartitionGridBinary on it. A split grid numbers its values anew, so the mask kernel reads each voxel's label through the source leaf at the same origin. A range of labels without voxels is cut at its middle label, which keeps the recursion shallow. The binary class becomes PartitionGridBinary: the n-ary class is the general case, and the binary one is its building block. Both classes follow ConnectedComponents and PruneGrid: a static_assert for ValueOnIndex, the shared message for grids with active tiles, the same constructor, setVerbose and setChecksum docs, build() beside getHandles, and 183-character separators. The n-ary outputs are appended in label order rather than assigned, so an output buffer type that is not default constructible, such as the one build() returns, works. Add TestNanoVDBCUDA.PartitionGrid_ValueOnIndex: every output is byte-identical to PruneGrid with the mask of its label, for uneven labels including one with no voxels, for 13 labels and for a single label. The binary test becomes PartitionGridBinary_ValueOnIndex, and the pending change notes list both classes. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: JaeHyun Lee <jaehlee@nvidia.com>
…s and the performance Add updates for 2026-10-28 and 2026-10-29: the n-ary implementation and how it maps to §7, why the classes are now PartitionGridBinary and PartitionGrid, and the measured cost of both against PruneGrid. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: JaeHyun Lee <jaehlee@nvidia.com>
…ve regions - Take the SDF sidecar itself (const ValueT*, ValueT floating point, default float) instead of an int8_t sign array: a voxel is inside if its SDF value is negative, zero counts as outside. - Require a ValueOnIndex grid. leafSeedSlice reads each value index once, and an inactive voxel's index 0 marks it as never a seed. - Name the outputs invert masks, returned as a PropagateSign::InvertMasks struct instead of a tuple. A bit is on for an inactive element whose value is the background with its sign inverted, so leaf masks no longer flag active voxels; face seeding still uses the whole interior. Document the definition, its invariance under complementing the level set, and restate precondition 3 relative to the background's sign. - Rename the signing namespace to floodfill, and state in the stage comment why the node-by-node, finest-first fill is exact under the preconditions. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
…sely at the root - propagate() takes the caller's leaf, lower and upper mask arrays, which need not be initialized, and returns the origins of the inside absent root tiles and whether there is any, instead of an InvertMasks struct. Node counts come from a host copy of the tree header. - Replace the dense root flood with the sparse root propagation proposed in the design notes (2026-10-29). A seed at the root sets a face bit of the root-table entry that holds the seeding node, and stage 4 reads these bits back and stops if none is set. Otherwise it grows the known positions in rounds of sort, unique, binary search against the sorted known positions and flagged select, merged with cub::DeviceMerge::MergeKeys, within the hull of the entries. This removes the dense root mask, the active-bbox read and the limit on the number of root-tile positions in the active bounds. - Deviations from the proposal: no thrust merge or set_difference, whose scratch would bypass ResourceT; a neighbor outside the hull is not added and trips NANOVDB_ASSERT instead of throwing. - seedAcrossFace descends root -> upper -> lower itself and seeds where the descent ends; TerminalSlot and probeTerminalSlot are removed. - Name a set of faces FaceSet, and call cub through CALL_CUBS with a TempPool, as ConnectedComponents does. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: JaeHyun Lee <jaehlee@nvidia.com>
…SeedSlice - Spell the ValueOnIndex check BuildTraits<BuildT>::is_onindex, as ConnectedComponents and TopologyBuilder do (same condition). - leafSeedSlice: state that it is warp-collective (one warp per leaf, all 32 lanes share the work), that it returns the leaf's active voxels with a negative SDF value, and that this Mask<3> comes back as one Mask2D<3> x-slice per lane; explain how each ballot fills half a slice. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
…arify comments - initializeMasks() throws if the grid has active tiles or if the root table has more entries than there are upper nodes (root value tiles, active or not), so seedAcrossFace only asserts tile->isChild(): without the check, a root value tile would make the walk alias the root as an upper node and write outside the masks in release builds. - seedAcrossFace asserts that the seeding node's own root entry exists. - A set of faces is a plain uint32_t: atomicOr and the warp shuffles fix the width at 32 bits, so the FaceSet alias added nothing. - Comments: BitFloodFunctor is warp-collective and hasSlice says which lanes hold slices; the flood update cites precondition 2; the sweep bound is the node's Manhattan diameter; the same-level returns in seedAcrossFace say that a node of our own level fills itself. - Design notes: an update for 2026-10-30 (interface, the sparse root as implemented, this review pass, still open items, an optional simulator check of seedAcrossFace on synthetic trees). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
The regression harness's one-sign synthetic cases break precondition 3; the sparse root's hull assert catches them in debug builds. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
…nVDB The example reads an OpenVDB level set, strips it to a narrow-band-only ValueOnIndex grid plus SDF sidecar (dropping every value tile and inactive voxel value, and rebuilding the topology from the active voxels), runs PropagateSign, and checks every OpenVDB off value (voxels, tiles, root tiles) against the invert state of the NanoVDB element containing it, plus every returned root tile against OpenVDB in the reverse direction. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
…l the header - fillRootInvertTiles() throws when a neighbor falls outside the hull of the root-table entries, i.e. the input breaks preconditions 2 or 3, as the design notes' guard asks for release builds; it used to drop the neighbor and only assert after the loop. - Stage 4: declare the search state and each round's scratch buffers together, with what each holds, and split the two-int counts buffer, whose first slot held two different counts, into uniqueCount, newCount and escaped. - Add PropagateSign.cuh to NANOVDB_INCLUDE_TOOLS_CUDA_FILES, so that it is installed and compiled on its own by nanovdb_cuda_header_check. - pendingchanges: describe the SDF input and the caller-allocated masks. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: JaeHyun Lee <jaehlee@nvidia.com>
…t tests - ex_propagate_sign_nanovdb_cuda reads an OpenVDB level set, strips it to its narrow band (only the active leaf voxels are kept, so every value tile and every inactive voxel value is dropped), converts the band to a ValueOnIndex grid whose blind data channel is the SDF sidecar, and runs PropagateSign. Every off value of the level set, voxels and tiles at every level, is then checked against the invert bit of the element that contains it, and every returned root tile against the level set, as planned in the design notes. - TestNanoVDB.cu, PropagateSign_Sphere: on the narrow band of a sphere from createLevelSetSphere, every leaf, lower and upper mask bit is compared with the analytic answer. - TestMemoryResource.cu: PropagateSign_InjectedResourceSeam and PropagateSign_NonBlockingStream, as for ConnectedComponents. - Design notes: record the previous commit's review pass, stage 4 on valid inputs, the unit tests (an analytic sphere without OpenVDB, which leaves the comparison with OpenVDB to the example) and the example's results. - pendingchanges: mention the example and the unit tests. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: JaeHyun Lee <jaehlee@nvidia.com>
…and simplify the tests - ex_partition_grid_cuda reads several .vdb files, places their active voxels side by side along x so that inputs never touch, labels the connected components with ConnectedComponents and splits them with PartitionGrid. Each output is compared byte for byte with PruneGrid given the mask of its component, and with the OpenVDB segmentActiveVoxels segment that holds one of its voxels. - The tests are now unit tests on five voxels: each output must hold exactly the voxels of its side or label. They replace the byte comparisons with PruneGrid, and the class documentation and pending change notes no longer promise outputs identical to PruneGrid's. - Design notes: record the 2026-10-30 review, the new tests and the example, and what is left. - pendingchanges: mention the example and the unit tests. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: JaeHyun Lee <jaehlee@nvidia.com>
…h-to-sdf-pr MeshToSDF builds on PropagateSign, which is not in develop/nanovdb yet. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: JaeHyun Lee <jaehlee@nvidia.com>
…h-to-sdf-pr MeshToSDF builds on PartitionGrid, which is not in develop/nanovdb yet. The conflicts in examples/CMakeLists.txt and TestNanoVDB.cu keep both sides. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: JaeHyun Lee <jaehlee@nvidia.com>
…nate - isInvert(grid, ijk, leafMasks, lowerMasks, upperMasks, rootOrigins, rootOriginCount), a host and device free function at the end of the header: it descends from the root to the first element without a child that holds ijk and returns its invert bit. For an absent root tile, that is whether its origin is among the returned root-tile origins, found by bisection since they come sorted. An active voxel's bit is off, so it returns false there; its sign is that of its own SDF value. - The class documentation points to it, and pendingchanges mentions it. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: JaeHyun Lee <jaehlee@nvidia.com>
…h-to-sdf-pr Brings in isInvert, which MeshToSDF uses to test whether a surface lies inside another. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: JaeHyun Lee <jaehlee@nvidia.com>
…grids
- tools::cuda::MeshToSDF rasterizes the unsigned distance of the mesh with
MeshToGrid, splits the band into its closed surfaces with
ConnectedComponents and PartitionGrid, and signs each surface: the voxels
outside the barrier around the surface by the connected components of the
barrier-free band, and the barrier voxels by the OpenVDB intersecting-voxel
heuristic or all as interior. Nested surfaces are composed under an
even-odd or solid rule, testing a voxel of each other surface with
PropagateSign's isInvert. setIsoValue() signs the offset surface
{ udf == isoValue }. getHandleAndSDF() returns the grid and the signed
distances of its active voxels; PropagateSign extends the sign beyond them.
- MeshToGrid: add getHandleAndUDFAndIndex, which also returns the nearest
triangle of each active voxel from one packed atomic minimum, and share the
topology pipeline with getHandleAndUDF in a private buildPrunedGrid.
- Rasterization.cuh: add ComputeUDFAndIndexFunctor for it.
- Install MeshToSDF.cuh with the other CUDA tool headers, and add a
pendingchanges entry.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Signed-off-by: JaeHyun Lee <jaehlee@nvidia.com>
This branch has not been deployed
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.
Draft: work in progress.
Summary
Adds
tools::cuda::MeshToSDF(nanovdb/tools/cuda/MeshToSDF.cuh), which converts closed triangle meshes to anarrow-band signed distance field on a
ValueOnIndexgrid on the GPU.getHandleAndSDF()returns the grid anda sidecar with the signed distance of each active voxel (slot 0 is the positive background);
tools::cuda::PropagateSignextends the sign to the inactive voxels and tiles.Pipeline
MeshToGrid::getHandleAndUDFAndIndexbuilds the narrow band, with the unsigned distance and thenearest triangle of each active voxel.
ConnectedComponentslabels the closed surfaces of the band, andPartitionGridgives each itsown grid.
PruneGrid);ConnectedComponents: the component of the minimum-x voxel is exterior,the others interior;
PropagateSign) andisInverttell which surfacesenclose which, and an even-odd or solid rule gives the final sign.
setIsoValue(), the offset surface{ udf == isoValue }is signed instead, and the distances are takento it.
Other changes
MeshToGrid: addgetHandleAndUDFAndIndex, which also returns the nearest triangle of each active voxel fromone packed atomic minimum; its topology pipeline is shared with
getHandleAndUDFin a privatebuildPrunedGrid.Rasterization.cuh: addComputeUDFAndIndexFunctorfor it.Dependencies
Builds on #2361 (
PropagateSign) and #2375 (PartitionGrid), which are not indevelop/nanovdbyet. This branchmerges them, so their files appear in the diff until they land.
TODO
unittest/TestNanoVDB.cu).nanovdb/examples.