NanoVDB: CUDA connected-components labeling for ValueOnIndex grids by sifakis · Pull Request #2261 · AcademySoftwareFoundation/openvdb

sifakis · 2026-07-23T21:04:08Z

Summary

Adds nanovdb::tools::cuda::ConnectedComponents — a CUDA connected-components labeling of the
active voxels of a NanoVDB ValueOnIndex grid (6-connectivity). Two active voxels share a label
iff they are connected through a path of adjacent active voxels. The public
getVoxelLabelsAndCount() returns a per-active-voxel dense component id in [0, N) plus the
component count N.

Algorithm

Hierarchical, exploiting the 8³ leaf structure:

Per-leaf Shiloach–Vishkin union-find in shared memory (one block per leaf), producing each
leaf-local component's voxel Mask<3> and six face bitmasks.
Cross-leaf edge detection by intersecting the touching face masks over each leaf's +X/+Y/+Z
neighbours (each undirected boundary visited once).
Global lock-free union-find over the resulting component graph; the representative is each
class's minimum slot, so the result is order-independent.

Also included

ex_connected_components_cuda — rasterizes an .obj to a narrow band (MeshToGrid), drops the
√3/2-voxel surface shell (PruneGrid), then labels; a CPU union-find oracle independently
verifies the GPU component count and per-voxel partition.
ConnectedComponentsMultiSphere unit test (analytic ground truth: N separated spheres → 2N
shells; overlapping pair → 2).
A pendingchanges/ entry.

Notes

Depends only on existing NanoVDB CUDA tools (MeshToGrid, PruneGrid, voxelsToGrid) + CUB;
no new third-party dependencies, no OpenVDB requirement for the example.
Branch is based on current master.

Opening as a draft to run CI and for review.

🤖 Generated with Claude Code

linux-foundation-easycla · 2026-07-24T01:33:30Z

The committers listed above are authorized under a signed CLA.

✅ login: JaeHyunLee94 / name: JaeHyun Lee (0bcbb46)
✅ login: sifakis / name: Efty Sifakis (a2e1c20)

sifakis · 2026-08-04T21:37:12Z

+// SV root hook: every vertex v whose smallest active neighbor label m is below parent[v]
+// lowers the slot of v's *parent* (its tree root, once flattened) toward m, via atomicMin.
+// Sets *changed (when non-null) iff some root slot was actually lowered.
+__device__ inline void ccHook(int*& cur, int*& nxt, int n, int* changed)


Is changed a shared memory variable? If so update the comment to indicate so.

sifakis · 2026-08-04T22:01:26Z

+    }
+}
+
+template <typename BuildT>


Document this class please (i.e. what is the operator() expected to do)

sifakis · 2026-08-04T22:05:58Z

+        if (tID == 0) sEdges = 0;
+        __syncthreads();
+        ccForEachCrossLeafEdge<BuildT>(d_grid, d_offsets, d_faces, leafID, tID, blockDim.x,
+            [&] __device__ (uint32_t, uint32_t) { atomicAdd_block(&sEdges, 1); });


Make a comment that the lambda here is normally called with the component indices as arguments, but we're only counting the edges, acting on them in any other way

sifakis · 2026-08-04T22:07:42Z

+    const auto&    leaf   = d_grid->tree().template getFirstNode<0>()[leafID];
+    const uint64_t baseL  = d_offsets[leafID];
+    const int      countL = int(d_offsets[leafID + 1] - baseL);
+    if (countL == 0) return;


What do "L" and "N" postfixes mean here? i.e. in countL, faceL, faceN
(might be useful for clarity to spell them out as Leaf/Neighbor)

Add nanovdb::tools::cuda::ConnectedComponents, which labels the active voxels of a NanoVDB ValueOnIndex grid by 6-connectivity: two active voxels share a label iff they are connected through a path of adjacent active voxels. The public getVoxelLabelsAndCount() returns a per-active-voxel dense component id in [0,N) plus the component count N. The implementation is hierarchical: - per-leaf Shiloach-Vishkin union-find in shared memory (one block per leaf), yielding each leaf-local component's voxel Mask<3> and six face bitmasks; - cross-leaf edge detection by intersecting touching face masks over +X/+Y/+Z neighbours; - a global lock-free union-find over the resulting component graph. Also add the ex_connected_components_cuda example (rasterize an .obj to a narrow band via MeshToGrid, drop the sqrt(3)/2-voxel surface shell via PruneGrid, then label; a CPU union-find oracle verifies the result) and a ConnectedComponentsMultiSphere unit test. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Signed-off-by: Efty Sifakis <esifakis@nvidia.com>

