Add unit tests for MeshToGrid: empty mesh and unit tetrahedron

MeshToGrid_EmptyMesh: verifies that getHandle() returns a valid grid
with zero active voxels when the triangle list is empty.

MeshToGrid_UnitTetrahedron: constructs a unit tetrahedron (4 vertices,
4 faces) with voxel size 0.1. Tests both getHandle() (topology-only)
and getHandleAndUDF() (with sidecar). Checks:
  - Topology: checksum from getHandle() equals that from getHandleAndUDF(),
    confirming both paths produce identical grids.
  - Per-voxel UDF: sidecar values match CPU brute-force (index-space
    pointToTriangleDistSqr * voxelSize) within 1e-3 voxels.
  - No false positives: every active voxel lies within the narrow band.
  - No false negatives: every voxel with CPU UDF strictly inside the band
    is active (checked via leaf-iteration active set).

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>

Author: sifakis

nanovdb/nanovdb/unittest/TestNanoVDB.cu

@@ -23,6 +23,8 @@
 #include <nanovdb/tools/cuda/CoarsenGrid.cuh>
 #include <nanovdb/tools/cuda/RefineGrid.cuh>
 #include <nanovdb/util/cuda/Injection.cuh>
+#include <nanovdb/tools/cuda/MeshToGrid.cuh>
+#include <nanovdb/math/Proximity.h>
 #include <nanovdb/util/cuda/Timer.h>
 #include <nanovdb/util/Timer.h>
 #include <nanovdb/io/IO.h>
@@ -34,6 +36,7 @@
 #include <cuda_runtime_api.h>
 #include <gtest/gtest.h>
 #include <algorithm>// for std::sort
+#include <unordered_set>
 #include <iomanip> // for std::setw, std::setfill
 #include <thread> // for std::thread
 
@@ -3746,3 +3749,178 @@ TEST(TestNanoVDBCUDA, GridHandle_from_HostBuffer)
     }
 }
 
