Skip to content

Commit 4e02cba

Browse files
matthewdcongsifakis
matthewdcong
and
sifakis
authored
Adding coarsen/refine operators to NanoVDB (#2089)
* Adding coarsen/refine operators to NanoVDB * namespace adjustments, stricter const-correctness * adding grid coarsening/refinement operators in NanoVDB * adding unittest for refine/coarsen ops * ws fix * better forwarding to emplace_back * non-tbb version (refine), TBB guards, cleanup * ws fix * stricter const-ness for DeviceGridTraits methods * moving synchronization-sensitive initialization to process methods Signed-off-by: Matthew Cong <mcong@nvidia.com> * making helpers private; inlining/cleanup Signed-off-by: Efty Sifakis <esifakis@nvidia.com> --------- Signed-off-by: Matthew Cong <mcong@nvidia.com> Signed-off-by: Efty Sifakis <esifakis@nvidia.com> Co-authored-by: Eftychios Sifakis <esifakis@nvidia.com>
1 parent ddaf9f4 commit 4e02cba

17 files changed

Lines changed: 1396 additions & 156 deletions

‎nanovdb/nanovdb/examples/CMakeLists.txt‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -111,6 +111,8 @@ nanovdb_example(NAME "ex_raytrace_fog_volume")
111111
nanovdb_example(NAME "ex_raytrace_level_set")
112112
nanovdb_example(NAME "ex_dilate_nanovdb_cuda" OPENVDB)
113113
nanovdb_example(NAME "ex_merge_nanovdb_cuda" OPENVDB)
114+
nanovdb_example(NAME "ex_refine_nanovdb_cuda" OPENVDB)
115+
nanovdb_example(NAME "ex_coarsen_nanovdb_cuda" OPENVDB)
114116

115117
if(CUDAToolkit_FOUND)
116118
nanovdb_example(NAME "ex_make_mgpu_nanovdb") # requires cuRAND

‎nanovdb/nanovdb/examples/ex_coarsen_nanovdb_cuda/coarsen_nanovdb_cuda.cpp‎

Lines changed: 150 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,150 @@
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/Composite.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<class CoordT>
14+
inline CoordT
15+
coarsenCoord(const CoordT& coord)
16+
{
17+
auto coarsenComponent = [](const typename CoordT::ValueType n) {return (n>=0) ? (n>>1) : -((-n+1)>>1);};
18+
CoordT result;
19+
result[0] = coarsenComponent(coord[0]);
20+
result[1] = coarsenComponent(coord[1]);
21+
result[2] = coarsenComponent(coord[2]);
22+
return result;
23+
}
24+
25+
template<typename BuildT>
26+
void mainCoarsenGrid(
27+
nanovdb::NanoGrid<BuildT> *deviceGridOriginal,
28+
nanovdb::NanoGrid<BuildT> *deviceGridCoarsened,
29+
nanovdb::NanoGrid<BuildT> *indexGridOriginal,
30+
nanovdb::NanoGrid<BuildT> *indexGridCoarsened,
31+
uint32_t benchmark_iters
32+
);
33+
34+
/// @brief This example depends on OpenVDB, NanoVDB, and CUDA
35+
int main(int argc, char *argv[])
36+
{
37+
using GridT = openvdb::FloatGrid;
38+
using BuildT = nanovdb::ValueOnIndex;
39+
40+
openvdb::util::CpuTimer cpuTimer;
41+
const bool printGridDiagnostics = true;
42+
43+
try {
44+
if (argc<2) OPENVDB_THROW(openvdb::ValueError, "usage: "+std::string(argv[0])+" input.vdb [<iterations>]\n");
45+
int benchmark_iters = 10;
46+
if (argc > 2) sscanf(argv[2], "%d", &benchmark_iters);
47+
48+
// Read the initial level set from file
49+
50+
cpuTimer.start("Read input VDB file");
51+
openvdb::initialize();
52+
openvdb::io::File inFile(argv[1]);
53+
inFile.open(false); // disable delayed loading
54+
auto baseGrids = inFile.getGrids();
55+
inFile.close();
56+
auto grid = openvdb::gridPtrCast<GridT>(baseGrids->at(0));
57+
if (!grid) OPENVDB_THROW(openvdb::ValueError, "First grid is not a FloatGrid\n");
58+
cpuTimer.stop();
59+
60+
// Convert to indexGrid (original, un-coarsened)
61+
cpuTimer.start("Converting openVDB input to indexGrid (original version)");
62+
auto handleOriginal = nanovdb::tools::openToIndexVDB<BuildT, nanovdb::cuda::DeviceBuffer>(
63+
grid,
64+
0u, // Don't copy data channel
65+
false, // No stats
66+
false, // No tiles
67+
1 // Verbose mode
68+
);
69+
auto *indexGridOriginal = handleOriginal.grid<BuildT>();
70+
cpuTimer.stop();
71+
72+
if (printGridDiagnostics) {
73+
std::cout << "============ Original Grid ===========" << std::endl;
74+
std::cout << "Allocated values [valueCount()] : " << indexGridOriginal->valueCount() << std::endl;
75+
std::cout << "Active voxels [activeVoxelCount()] : " << indexGridOriginal->activeVoxelCount() << std::endl;
76+
auto minCorner = indexGridOriginal->indexBBox().min(), maxCorner = indexGridOriginal->indexBBox().max();
77+
std::cout << "Index-space bounding box : [" << minCorner.x() << "," << minCorner.y() << "," << minCorner.z()
78+
<< "] -> [" << maxCorner.x() << "," << maxCorner.y() << "," << maxCorner.z() << "]" << std::endl;
79+
std::cout << "Leaf nodes : " << indexGridOriginal->tree().nodeCount(0) << std::endl;
80+
std::cout << "Lower internal nodes : " << indexGridOriginal->tree().nodeCount(1) << std::endl;
81+
std::cout << "Upper internal nodes : " << indexGridOriginal->tree().nodeCount(2) << std::endl;
82+
std::cout << "Leaf-level occupancy : "
83+
<< 100.f * (float)(indexGridOriginal->activeVoxelCount())/(float)(indexGridOriginal->tree().nodeCount(0) * 512)
84+
<< "%" << std::endl;
85+
std::cout << "Memory usage : " << indexGridOriginal->gridSize() << " bytes" << std::endl;
86+
}
87+
// Coarsening (CPU/OpenVDB version)
88+
cpuTimer.start("Coarsening OpenVDB (on CPU)");
89+
using TreeT = GridT::TreeType;
90+
using LeafManagerT = openvdb::tree::LeafManager<const TreeT>;
91+
LeafManagerT leafMgr(grid->tree());
92+
auto coarsenedGrid = openvdb::FloatGrid::create(grid->background());
93+
coarsenedGrid->setTransform(grid->transform().copy());
94+
coarsenedGrid->setName(grid->getName());
95+
auto dstAcc = coarsenedGrid->getAccessor();
96+
for (std::size_t leafID = 0; leafID < leafMgr.leafCount(); ++leafID) {
97+
const auto& srcLeaf = leafMgr.leaf(leafID);
98+
for (auto iter = srcLeaf.cbeginValueOn(); iter; ++iter) {
99+
const auto dstCoord = coarsenCoord(iter.getCoord());
100+
if (!dstAcc.isValueOn(dstCoord))
101+
dstAcc.setValue(dstCoord, iter.getValue());
102+
}
103+
}
104+
cpuTimer.stop();
105+
106+
// Convert to indexGrid (coarsened)
107+
cpuTimer.start("Converting openVDB input to indexGrid (coarsened version)");
108+
auto handleCoarsened = nanovdb::tools::openToIndexVDB<BuildT, nanovdb::cuda::DeviceBuffer>(
109+
coarsenedGrid,
110+
0u, // Don't copy data channel
111+
false, // No stats
112+
false, // No tiles
113+
1 // Verbose mode
114+
);
115+
cpuTimer.stop();
116+
117+
auto *indexGridCoarsened = handleCoarsened.grid<BuildT>();
118+
119+
if (printGridDiagnostics) {
120+
std::cout << "=========== Coarsened Grid ===========" << std::endl;
121+
std::cout << "Allocated values [valueCount()] : " << indexGridCoarsened->valueCount() << std::endl;
122+
std::cout << "Active voxels [activeVoxelCount()] : " << indexGridCoarsened->activeVoxelCount() << std::endl;
123+
auto minCorner = indexGridCoarsened->indexBBox().min(), maxCorner = indexGridCoarsened->indexBBox().max();
124+
std::cout << "Index-space bounding box : [" << minCorner.x() << "," << minCorner.y() << "," << minCorner.z()
125+
<< "] -> [" << maxCorner.x() << "," << maxCorner.y() << "," << maxCorner.z() << "]" << std::endl;
126+
std::cout << "Leaf nodes : " << indexGridCoarsened->tree().nodeCount(0) << std::endl;
127+
std::cout << "Lower internal nodes : " << indexGridCoarsened->tree().nodeCount(1) << std::endl;
128+
std::cout << "Upper internal nodes : " << indexGridCoarsened->tree().nodeCount(2) << std::endl;
129+
std::cout << "Leaf-level occupancy : "
130+
<< 100.f * (float)(indexGridCoarsened->activeVoxelCount())/(float)(indexGridCoarsened->tree().nodeCount(0) * 512)
131+
<< "%" << std::endl;
132+
std::cout << "Memory usage : " << indexGridCoarsened->gridSize() << " bytes" << std::endl;
133+
}
134+
135+
// Copy both NanoVDB grids to GPU
136+
handleOriginal.deviceUpload();
137+
handleCoarsened.deviceUpload();
138+
auto* deviceGridOriginal = handleOriginal.deviceGrid<BuildT>();
139+
auto* deviceGridCoarsened = handleCoarsened.deviceGrid<BuildT>();
140+
if (!deviceGridOriginal || !deviceGridCoarsened)
141+
OPENVDB_THROW(openvdb::RuntimeError, "Failure while uploading indexGrids to GPU");
142+
143+
// Launch benchmark
144+
mainCoarsenGrid( deviceGridOriginal, deviceGridCoarsened, indexGridOriginal, indexGridCoarsened, benchmark_iters );
145+
}
146+
catch (const std::exception& e) {
147+
std::cerr << "An exception occurred: \"" << e.what() << "\"" << std::endl;
148+
}
149+
return 0;
150+
}

‎nanovdb/nanovdb/examples/ex_coarsen_nanovdb_cuda/coarsen_nanovdb_cuda_kernels.cu‎

Lines changed: 56 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,56 @@
1+
// Copyright Contributors to the OpenVDB Project
2+
// SPDX-License-Identifier: Apache-2.0
3+
4+
#include <nanovdb/tools/cuda/CoarsenGrid.cuh>
5+
6+
template<typename T>
7+
bool bufferCheck(const T* deviceBuffer, const T* hostBuffer, size_t elem_count) {
8+
T* tmpBuffer = new T[elem_count];
9+
cudaCheck(cudaMemcpy(tmpBuffer, deviceBuffer, elem_count * sizeof(T), cudaMemcpyDeviceToHost));
10+
bool same = true;
11+
for (int i=0; same && i< elem_count; ++i) { same = (tmpBuffer[i] == hostBuffer[i]); }
12+
delete [] tmpBuffer;
13+
return same;
14+
}
15+
16+
template<typename BuildT>
17+
void mainCoarsenGrid(
18+
nanovdb::NanoGrid<BuildT> *deviceGridOriginal,
19+
nanovdb::NanoGrid<BuildT> *deviceGridCoarsened,
20+
nanovdb::NanoGrid<BuildT> *indexGridOriginal,
21+
nanovdb::NanoGrid<BuildT> *indexGridCoarsened,
22+
uint32_t benchmark_iters)
23+
{
24+
nanovdb::util::cuda::Timer gpuTimer;
25+
26+
// Initialize coarsener
27+
nanovdb::tools::cuda::CoarsenGrid<BuildT> coarsener( deviceGridOriginal );
28+
coarsener.setChecksum(nanovdb::CheckMode::Default);
29+
coarsener.setVerbose(1);
30+
31+
auto handle = coarsener.getHandle();
32+
auto dstGrid = handle.template deviceGrid<BuildT>();
33+
34+
// Check for correctness
35+
if (bufferCheck((char*)dstGrid, (char*)indexGridCoarsened->data(), indexGridCoarsened->gridSize()))
36+
std::cout << "Result of CoarsenGrid check out CORRECT against reference" << std::endl;
37+
else
38+
std::cout << "Result of CoarsenGrid compares INCORRECT against reference" << std::endl;
39+
40+
// Re-run warm-started iterations
41+
coarsener.setVerbose(0);
42+
for (int i = 0; i < benchmark_iters; i++) {
43+
gpuTimer.start("Re-running entire coarsening after warmstart");
44+
auto dummyHandle = coarsener.getHandle();
45+
gpuTimer.stop();
46+
}
47+
}
48+
49+
template
50+
void mainCoarsenGrid(
51+
nanovdb::NanoGrid<nanovdb::ValueOnIndex> *deviceGridOriginal,
52+
nanovdb::NanoGrid<nanovdb::ValueOnIndex> *deviceGridCoarsened,
53+
nanovdb::NanoGrid<nanovdb::ValueOnIndex> *indexGridOriginal,
54+
nanovdb::NanoGrid<nanovdb::ValueOnIndex> *indexGridCoarsened,
55+
uint32_t benchmark_iters
56+
);

0 commit comments

Comments
 (0)