NanoVDB: encode points for any resource in PointsToGrid (CUDA) (#2244)

* NanoVDB: encode points for any resource in PointsToGrid (CUDA)

The point-encode step in tools::cuda::PointsToGrid was only reachable for
the default DeviceResource, because it lived in the BuildT=Point full
specialization of processPoints (i.e. <Point, DeviceResource>). A Point
grid built with a custom resource fell through to the generic no-op and
was silently left with unencoded point data.

Fold the encode into the generic processPoints and gate it with
if constexpr(is_same<BuildT, Point>). The encode kernels run on device
data and are independent of the resource; only the trailing d_indx
deallocation routes through ResourceT. Non-Point builds and default-
resource Point builds are unchanged.

Add a regression test that builds a NanoGrid<Point> through a custom
resource and asserts every input point is recoverable from the encoded
per-voxel data.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>

* NanoVDB: move the point-encode kernels out of the if-constexpr block

nvcc rejects an extended lambda defined directly inside the block of an
if constexpr statement on some host compilers, so the nine encode
launches failed to compile against MSVC while building fine with gcc and
clang. Give them their own member function and have the branch call it.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>

---------

Signed-off-by: Mark Harris <mharris@nvidia.com>
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>

Author: harrism

nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh

@@ -373,6 +373,13 @@ public:
     template<typename PtrT>
     void processPoints(const PtrT points, size_t pointCount);
 
+    // Only instantiated when BuildT == Point, from the corresponding branch of
+    // processPoints. Holds the encode kernels because nvcc rejects an extended
+    // lambda defined directly inside the block of an if constexpr statement on
+    // some host compilers.
+    template<typename PtrT>
+    void encodePoints(const PtrT points, size_t pointCount);
+
     void processBBox();
 
     // the following methods are only defined when BuildT == Point
@@ -1164,23 +1171,24 @@ inline void PointsToGrid<BuildT, ResourceT>::processLeafNodes(size_t pointCount)
 
 //-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------
 
+// The point-encode kernels operate on device data and are independent of the
+// resource type; only the trailing d_indx deallocation routes through ResourceT.
+// Dispatching the encode with if constexpr keeps it reachable for any ResourceT.
 template<typename BuildT, typename ResourceT>
 template<typename PtrT>
-inline void PointsToGrid<BuildT, ResourceT>::processPoints(const PtrT, size_t pointCount)
+inline void PointsToGrid<BuildT, ResourceT>::processPoints(const PtrT points, size_t pointCount)
 {
+    if constexpr(util::is_same<BuildT, Point>::value) this->encodePoints(points, pointCount);
     mResource->deallocate_async(mData.d_indx, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream);
-}
+}// PointsToGrid<BuildT, ResourceT>::processPoints
 
-//-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------
-
-// Template specialization with BuildT = Point
-template<>
+template<typename BuildT, typename ResourceT>
 template<typename PtrT>
-inline void PointsToGrid<Point>::processPoints(const PtrT points, size_t pointCount)
+inline void PointsToGrid<BuildT, ResourceT>::encodePoints(const PtrT points, size_t pointCount)
 {
     switch (mPointType){
     case PointType::Disable:
-        throw std::runtime_error("PointsToGrid<Point>::processPoints: mPointType == PointType::Disable\n");
+        throw std::runtime_error("PointsToGrid<Point, ResourceT>::encodePoints: mPointType == PointType::Disable\n");
     case PointType::PointID:
         util::cuda::lambdaKernel<<<numBlocks(pointCount), mNumThreads, 0, mStream>>>(pointCount, [=] __device__(size_t tid, PointsToGridData<Point> *d_data) {
             d_data->template getPoint<uint32_t>(tid) = d_data->d_indx[tid];
@@ -1227,12 +1235,9 @@ inline void PointsToGrid<Point>::processPoints(const PtrT points, size_t pointCo
         }, mDeviceData); cudaCheckError();
         break;
     default:
-        printf("Internal error in PointsToGrid<Point>::processPoints\n");
+        printf("Internal error in PointsToGrid<Point, ResourceT>::encodePoints\n");
     }
-    // This is the PointsToGrid<Point> (== <Point, DeviceResource>) member specialization,
-    // so ResourceT is not a name here; mResource is a DeviceResource* and the alignment is concrete.
-    mResource->deallocate_async(mData.d_indx, pointCount*sizeof(uint32_t), nanovdb::cuda::DeviceResource::DEFAULT_ALIGNMENT, mStream);
-}// PointsToGrid<Point>::processPoints
+}// PointsToGrid<BuildT, ResourceT>::encodePoints
 
 //-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------
 