+TEST(TestNanoVDBCUDA, MeshToGrid_EmptyMesh)
+{
+    using BuildT = nanovdb::ValueOnIndex;
+
+    nanovdb::Map map;
+    map.set(0.1, nanovdb::Vec3d(0.0));
+
+    nanovdb::tools::cuda::MeshToGrid<BuildT> converter(nullptr, 0u, nullptr, 0u, map);
+    converter.setVerbose(0);
+    auto handle = converter.getHandle();
+    handle.deviceDownload();
+    const auto* grid = handle.grid<BuildT>();
+    ASSERT_NE(grid, nullptr);
+    EXPECT_EQ(grid->tree().activeVoxelCount(), 0);
+}// MeshToGrid_EmptyMesh
+
+TEST(TestNanoVDBCUDA, MeshToGrid_UnitTetrahedron)
+{
+    using BuildT = nanovdb::ValueOnIndex;
+
+    // Unit tetrahedron: four vertices, four triangular faces.
+    // Vertex coordinates are in world space.
+    const std::vector<nanovdb::Vec3f> hostPoints = {
+        {0.f, 0.f, 0.f},  // p0
+        {1.f, 0.f, 0.f},  // p1
+        {0.f, 1.f, 0.f},  // p2
+        {0.f, 0.f, 1.f},  // p3
+    };
+    // t0: z=0 face, t1: x=0 face, t2: y=0 face, t3: diagonal face (x+y+z=1)
+    const std::vector<nanovdb::Vec3i> hostTriangles = {
+        {0, 1, 2},
+        {0, 2, 3},
+        {0, 3, 1},
+        {1, 2, 3},
+    };
+
+    auto nPoints = hostPoints.size();
+    auto nTriangles = hostTriangles.size();
+
+    // Upload points and triangles to device
+    auto pointsBuf = nanovdb::cuda::DeviceBuffer::create(nPoints * sizeof(nanovdb::Vec3f), nullptr, false);
+    ASSERT_TRUE(pointsBuf.deviceData());
+    cudaCheck(cudaMemcpy(pointsBuf.deviceData(), hostPoints.data(),
+                         nPoints * sizeof(nanovdb::Vec3f), cudaMemcpyHostToDevice));
+
+    auto trisBuf = nanovdb::cuda::DeviceBuffer::create(nTriangles * sizeof(nanovdb::Vec3i), nullptr, false);
+    ASSERT_TRUE(trisBuf.deviceData());
+    cudaCheck(cudaMemcpy(trisBuf.deviceData(), hostTriangles.data(),
+                         nTriangles * sizeof(nanovdb::Vec3i), cudaMemcpyHostToDevice));
+
+    auto dPoints = static_cast<const nanovdb::Vec3f*>(pointsBuf.deviceData());
+    auto dTriangles = static_cast<const nanovdb::Vec3i*>(trisBuf.deviceData());
+
+    // Uniform-scale map: world = dx * index
+    const double dx = 0.1;
+    nanovdb::Map map;
+    map.set(dx, nanovdb::Vec3d(0.0));
+
+    const float bandWidth = 3.0f;  // default, in voxels
+    const float bandWidthWorld = bandWidth * float(dx);
+
+    // CPU brute-force UDF in index space (matches GPU arithmetic exactly):
+    // transform world-space verts to index space, compute pointToTriangleDistSqr
+    // with integer voxel centers, scale result to world space.
+    std::array<nanovdb::Vec3f, 4> idxVerts;
+    for (uint32_t i = 0; i < nPoints; ++i)
+        idxVerts[i] = map.applyInverseMap(hostPoints[i]);
+
+    auto cpuUDF = [&](int ix, int iy, int iz) -> float {
+        const nanovdb::Vec3f p{float(ix), float(iy), float(iz)};
+        float minDistSqr = std::numeric_limits<float>::max();
+        for (const auto& tri : hostTriangles) {
+            const float d = nanovdb::math::pointToTriangleDistSqr<nanovdb::Vec3f>(
+                idxVerts[tri[0]], idxVerts[tri[1]], idxVerts[tri[2]], p);
+            minDistSqr = std::min(minDistSqr, d);
+        }
+        return std::sqrt(minDistSqr) * float(dx);  // world-space distance
+    };
+
+    // --- Topology-only path ---
+    uint64_t topoChecksum = 0;
+    {
+        nanovdb::tools::cuda::MeshToGrid<BuildT> conv(dPoints, nPoints, dTriangles, nTriangles, map);
+        conv.setVerbose(0);
+        conv.setChecksum(nanovdb::CheckMode::Full);
+        auto handle = conv.getHandle();
+        handle.deviceDownload();
+        const auto* topoGrid = handle.grid<BuildT>();
+        ASSERT_NE(topoGrid, nullptr);
+        topoChecksum = topoGrid->mChecksum.full();
+        EXPECT_GT(topoGrid->tree().activeVoxelCount(), uint64_t(0));
+    }
+
+    // --- UDF path ---
+    nanovdb::tools::cuda::MeshToGrid<BuildT> converter(dPoints, nPoints, dTriangles, nTriangles, map);
+    converter.setVerbose(0);
+    converter.setChecksum(nanovdb::CheckMode::Full);
+    auto [handle, sidecarBuf] = converter.getHandleAndUDF();
+
+    handle.deviceDownload();
+    const auto* grid = handle.grid<BuildT>();
+    ASSERT_NE(grid, nullptr);
+
+    const uint64_t sidecarCount = sidecarBuf.size() / sizeof(float);
+    std::vector<float> hostSidecar(sidecarCount);
+    cudaCheck(cudaMemcpy(hostSidecar.data(), sidecarBuf.deviceData(),
+                         sidecarBuf.size(), cudaMemcpyDeviceToHost));
+
+    // Per-voxel UDF correctness: every active voxel must lie inside the narrow band
+    // and its sidecar distance must match the CPU brute-force value within 1e-3 voxels.
+    const uint32_t nLeaves = grid->tree().nodeCount(0);
+    const auto* leaves  = grid->tree().getFirstLeaf();
+    uint64_t activeCount = 0;
+
+    for (uint32_t li = 0; li < nLeaves; ++li) {
+        const auto& leaf = leaves[li];
+        const auto  org  = leaf.origin();
+        for (int vi = 0; vi < 512; ++vi) {
+            if (!leaf.isActive(vi)) continue;
+            ++activeCount;
+
+            const int lx = vi & 7, ly = (vi >> 3) & 7, lz = (vi >> 6) & 7;
+            const int ix = org[0]+lx, iy = org[1]+ly, iz = org[2]+lz;
+
+            const float exactUDF = cpuUDF(ix, iy, iz);
+            ASSERT_LE(exactUDF, bandWidthWorld * (1.f + 1e-5f))
+                << "Active voxel at (" << ix << "," << iy << "," << iz
+                << ") is outside the narrow band (distance=" << exactUDF/float(dx) << " voxels)";
+
+            const uint64_t sIdx = leaf.getValue(vi);
+            ASSERT_LT(sIdx, sidecarCount) << "Sidecar index out of range";
+
+            const float ourUDF    = hostSidecar[sIdx];
+            const float errVoxels = std::abs(ourUDF - exactUDF) / float(dx);
+            EXPECT_LT(errVoxels, 1e-3f)
+                << "UDF error at (" << ix << "," << iy << "," << iz
+                << "): ours=" << ourUDF/float(dx) << " exact=" << exactUDF/float(dx) << " voxels";
+        }
+    }
+    EXPECT_GT(activeCount, uint64_t(0));
+
+    // Build a flat set of active coord keys from the leaf iteration
+    // Encode each (ix,iy,iz) as a uint64_t with 21 bits per axis, offset by 2^20.
+    auto encodeCoord = [](int x, int y, int z) -> uint64_t {
+        return (uint64_t(x + (1<<20)))
+             | (uint64_t(y + (1<<20)) << 21)
+             | (uint64_t(z + (1<<20)) << 42);
+    };
+    std::unordered_set<uint64_t> activeSet;
+    activeSet.reserve(activeCount);
+    for (uint32_t li = 0; li < nLeaves; ++li) {
+        const auto& leaf = leaves[li];
+        const auto  org  = leaf.origin();
+        for (int vi = 0; vi < 512; ++vi) {
+            if (!leaf.isActive(vi)) continue;
+            const int lx = vi & 7, ly = (vi >> 3) & 7, lz = (vi >> 6) & 7;
+            activeSet.insert(encodeCoord(org[0]+lx, org[1]+ly, org[2]+lz));
+        }
+    }
+
+    // No false negatives: every voxel with CPU UDF strictly inside the band must be active.
+    const int ilo = (int)std::floor(-bandWidth) - 1;
+    const int ihi = (int)std::ceil(1.0 / dx + bandWidth) + 1;
+    uint64_t missedCount = 0;
+    for (int ix = ilo; ix <= ihi; ++ix)
+        for (int iy = ilo; iy <= ihi; ++iy)
+            for (int iz = ilo; iz <= ihi; ++iz)
+                if (cpuUDF(ix, iy, iz) < bandWidthWorld)
+                    if (activeSet.count(encodeCoord(ix, iy, iz)) == 0)
+                        ++missedCount;
+    EXPECT_EQ(missedCount, uint64_t(0));
+
+    // getHandle() and getHandleAndUDF() must produce identical grids.
+    EXPECT_EQ(grid->mChecksum.full(), topoChecksum);
+}// MeshToGrid_UnitTetrahedron