…-voxels Accept one or more .obj files (concatenated into a single mesh) and add a --discard-surface-voxels switch: by default label the full narrow band (one component per closed surface), or prune the sqrt(3)/2-voxel barrier shell first when the switch is given. --voxel-size / --band-width replace the old positional args. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Signed-off-by: JaeHyun Lee <jaehlee@nvidia.com>

Use the same #ifndef-define / #ifdef-#undef pattern as TopologyBuilder.cuh and PointsToGrid.cuh so the helper macro neither leaks nor collides. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Signed-off-by: Efty Sifakis <esifakis@nvidia.com>

Comment-only: +/-, x, ->, - in place of the plus-minus, times, arrow, and em-dash glyphs. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Signed-off-by: Efty Sifakis <esifakis@nvidia.com>

Scope neighborMin/hook/compress and the INACTIVE sentinel under LeafUnionFind, dropping the redundant cc prefixes, and add solve() so the warm-up plus convergence loop lives in one place instead of being copied into both leaf functors. Also drop the single-use ParentsT alias and LEAF_DIM constant, and the redundant `current` argument of neighborMin (it was always parents[n]). Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Signed-off-by: Efty Sifakis <esifakis@nvidia.com>

The per-leaf loop runs Liu & Tarjan's algorithm P, which has no proven step bound, so its 64-round cap rested on a heuristic. Restrict the hook to root parents after SwitchToRootAfter rounds, which turns the tail into their algorithm R and its O(lg n) bound. Leaves are observed to converge in far fewer rounds than the switch, so it normally never runs. The warm-up gains a fourth compress: leaving the forest flat rather than merely shallow is what keeps the root-restricted hook free. A leaf that did exhaust the cap was abandoned mid-solve and reported too many components, which nothing downstream could tell from genuine fragmentation. solve() now reports whether it converged, and leaves that did not are counted and warned about on stderr. The counter is read back alongside the component offsets, so it costs no extra synchronization. Also from review: spell out the L/N suffixes in the cross-leaf enumerator as Leaf/Neighbor, document the two cross-leaf edge functors and the unnamed callback parameters, and note that the primitives' `changed` flag lives in shared memory. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Signed-off-by: JaeHyun Lee <jaehlee@nvidia.com>

@PARAM

Name each primitive by its union-find operation and state the Liu & Tarjan correspondence separately: hook is parent-connect, or parent-root-connect when rootsOnly is set, and compress is shortcut. Previously the summary lines said "SV" while the detail below them used the paper's vocabulary. The rootsOnly @PARAM no longer carries the naming and now only explains why the restriction buys monotonicity. Note in compress that parent[v] <- parent[parent[v]] is pointer jumping applied to every node at once, so it is not mistaken for the sequential path halving it resembles. Also drop a forward pointer to the cross-leaf stage from the class comment, drop an orphan "Phase A" label in hook (there is no phase B), and split an overloaded comment in compress so the branch condition and the grandparent read are documented separately. Comment-only; no functional change. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Signed-off-by: Efty Sifakis <esifakis@nvidia.com>

Give ConnectedComponents a ResourceT template parameter, defaulted to cuda::DeviceResource, and move its device buffers from the dual cuda::DeviceBuffer onto the single-space cuda::Buffer, borrowing the injected resource through a ResourceRef so all traffic reaches the caller's instance rather than a copy. None of the buffers is read on the host, so none needs DeviceBuffer's host pointer or per-device array. Typing each by its element removes fifteen casts, and sizes are now counts of elements rather than hand-computed byte products. Two defects go away with the representation rather than by hand: - DeviceBuffer::create(bytes, nullptr, false) bound the (int device) overload, because the (bool host) overload has no default for its stream argument. The false was therefore a device id, pinning every buffer to device 0 and the null stream instead of the current device and mStream. cuda::Buffer takes no device argument at all and its stream is a mandatory constructor parameter. - The rank buffer in processVoxelLabels was freed at scope exit on the null stream while the voxel-label scatter was still reading it on mStream, which is safe only for a blocking stream. cuda::Buffer retains its allocation stream and frees on that, so the free is ordered behind the scatter. getVoxelLabelsAndCount now returns the label buffer by value instead of a raw pointer into a member. Ownership passes to the caller, so the labels no longer require the operator to outlive them -- only the caller's own resource -- and the buffer carries its element count and frees itself on its retained stream. The stream is synchronized before returning, so the contents are complete. Add three tests: the injected resource observes every allocation with no leak, the labels stay valid after the operator is destroyed, and the pipeline runs on a non-blocking stream, a path the suite did not exercise before. Signed-off-by: Efty Sifakis <esifakis@nvidia.com>