nanovdb/nanovdb/unittest/TestMemoryResource.cu

@@ -209,4 +209,63 @@ TEST(TestMemoryResource, PointsToGrid_InjectedInstanceSeam)
     ASSERT_EQ(cudaFree(d_coords), cudaSuccess);
 }
 
+//======================================================================
+// A Point grid built through a custom (non-default) resource must still
+// encode its point coordinates. The encode path must not depend on the
+// resource type; only allocation does.
+//======================================================================
+
+TEST(TestMemoryResource, PointsToGrid_PointEncodedWithCustomResource)
+{
+    using BuildT = nanovdb::Point;
+    using Vec3T  = nanovdb::Vec3d;
+
+    const size_t pointCount = 256;
+    std::vector<Vec3T> points;
+    points.reserve(pointCount);
+    std::srand(98765);
+    const int max = 128, min = -max;
+    auto op = [&]() { return double(std::rand() % (max - min) + min); };
+    while (points.size() < pointCount) points.emplace_back(op(), op(), op());
+
+    Vec3T* d_points = nullptr;
+    const size_t pointSize = points.size() * sizeof(Vec3T);
+    ASSERT_EQ(cudaMalloc(&d_points, pointSize), cudaSuccess);
+    ASSERT_EQ(cudaMemcpy(d_points, points.data(), pointSize, cudaMemcpyHostToDevice), cudaSuccess);
+
+    const double voxelSize = 8.0;
+    CountingResource res;
+    nanovdb::tools::cuda::PointsToGrid<BuildT, CountingResource> converter(voxelSize, nanovdb::Vec3d(0.0), 0, res);
+    auto handle = converter.getHandle(d_points, pointCount);
+    ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess);
+    ASSERT_EQ(cudaFree(d_points), cudaSuccess);
+
+    EXPECT_TRUE(handle.deviceData());
+    EXPECT_TRUE(handle.deviceGrid<BuildT>());
+
+    handle.deviceDownload();
+    auto* grid = handle.grid<BuildT>();
+    ASSERT_TRUE(grid);
+    EXPECT_EQ(pointCount, grid->pointCount());
+
+    // Every input point must be recoverable from the encoded per-voxel point data.
+    nanovdb::PointAccessor<Vec3T, BuildT> acc(*grid);
+    ASSERT_TRUE(acc);
+    for (size_t i = 0; i < points.size(); ++i) {
+        const nanovdb::Coord ijk = grid->worldToIndex(points[i]).round();
+        ASSERT_TRUE(acc.probeLeaf(ijk) != nullptr);
+        ASSERT_TRUE(acc.isActive(ijk));
+        const Vec3T *start = nullptr, *stop = nullptr;
+        const uint64_t count = acc.voxelPoints(ijk, start, stop);
+        ASSERT_TRUE(start);
+        ASSERT_TRUE(stop);
+        bool found = false;
+        for (uint64_t j = 0; !found && j < count; ++j)
+            found = nanovdb::math::isApproxZero<double>((points[i] - start[j]).lengthSqr());
+        EXPECT_TRUE(found);
+    }
+
+    EXPECT_GT(res.allocs, 0);
+}
+
 } // unnamed namespace