ex_connected_components_cuda: merge multiple OBJs + --discard-surface-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>

Author: JaeHyunLee94

nanovdb/nanovdb/examples/ex_connected_components_cuda/connected_components_cuda.cpp

@@ -5,9 +5,10 @@
 ///
 /// @brief Host driver for the connected-components example (NanoVDB / CUDA only, no OpenVDB).
 ///
-///        Reads a triangle mesh from a Wavefront .obj file, builds the index<->world transform,
-///        and hands the mesh to the CUDA side, which rasterizes it into a ValueOnIndex narrow-band
-///        grid, discards the surface/barrier shell, and runs connected-components labeling
+///        Reads one or more Wavefront .obj meshes (concatenated into a single vertex/triangle list),
+///        builds the index<->world transform, and hands the mesh to the CUDA side, which rasterizes
+///        it into a ValueOnIndex narrow-band grid, optionally discards the surface/barrier shell
+///        (--discard-surface-voxels), and runs connected-components labeling
 ///        (nanovdb::tools::cuda::ConnectedComponents) on the result. The device side also runs a
 ///        CPU union-find oracle that independently verifies the GPU labeling.
 ///
@@ -27,14 +28,17 @@
 
 // ---- Host/device seam (implemented in connected_components_cuda_kernels.cu) -----------------------
 //
-// Rasterize the mesh -> ValueOnIndex narrow band, discard the barrier shell (unsigned distance
-// within sqrt(3)/2 voxels of the surface), and run connected-components labeling. Prints topology
-// diagnostics, the component count, and a CPU-oracle PASS/FAIL, and returns the number of
-// connected components.
+// Rasterize the mesh -> ValueOnIndex narrow band and run connected-components labeling. When
+// @a discardSurfaceVoxels is true, the surface/barrier shell (unsigned distance within sqrt(3)/2
+// voxels of the surface) is pruned first, splitting each closed surface's band into disjoint
+// inner/outer shells; otherwise labeling runs on the full narrow band (one component per closed
+// surface). Prints topology diagnostics, the component count, and a CPU-oracle PASS/FAIL, and
+// returns the number of connected components.
 uint64_t connectedComponentsFromMesh(const std::vector<nanovdb::Vec3f>& points,
                                      const std::vector<nanovdb::Vec3i>& triangles,
                                      const nanovdb::Map&                map,
-                                     float                              bandWidth);
+                                     float                              bandWidth,
+                                     bool                               discardSurfaceVoxels);
 
 /// @brief Minimal Wavefront .obj reader (vertices + faces) using NanoVDB types.
 ///