The per-voxel labels, the dense count N, and the global slots stored in CrossLeafEdge were each spelled uint32_t independently. Give them one namespace-scope alias, ComponentLabelT, since their capacities provably coincide: when every active voxel is isolated no fragments merge, so N equals the leaf-local component count K, and neither use can be narrowed relative to the other. The two label functors take the type by deduction, so nothing outside the class has to name it. The three cub scans took their item counts as int, which halved the addressable range and pinned the scans to 32-bit offsets for no reason: cub deduces its offset type from the argument's size, so passing leafCount and K in their natural types is a deletion rather than an addition. Measured on a 1M-voxel checkerboard, the now 64-bit rank scan costs 0.025 ms of a 4.35 ms pipeline. That leaves ComponentLabelT as the only thing bounding the addressable range, which MaxLeafComponents now states directly. K is carried as uint64_t precisely so a value that fails the bound survives to be tested, and it is tested once the leaf-local components have been counted -- before any global slot is stored, and before N exists. Past that point every narrower type is known to fit, which is what lets N share the label type. Also add the missing timer around the rank scan, the largest of the three, and an alignment assert for the word-sized scratch. Signed-off-by: Efty Sifakis <esifakis@nvidia.com>

Complete the ComponentLabelT unification. The enumerator that produces the global slots still spelled them uint32_t, so widening the alias would have raised MaxComponentCount past 2^32 while the cast kept truncating, leaving the bound and the cast to disagree silently. Name the alias at the cast and in both consuming callbacks, and record at the cast itself why narrowing from the uint64_t offsets is safe, since the check that licenses it lives in another function. Move LeafNeighborTap into cc_detail. Every use of it in code is already there, and at namespace scope it placed six very generic names -- minusX, plusX, minusY, plusY, minusZ, plusZ -- into the namespace shared by every CUDA tool. Rename MaxLeafComponents to MaxComponentCount: "leaf-local" reads as a per-leaf ceiling, which is 256, rather than the grid-wide aggregate the constant actually bounds. Its comment also leaned on K and N, neither of which is defined at that point, so cut it to one line; the reasoning it carried is already stated at the check. Describe the returned labels as a per-voxel sidecar instead of spelling out its length and leaf.getValue(n) indexing, both of which the term implies. Say that the resource is needed only to free the buffer, not to read it, and that the default one has program lifetime; as written it read as a standing obligation on the caller. Finally, note that no scratch buffer is released before the operator is destroyed, though the face masks and the cross-leaf edges each die a stage early. Comment and naming only; no functional change. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Signed-off-by: Efty Sifakis <esifakis@nvidia.com>

swahtz added the nanovdb label Jul 23, 2026

swahtz assigned sifakis Jul 23, 2026

swahtz assigned sifakis and unassigned sifakis Jul 24, 2026

sifakis commented Aug 4, 2026

View reviewed changes

sifakis force-pushed the connected-components branch from 0bcbb46 to c3ee877 Compare August 5, 2026 20:45

sifakis force-pushed the connected-components branch 2 times, most recently from 4c14491 to d364d90 Compare September 3, 2026 18:08

sifakis and others added 8 commits September 3, 2026 14:27

nanovdb: replace unicode characters with ASCII in ConnectedComponents

cf8effe

Comment-only: +/-, x, ->, - in place of the plus-minus, times, arrow, and em-dash glyphs. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Signed-off-by: Efty Sifakis <esifakis@nvidia.com>

sifakis force-pushed the connected-components branch from d364d90 to cac85b3 Compare September 3, 2026 20:48

sifakis and others added 2 commits September 3, 2026 17:53

Navigation Menu

Uh oh!

NanoVDB: CUDA connected-components labeling for ValueOnIndex grids - #2261

NanoVDB: CUDA connected-components labeling for ValueOnIndex grids#2261
sifakis wants to merge 10 commits into
AcademySoftwareFoundation:masterfrom
sifakis:connected-components

sifakis commented Jul 23, 2026

Uh oh!

linux-foundation-easycla Bot commented Jul 24, 2026 •

edited

Loading

Uh oh!

sifakis Aug 4, 2026

Uh oh!

sifakis Aug 4, 2026

Uh oh!

sifakis Aug 4, 2026

Uh oh!

sifakis Aug 4, 2026 •

edited

Loading

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Uh oh!

Conversation

sifakis commented Jul 23, 2026

Summary

Algorithm

Also included

Notes

Uh oh!

linux-foundation-easycla Bot commented Jul 24, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

sifakis Aug 4, 2026

Choose a reason for hiding this comment

Uh oh!

sifakis Aug 4, 2026

Choose a reason for hiding this comment

Uh oh!

sifakis Aug 4, 2026

Choose a reason for hiding this comment

Uh oh!

sifakis Aug 4, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

linux-foundation-easycla Bot commented Jul 24, 2026 •

edited

Loading

sifakis Aug 4, 2026 •

edited

Loading