Skip to content

Commit 06f1a38

Browse files
swahtzmatthewdcong
swahtz
authored and
matthewdcong
committed
Cumulative NanoVDB updates from feature/fvdb
--------- Signed-off-by: Jonathan Swartz <jonathan@jswartz.info> Signed-off-by: Matthew Cong <mcong@nvidia.com>
1 parent 7969cd8 commit 06f1a38

25 files changed

Lines changed: 7816 additions & 55 deletions

‎nanovdb/nanovdb/PNanoVDB.h‎

Lines changed: 342 additions & 26 deletions
Large diffs are not rendered by default.

‎nanovdb/nanovdb/PNanoVDB2.h‎

Lines changed: 1301 additions & 0 deletions
Large diffs are not rendered by default.

‎nanovdb/nanovdb/examples/CMakeLists.txt‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -109,6 +109,8 @@ nanovdb_example(NAME "ex_bump_pool_buffer")
109109
nanovdb_example(NAME "ex_collide_level_set")
110110
nanovdb_example(NAME "ex_raytrace_fog_volume")
111111
nanovdb_example(NAME "ex_raytrace_level_set")
112+
nanovdb_example(NAME "ex_dilate_nanovdb_cuda" OPENVDB)
113+
nanovdb_example(NAME "ex_merge_nanovdb_cuda" OPENVDB)
112114

113115
if(CUDAToolkit_FOUND)
114116
nanovdb_example(NAME "ex_make_mgpu_nanovdb") # requires cuRAND

‎nanovdb/nanovdb/examples/ex_dilate_nanovdb_cuda/dilate_nanovdb_cuda.cpp‎

