From 9f73006b975a2efdde975277ae14092bdb5a6afe Mon Sep 17 00:00:00 2001 From: Andreas Henne Date: Fri, 14 Mar 2025 15:56:10 +0100 Subject: [PATCH] Fixed an issue in nanovdb::tools::cuda::indexToGrid that occured when building a grid with internal nodes that differed in size from the size of internal nodes of the index grid. --- nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh | 12 +++- nanovdb/nanovdb/unittest/TestNanoVDB.cu | 78 ++++++++++++---------- 2 files changed, 55 insertions(+), 35 deletions(-) diff --git a/nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh b/nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh index 9e8431260c..6db3108e25 100644 --- a/nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh @@ -183,7 +183,17 @@ __global__ void processRootTilesKernel(typename IndexToGrid::NodeAcce auto &dstTile = *nodeAcc->template dstRoot().tile(tid); dstTile.key = srcTile.key; if (srcTile.child) { - dstTile.child = sizeof(NanoRoot) + sizeof(NanoRoot::Tile)*((srcTile.child - sizeof(NanoRoot))/sizeof(NanoRoot::Tile)); + using SrcTileT = typename NanoRoot::Tile; + using DstTileT = typename NanoRoot::Tile; + using SrcChildT = typename NanoRoot::ChildNodeType; + using DstChildT = typename NanoRoot::ChildNodeType; + const uint64_t nodeSkip = nodeAcc->nodeCount[3]; + const uint64_t srcOff = sizeof(SrcTileT)*nodeSkip + sizeof(NanoRoot); + const uint64_t dstOff = sizeof(DstTileT)*nodeSkip + sizeof(NanoRoot); + + const uint64_t childID = (srcTile.child - srcOff)/sizeof(SrcChildT); + dstTile.child = dstOff + childID*sizeof(DstChildT); + dstTile.value = srcValues[0];// set to background dstTile.state = false; } else { diff --git a/nanovdb/nanovdb/unittest/TestNanoVDB.cu b/nanovdb/nanovdb/unittest/TestNanoVDB.cu index 6896c6ca09..eda8483dc5 100644 --- a/nanovdb/nanovdb/unittest/TestNanoVDB.cu +++ b/nanovdb/nanovdb/unittest/TestNanoVDB.cu @@ -1062,41 +1062,51 @@ TEST(TestNanoVDBCUDA, CudaIndexGridToGrid_ValueOnIndex) idxHdl.deviceUpload(); auto *idxGrid = idxHdl.grid(); EXPECT_TRUE(idxGrid); - //timer.restart("Create value list on CPU"); - float *values = new float[idxGrid->valueCount()], *d_values = nullptr; - values[0] = floatGrid->tree().root().background(); - for (auto it = floatGrid->indexBBox().begin(); it; ++it) { - EXPECT_EQ(acc.isActive(*it), idxGrid->tree().isActive(*it)); - if (acc.isActive(*it)) { - const uint64_t idx = idxGrid->tree().getValue(*it); - EXPECT_TRUE(idx < idxGrid->valueCount()); - values[idx] = acc.getValue(*it); + + auto checkConvertedValueGrid = [&](const auto& fromOriginalData, const auto& toOriginalData) { + using FloatType = decltype(fromOriginalData(0.0f)); + + //timer.restart("Create value list on CPU"); + FloatType *values = new FloatType[idxGrid->valueCount()], *d_values = nullptr; + values[0] = fromOriginalData(floatGrid->tree().root().background()); + for (auto it = floatGrid->indexBBox().begin(); it; ++it) { + EXPECT_EQ(acc.isActive(*it), idxGrid->tree().isActive(*it)); + if (acc.isActive(*it)) { + const uint64_t idx = idxGrid->tree().getValue(*it); + EXPECT_TRUE(idx < idxGrid->valueCount()); + values[idx] = fromOriginalData(acc.getValue(*it)); + } } - } - //timer.restart("Allocate and copy values from CPU to GPU"); - cudaCheck(cudaMalloc((void**)&d_values, idxGrid->valueCount()*sizeof(float))); - cudaCheck(cudaMemcpy(d_values, values, idxGrid->valueCount()*sizeof(float), cudaMemcpyHostToDevice)); - EXPECT_FALSE(idxHdl.deviceGrid()); - auto *d_idxGrid = idxHdl.deviceGrid(); - EXPECT_TRUE(d_idxGrid); - //timer.restart("Call CudaIndexToGrid"); - auto hdl = nanovdb::tools::cuda::indexToGrid(d_idxGrid, d_values); - //timer.restart("unit-test"); - EXPECT_FALSE(hdl.grid());// no host grid - EXPECT_TRUE(hdl.deviceGrid()); - hdl.deviceDownload(); - auto *floatGrid2 = hdl.grid(); - EXPECT_TRUE(floatGrid2); - auto acc2 = floatGrid2->getAccessor(); - EXPECT_EQ(floatGrid->indexBBox(), floatGrid2->indexBBox()); - EXPECT_EQ(floatGrid->worldBBox(), floatGrid2->worldBBox()); - EXPECT_EQ(floatGrid->tree().root().background(), floatGrid2->tree().root().background()); - for (auto it = floatGrid->indexBBox().begin(); it; ++it) { - EXPECT_EQ(acc.isActive(*it), acc2.isActive(*it)); - if (acc.isActive(*it)) EXPECT_EQ(acc.getValue(*it), acc2.getValue(*it)); - } - //timer.stop(); - cudaFree(d_values); + //timer.restart("Allocate and copy values from CPU to GPU"); + cudaCheck(cudaMalloc((void**)&d_values, idxGrid->valueCount()*sizeof(FloatType))); + cudaCheck(cudaMemcpy(d_values, values, idxGrid->valueCount()*sizeof(FloatType), cudaMemcpyHostToDevice)); + EXPECT_FALSE(idxHdl.deviceGrid()); + auto *d_idxGrid = idxHdl.deviceGrid(); + EXPECT_TRUE(d_idxGrid); + //timer.restart("Call CudaIndexToGrid"); + auto hdl = nanovdb::tools::cuda::indexToGrid(d_idxGrid, d_values); + //timer.restart("unit-test"); + EXPECT_FALSE(hdl.template grid());// no host grid + EXPECT_TRUE(hdl.template deviceGrid()); + hdl.deviceDownload(); + auto *floatGrid2 = hdl.template grid(); + EXPECT_TRUE(floatGrid2); + auto acc2 = floatGrid2->getAccessor(); + EXPECT_EQ(floatGrid->indexBBox(), floatGrid2->indexBBox()); + EXPECT_EQ(floatGrid->worldBBox(), floatGrid2->worldBBox()); + EXPECT_EQ(floatGrid->tree().root().background(), toOriginalData(floatGrid2->tree().root().background())); + for (auto it = floatGrid->indexBBox().begin(); it; ++it) { + EXPECT_EQ(acc.isActive(*it), acc2.isActive(*it)); + if (acc.isActive(*it)) EXPECT_EQ(acc.getValue(*it), toOriginalData(acc2.getValue(*it))); + } + //timer.stop(); + cudaFree(d_values); + }; + checkConvertedValueGrid([](float x) { return x; }, [](float x) { return x; }); + + // Convert index grid to grid of Vec3fs, whereat the original float data is just stored in all components of a Vec3f. + // This test covers code in indexToGrid that is only relevant if the size of grid data changes and does not align. + checkConvertedValueGrid([](float x) { return nanovdb::Vec3f(x); }, [](const nanovdb::Vec3f& x) { return x[0]; }); }// CudaPointToGrid_ValueOnIndex TEST(TestNanoVDBCUDA, CudaSignedFloodFill)