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>

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 and others added 2 commits August 5, 2026 15:40

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

sifakis and others added 5 commits August 5, 2026 16:12

nanovdb: replace unicode characters with ASCII in ConnectedComponents

36efbbc

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>

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 7 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