Lines changed: 133 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,133 @@
1+
// Copyright Contributors to the OpenVDB Project
2+
// SPDX-License-Identifier: Apache-2.0
3+
4+
// the following files are from OpenVDB
5+
#include <openvdb/tools/Morphology.h>
6+
#include <openvdb/util/CpuTimer.h>
7+
8+
// the following files are from NanoVDB
9+
#include <nanovdb/NanoVDB.h>
10+
#include <nanovdb/cuda/DeviceBuffer.h>
11+
#include <nanovdb/tools/CreateNanoGrid.h>
12+
13+
template<typename BuildT>
14+
void mainDilateGrid(
15+
nanovdb::NanoGrid<BuildT> *deviceGridOriginal,
16+
nanovdb::NanoGrid<BuildT> *deviceGridDilated,
17+
nanovdb::NanoGrid<BuildT> *indexGridOriginal,
18+
nanovdb::NanoGrid<BuildT> *indexGridDilated,
19+
uint32_t nnType,
20+
uint32_t benchmark_iters
21+
);
22+
23+
/// @brief This example depends on OpenVDB, NanoVDB, and CUDA
24+
int main(int argc, char *argv[])
25+
{
26+
using GridT = openvdb::FloatGrid;
27+
using BuildT = nanovdb::ValueOnIndex;
28+
29+
// Select the type of dilation here. The NN_EDGE case supports leaf dilation too (currently)
30+
// openvdb::tools::NearestNeighbors nnType = openvdb::tools::NN_FACE_EDGE_VERTEX;
31+
openvdb::tools::NearestNeighbors nnType = openvdb::tools::NN_FACE;
32+
33+
openvdb::util::CpuTimer cpuTimer;
34+
const bool printGridDiagnostics = true;
35+
36+
try {
37+
38+
if (argc<2) OPENVDB_THROW(openvdb::ValueError, "usage: "+std::string(argv[0])+" input.vdb [<iterations>]\n");
39+
int benchmark_iters = 10;
40+
if (argc > 2) sscanf(argv[2], "%d", &benchmark_iters);
41+
42+
// Read the initial level set from file
43+
44+
cpuTimer.start("Read input VDB file");
45+
openvdb::initialize();
46+
openvdb::io::File inFile(argv[1]);
47+
inFile.open(false); // disable delayed loading
48+
auto baseGrids = inFile.getGrids();
49+
inFile.close();
50+
auto grid = openvdb::gridPtrCast<GridT>(baseGrids->at(0));
51+
openvdb::FloatGrid* ptr = grid.get(); // raw pointer
52+
if (!grid) OPENVDB_THROW(openvdb::ValueError, "First grid is not a FloatGrid\n");
53+
cpuTimer.stop();
54+
55+
// Convert to indexGrid (original, un-dilated)
56+
cpuTimer.start("Converting openVDB input to indexGrid (original version)");
57+
auto handleOriginal = nanovdb::tools::openToIndexVDB<BuildT, nanovdb::cuda::DeviceBuffer>(
58+
grid,
59+
0u, // Don't copy data channel
60+
false, // No stats
61+
false, // No tiles
62+
1 // Verbose mode
63+
);
64+
auto *indexGridOriginal = handleOriginal.grid<BuildT>();
65+
cpuTimer.stop();
66+
67+
if (printGridDiagnostics) {
68+
std::cout << "============ Original Grid ===========" << std::endl;
69+
std::cout << "Allocated values [valueCount()] : " << indexGridOriginal->valueCount() << std::endl;
70+
std::cout << "Active voxels [activeVoxelCount()] : " << indexGridOriginal->activeVoxelCount() << std::endl;
71+
auto minCorner = indexGridOriginal->indexBBox().min(), maxCorner = indexGridOriginal->indexBBox().max();
72+
std::cout << "Index-space bounding box : [" << minCorner.x() << "," << minCorner.y() << "," << minCorner.z()
73+
<< "] -> [" << maxCorner.x() << "," << maxCorner.y() << "," << maxCorner.z() << "]" << std::endl;
74+
std::cout << "Leaf nodes : " << indexGridOriginal->tree().nodeCount(0) << std::endl;
75+
std::cout << "Lower internal nodes : " << indexGridOriginal->tree().nodeCount(1) << std::endl;
76+
std::cout << "Upper internal nodes : " << indexGridOriginal->tree().nodeCount(2) << std::endl;
77+
std::cout << "Leaf-level occupancy : "
78+
<< 100.f * (float)(indexGridOriginal->activeVoxelCount())/(float)(indexGridOriginal->tree().nodeCount(0) * 512)
79+
<< "%" << std::endl;
80+
std::cout << "Memory usage : " << indexGridOriginal->gridSize() << " bytes" << std::endl;
81+
}
82+
83+
// Dilation (CPU/OpenVDB version)
84+
cpuTimer.start("Dilating openVDB (on CPU)");
85+
openvdb::tools::dilateActiveValues(grid->tree(), 1, nnType);
86+
cpuTimer.stop();
87+
88+
// Convert to indexGrid (dilated)
89+
cpuTimer.start("Converting openVDB input to indexGrid (dilated version)");
90+
auto handleDilated = nanovdb::tools::openToIndexVDB<BuildT, nanovdb::cuda::DeviceBuffer>(
91+
grid,
92+
0u, // Don't copy data channel
93+
false, // No stats
94+
false, // No tiles
95+
1 // Verbose mode
96+
);
97+
cpuTimer.stop();
98+
99+
auto *indexGridDilated = handleDilated.grid<BuildT>();
100+
101+
if (printGridDiagnostics) {
102+
std::cout << "============ Dilated Grid ============" << std::endl;
103+
std::cout << "Allocated values [valueCount()] : " << indexGridDilated->valueCount() << std::endl;
104+
std::cout << "Active voxels [activeVoxelCount()] : " << indexGridDilated->activeVoxelCount() << std::endl;
105+
auto minCorner = indexGridDilated->indexBBox().min(), maxCorner = indexGridDilated->indexBBox().max();
106+
std::cout << "Index-space bounding box : [" << minCorner.x() << "," << minCorner.y() << "," << minCorner.z()
107+
<< "] -> [" << maxCorner.x() << "," << maxCorner.y() << "," << maxCorner.z() << "]" << std::endl;
108+
std::cout << "Leaf nodes : " << indexGridDilated->tree().nodeCount(0) << std::endl;
109+
std::cout << "Lower internal nodes : " << indexGridDilated->tree().nodeCount(1) << std::endl;
110+
std::cout << "Upper internal nodes : " << indexGridDilated->tree().nodeCount(2) << std::endl;
111+
std::cout << "Leaf-level occupancy : "
112+
<< 100.f * (float)(indexGridDilated->activeVoxelCount())/(float)(indexGridDilated->tree().nodeCount(0) * 512)
113+
<< "%" << std::endl;
114+
std::cout << "Memory usage : " << indexGridDilated->gridSize() << " bytes" << std::endl;
115+
}
116+
117+
// Copy both NanoVDB grids to GPU
118+
handleOriginal.deviceUpload();
119+
handleDilated.deviceUpload();
120+
auto* deviceGridOriginal = handleOriginal.deviceGrid<BuildT>();
121+
auto* deviceGridDilated = handleDilated.deviceGrid<BuildT>();
122+
if (!deviceGridOriginal || !deviceGridDilated)
123+
OPENVDB_THROW(openvdb::RuntimeError, "Failure while uploading indexGrids to GPU");
124+
125+
// Launch benchmark
126+
mainDilateGrid( deviceGridOriginal, deviceGridDilated, indexGridOriginal, indexGridDilated, nnType, benchmark_iters );
127+
128+
}
129+
catch (const std::exception& e) {
130+
std::cerr << "An exception occurred: \"" << e.what() << "\"" << std::endl;
131+
}
132+
return 0;
133+
}