@@ -86,29 +90,59 @@ static void readOBJ(const std::string&            filename,
 int main(int argc, char* argv[])
 {
     try {
-        if (argc < 2)
-            throw std::runtime_error("usage: " + std::string(argv[0]) +
-                                     " <input.obj> [voxelSize] [bandWidth]");
+        // Parse args: any non-option token is an input .obj; options set the transform and the
+        // barrier-discard switch.
+        std::vector<std::string> objFiles;
+        float voxelSize            = 0.01f;
+        float bandWidth            = 3.0f;
+        bool  discardSurfaceVoxels = false;
 
-        const std::string inputFile = argv[1];
-        const float voxelSize = (argc > 2) ? std::stof(argv[2]) : 0.01f;
-        const float bandWidth = (argc > 3) ? std::stof(argv[3]) : 3.0f;
+        auto nextValue = [&](int& i, const char* opt) -> const char* {
+            if (i + 1 >= argc) throw std::runtime_error(std::string("missing value for ") + opt);
+            return argv[++i];
+        };
+        for (int i = 1; i < argc; ++i) {
+            const std::string a = argv[i];
+            if      (a == "--discard-surface-voxels") discardSurfaceVoxels = true;
+            else if (a == "--voxel-size")  voxelSize = std::stof(nextValue(i, "--voxel-size"));
+            else if (a == "--band-width")  bandWidth = std::stof(nextValue(i, "--band-width"));
+            else if (a.rfind("--", 0) == 0) throw std::runtime_error("unknown option: " + a);
+            else objFiles.push_back(a);
+        }
+        if (objFiles.empty())
+            throw std::runtime_error(
+                "usage: " + std::string(argv[0]) +
+                " <input.obj> [more.obj ...] [--voxel-size S] [--band-width W] [--discard-surface-voxels]");
 
+        // Read and merge all input meshes into one vertex/triangle list. Merging simply concatenates
+        // the meshes in their own coordinates (no repositioning), offsetting each mesh's triangle
+        // indices past the vertices already added.
         std::vector<nanovdb::Vec3f> points;
         std::vector<nanovdb::Vec3i> triangles;
-        std::cout << "Reading " << inputFile << "...\n";
-        readOBJ(inputFile, points, triangles);
-        std::cout << "Loaded " << points.size() << " vertices, "
-                  << triangles.size() << " triangles.\n";
+        for (const std::string& file : objFiles) {
+            std::vector<nanovdb::Vec3f> filePoints;
+            std::vector<nanovdb::Vec3i> fileTriangles;
+            std::cout << "Reading " << file << "...\n";
+            readOBJ(file, filePoints, fileTriangles);
+            const int offset = int(points.size());
+            points.insert(points.end(), filePoints.begin(), filePoints.end());
+            for (const nanovdb::Vec3i& t : fileTriangles)
+                triangles.emplace_back(t[0] + offset, t[1] + offset, t[2] + offset);
+        }
+        std::cout << "Loaded " << points.size() << " vertices, " << triangles.size()
+                  << " triangles from " << objFiles.size() << " mesh(es).\n";
         if (points.empty() || triangles.empty())
             throw std::runtime_error("mesh has no triangles");
 
         // Index<->world transform: uniform voxel size, no translation.
         nanovdb::Map map;
         map.set(double(voxelSize), nanovdb::Vec3d(0.0), 1.0);
 
+        std::cout << "Surface voxels: "
+                  << (discardSurfaceVoxels ? "discarded (barrier shell pruned)"
+                                           : "kept (full narrow band)") << "\n";
         const uint64_t numComponents =
-            connectedComponentsFromMesh(points, triangles, map, bandWidth);
+            connectedComponentsFromMesh(points, triangles, map, bandWidth, discardSurfaceVoxels);
         std::cout << "Connected components: " << numComponents << "\n";
 
         return 0;

nanovdb/nanovdb/examples/ex_connected_components_cuda/connected_components_cuda_kernels.cu

@@ -4,7 +4,7 @@
 /// @file  connected_components_cuda_kernels.cu
 ///
 /// @brief CUDA / NanoVDB side of the connected-components example. Rasterizes a triangle mesh into a
-///        ValueOnIndex narrow-band grid (nanovdb::tools::cuda::MeshToGrid), discards the
+///        ValueOnIndex narrow-band grid (nanovdb::tools::cuda::MeshToGrid), optionally discards the
 ///        surface/barrier shell (unsigned distance within sqrt(3)/2 voxels of the surface) with
 ///        nanovdb::tools::cuda::PruneGrid, and runs connected-components labeling with
 ///        nanovdb::tools::cuda::ConnectedComponents. A CPU union-find oracle independently verifies
@@ -172,7 +172,8 @@ bool validateAgainstOracle(const GridHandleT& derivedHandle, uint32_t leafCount,
 uint64_t connectedComponentsFromMesh(const std::vector<nanovdb::Vec3f>& points,
                                      const std::vector<nanovdb::Vec3i>& triangles,
                                      const nanovdb::Map&                map,
-                                     float                              bandWidth)
+                                     float                              bandWidth,
+                                     bool                               discardSurfaceVoxels)
 {
     const cudaStream_t stream = 0;
 
@@ -188,41 +189,46 @@ uint64_t connectedComponentsFromMesh(const std::vector<nanovdb::Vec3f>& points,
     auto [origHandle, udfSidecar] = converter.getHandleAndUDF();
     const auto* d_orig = origHandle.template deviceGrid<BuildT>();
 
-    // World-space voxel size from the map (uniform scale here, but read it generically).
-    const nanovdb::Vec3d w0 = map.applyMap(nanovdb::Vec3d(0.0, 0.0, 0.0));
-    const nanovdb::Vec3d wx = map.applyMap(nanovdb::Vec3d(1.0, 0.0, 0.0));
-    const float voxelSize   = float(wx[0] - w0[0]);
-
-    // ---- Step 2: discard the surface/barrier shell -> derived topology. ----
-    const float    barrierSqWorld = 0.75f * voxelSize * voxelSize;  // (sqrt(3)/2 * voxelSize)^2
-    const uint32_t srcLeafCount   = Traits::getTreeData(d_orig).mNodeCount[0];
-
-    auto  retainMask   = nanovdb::cuda::DeviceBuffer::create(
-        std::size_t(srcLeafCount) * sizeof(nanovdb::Mask<3>), nullptr, false);
-    auto* d_retainMask = static_cast<nanovdb::Mask<3>*>(retainMask.deviceData());
-
-    nanovdb::util::cuda::operatorKernel<UDFBarrierPruneMaskFunctor>
-        <<<srcLeafCount, UDFBarrierPruneMaskFunctor::MaxThreadsPerBlock, 0, stream>>>(
-            d_orig, static_cast<const float*>(udfSidecar.deviceData()), barrierSqWorld, d_retainMask);
-    cudaCheckError();
-
-    nanovdb::tools::cuda::PruneGrid<BuildT> pruner(d_orig, d_retainMask, stream);
-    auto        derivedHandle = pruner.getHandle();
-    const auto* d_derived     = derivedHandle.template deviceGrid<BuildT>();
+    // ---- Step 2 (optional): discard the surface/barrier shell -> derived topology. ----
+    // With the shell removed, each closed surface's band splits into disjoint inner/outer shells;
+    // without it, connected components run on the full narrow band (one component per closed surface).
+    GridHandleT                      derivedHandle;   // stays empty unless we prune
+    const nanovdb::NanoGrid<BuildT>* d_cc = d_orig;    // grid connected components will label
+    if (discardSurfaceVoxels) {
+        // World-space voxel size from the map (uniform scale here, but read it generically).
+        const nanovdb::Vec3d w0 = map.applyMap(nanovdb::Vec3d(0.0, 0.0, 0.0));
+        const nanovdb::Vec3d wx = map.applyMap(nanovdb::Vec3d(1.0, 0.0, 0.0));
+        const float    voxelSize      = float(wx[0] - w0[0]);
+        const float    barrierSqWorld = 0.75f * voxelSize * voxelSize;  // (sqrt(3)/2 * voxelSize)^2
+        const uint32_t srcLeafCount   = Traits::getTreeData(d_orig).mNodeCount[0];
+
+        auto  retainMask   = nanovdb::cuda::DeviceBuffer::create(
+            std::size_t(srcLeafCount) * sizeof(nanovdb::Mask<3>), nullptr, false);
+        auto* d_retainMask = static_cast<nanovdb::Mask<3>*>(retainMask.deviceData());
+
+        nanovdb::util::cuda::operatorKernel<UDFBarrierPruneMaskFunctor>
+            <<<srcLeafCount, UDFBarrierPruneMaskFunctor::MaxThreadsPerBlock, 0, stream>>>(
+                d_orig, static_cast<const float*>(udfSidecar.deviceData()), barrierSqWorld, d_retainMask);
+        cudaCheckError();
+
+        nanovdb::tools::cuda::PruneGrid<BuildT> pruner(d_orig, d_retainMask, stream);
+        derivedHandle = pruner.getHandle();
+        d_cc          = derivedHandle.template deviceGrid<BuildT>();
+    }
 
-    // ---- Step 3: connected-components labeling on the derived grid. ----
-    nanovdb::tools::cuda::ConnectedComponents<BuildT> cc(d_derived, stream);
+    // ---- Step 3: connected-components labeling on the selected grid. ----
+    nanovdb::tools::cuda::ConnectedComponents<BuildT> cc(d_cc, stream);
     auto [d_labels, numComponents] = cc.getVoxelLabelsAndCount();
     cudaCheck(cudaStreamSynchronize(stream));
 
-    // Diagnostics.
-    const uint64_t derivedActive = Traits::getActiveVoxelCount(d_derived);
-    const uint32_t derivedLeaves = Traits::getTreeData(d_derived).mNodeCount[0];
-    std::cout << "Derived (barrier-removed) grid: " << derivedActive << " active voxels, "
-              << derivedLeaves << " leaves.\n";
+    // Diagnostics + CPU-oracle self-check (on whichever grid was labeled).
+    const GridHandleT& ccHandle = discardSurfaceVoxels ? derivedHandle : origHandle;
+    const uint64_t     ccActive = Traits::getActiveVoxelCount(d_cc);
+    const uint32_t     ccLeaves = Traits::getTreeData(d_cc).mNodeCount[0];
+    std::cout << (discardSurfaceVoxels ? "Derived (barrier-removed) grid: " : "Full narrow-band grid: ")
+              << ccActive << " active voxels, " << ccLeaves << " leaves.\n";
 
-    // CPU-oracle self-check.
-    validateAgainstOracle(derivedHandle, derivedLeaves, derivedActive, d_labels, numComponents);
+    validateAgainstOracle(ccHandle, ccLeaves, ccActive, d_labels, numComponents);
 
     return numComponents;
 }