‎nanovdb/nanovdb/examples/ex_dilate_nanovdb_cuda/dilate_nanovdb_cuda_kernels.cu‎

Lines changed: 102 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,102 @@
1+
// Copyright Contributors to the OpenVDB Project
2+
// SPDX-License-Identifier: Apache-2.0
3+
4+
#include <nanovdb/tools/cuda/DilateGrid.cuh>
5+
#include <nanovdb/tools/cuda/PruneGrid.cuh>
6+
#include <nanovdb/util/cuda/Injection.cuh>
7+
8+
template<typename T>
9+
bool bufferCheck(const T* deviceBuffer, const T* hostBuffer, size_t elem_count) {
10+
T* tmpBuffer = new T[elem_count];
11+
cudaCheck(cudaMemcpy(tmpBuffer, deviceBuffer, elem_count * sizeof(T), cudaMemcpyDeviceToHost));
12+
bool same = true;
13+
for (int i=0; same && i< elem_count; ++i) { same = (tmpBuffer[i] == hostBuffer[i]); }
14+
delete [] tmpBuffer;
15+
return same;
16+
}
17+
18+
template<typename BuildT>
19+
void mainDilateGrid(
20+
nanovdb::NanoGrid<BuildT> *deviceGridOriginal,
21+
nanovdb::NanoGrid<BuildT> *deviceGridDilated,
22+
nanovdb::NanoGrid<BuildT> *indexGridOriginal,
23+
nanovdb::NanoGrid<BuildT> *indexGridDilated,
24+
uint32_t nnType,
25+
uint32_t benchmark_iters)
26+
{
27+
nanovdb::util::cuda::Timer gpuTimer;
28+
29+
// Initialize dilator
30+
nanovdb::tools::cuda::DilateGrid<BuildT> dilator( deviceGridOriginal );
31+
dilator.setOperation(nanovdb::tools::morphology::NearestNeighbors(nnType));
32+
dilator.setChecksum(nanovdb::CheckMode::Default);
33+
dilator.setVerbose(1);
34+
35+
auto handle = dilator.getHandle();
36+
auto dstGrid = handle.template deviceGrid<BuildT>();
37+
38+
// Check for correctness
39+
if (bufferCheck((char*)dstGrid, (char*)indexGridDilated->data(), indexGridDilated->gridSize()))
40+
std::cout << "Result of DilateGrid check out CORRECT against reference" << std::endl;
41+
else
42+
std::cout << "Result of DilateGrid compares INCORRECT against reference" << std::endl;
43+
44+
// Re-run warm-started iterations
45+
dilator.setVerbose(0);
46+
for (int i = 0; i < benchmark_iters; i++) {
47+
gpuTimer.start("Re-running entire dilation after warmstart");
48+
auto dummyHandle = dilator.getHandle();
49+
gpuTimer.stop();
50+
}
51+
52+
uint32_t dstLeafCount = nanovdb::util::cuda::DeviceGridTraits<BuildT>::getTreeData(dstGrid).mNodeCount[0];
53+
nanovdb::cuda::DeviceBuffer dstLeafMaskBuffer;
54+
nanovdb::Mask<3>* dstLeafMasks = nullptr;
55+
if (dstLeafCount) {
56+
dstLeafMaskBuffer = nanovdb::cuda::DeviceBuffer::create( std::size_t(dstLeafCount) * sizeof(nanovdb::Mask<3>), nullptr, false );
57+
dstLeafMasks = static_cast<nanovdb::Mask<3>*>(dstLeafMaskBuffer.deviceData());
58+
if (!dstLeafMasks) throw std::runtime_error("No GPU buffer for dstLeafMask");
59+
}
60+
61+
const unsigned int numThreads = 128;
62+
auto numBlocks = [numThreads] (unsigned int n) {return (n + numThreads - 1) / numThreads;};
63+
gpuTimer.start("Injecting un-dilated topology as a pruning mask");
64+
if (dstLeafCount)
65+
nanovdb::util::cuda::lambdaKernel<<<numBlocks(dstLeafCount), numThreads>>>(dstLeafCount,
66+
nanovdb::util::cuda::InjectGridMaskFunctor<BuildT>(),
67+
deviceGridOriginal, dstGrid, dstLeafMasks );
68+
gpuTimer.stop();
69+
70+
// Initialize pruner
71+
nanovdb::tools::cuda::PruneGrid<BuildT> pruner( dstGrid, dstLeafMasks );
72+
pruner.setChecksum(nanovdb::CheckMode::Default);
73+
pruner.setVerbose(1);
74+
75+
auto prunedHandle = pruner.getHandle();
76+
auto prunedGrid = prunedHandle.template deviceGrid<BuildT>();
77+
78+
// Check for correctness
79+
if (bufferCheck((char*)prunedGrid, (char*)indexGridOriginal->data(), indexGridOriginal->gridSize()))
80+
std::cout << "Result of PruneGrid check out CORRECT against reference" << std::endl;
81+
else
82+
std::cout << "Result of PruneGrid compares INCORRECT against reference" << std::endl;
83+
84+
// Re-run warm-started iterations
85+
pruner.setVerbose(0);
86+
for (int i = 0; i < benchmark_iters; i++) {
87+
gpuTimer.start("Re-running entire pruning after warmstart");
88+
auto dummyHandle = pruner.getHandle();
89+
gpuTimer.stop();
90+
}
91+
92+
}
93+
94+
template
95+
void mainDilateGrid(
96+
nanovdb::NanoGrid<nanovdb::ValueOnIndex> *deviceGridOriginal,
97+
nanovdb::NanoGrid<nanovdb::ValueOnIndex> *deviceGridDilated,
98+
nanovdb::NanoGrid<nanovdb::ValueOnIndex> *indexGridOriginal,
99+
nanovdb::NanoGrid<nanovdb::ValueOnIndex> *indexGridDilated,
100+
uint32_t nnType,
101+
uint32_t benchmark_iters
102+
);

0 commit comments

Comments
 (0)