Skip to content

Commit 72e53f8

Browse files
sifakissifakis
committed
Add MeshToGrid: CUDA triangle mesh to NanoVDB UDF/IndexGrid converter
Introduces nanovdb::tools::cuda::MeshToGrid<BuildT>, a GPU-accelerated rasterizer that converts a triangle soup (vertex list + index list) into a sparse NanoVDB indexGrid topology. The implementation uses a hierarchical 3-pass pipeline: (1) index-space triangle transformation, (2) root-tile / triangle AABB pair enumeration via CUB prefix-sum + scatter, and (3) recursive 8x subdivision with triangle-AABB / SAT intersection tests down to leaf resolution. The current endpoint is an indexGrid topology (no sidecar distance data); UDF and SDF outputs are planned. Includes the ex_mesh_to_grid_cuda example and a minimal addition to nanovdb/nanovdb/examples/CMakeLists.txt. Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
1 parent d23549f commit 72e53f8

4 files changed

Lines changed: 1021 additions & 0 deletions

File tree

‎nanovdb/nanovdb/examples/CMakeLists.txt‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -113,6 +113,7 @@ nanovdb_example(NAME "ex_dilate_nanovdb_cuda" OPENVDB)
113113
nanovdb_example(NAME "ex_merge_nanovdb_cuda" OPENVDB)
114114
nanovdb_example(NAME "ex_refine_nanovdb_cuda" OPENVDB)
115115
nanovdb_example(NAME "ex_coarsen_nanovdb_cuda" OPENVDB)
116+
nanovdb_example(NAME "ex_mesh_to_grid_cuda" OPENVDB)
116117

117118
if(CUDAToolkit_FOUND)
118119
nanovdb_example(NAME "ex_make_mgpu_nanovdb") # requires cuRAND

‎nanovdb/nanovdb/examples/ex_mesh_to_grid_cuda/mesh_to_grid_cuda.cpp‎

Lines changed: 181 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,181 @@
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+
#include <openvdb/tools/MeshToVolume.h>
8+
9+
// the following files are from NanoVDB
10+
#include <nanovdb/NanoVDB.h>
11+
#include <nanovdb/cuda/DeviceBuffer.h>
12+
#include <nanovdb/tools/CreateNanoGrid.h>
13+
14+
#include <thrust/universal_vector.h>
15+
16+
template<typename BuildT>
17+
void mainMeshToGrid(
18+
const nanovdb::Vec3f *devicePoints,
19+
const int pointCount,
20+
const nanovdb::Vec3i *deviceTriangles,
21+
const int triangleCount,
22+
const nanovdb::Map map);
23+
24+
void readOBJ(const std::string& filename,
25+
std::vector<openvdb::Vec3s>& points,
26+
std::vector<openvdb::Vec3I>& triangles,
27+
std::vector<openvdb::Vec4I>& quads)
28+
{
29+
std::ifstream file(filename);
30+
if (!file.is_open()) {
31+
OPENVDB_THROW(openvdb::IoError, "Failed to open OBJ file: " + filename);
32+
}
33+
34+
std::string line;
35+
int lineNumber = 0;
36+
37+
while (std::getline(file, line)) {
38+
lineNumber++;
39+
std::istringstream iss(line);
40+
std::string type;
41+
iss >> type;
42+
43+
if (type == "v") {
44+
float x, y, z;
45+
iss >> x >> y >> z;
46+
points.push_back(openvdb::Vec3s(x, y, z));
47+
} else if (type == "f") {
48+
std::vector<int> faceIndices;
49+
std::string vertexData;
50+
51+
while (iss >> vertexData) {
52+
// Isolate the vertex index (everything before the first slash)
53+
size_t slashPos = vertexData.find('/');
54+
std::string indexStr = vertexData.substr(0, slashPos);
55+
56+
if (indexStr.empty()) continue;
57+
58+
int raw_idx = std::stoi(indexStr);
59+
int actual_idx = 0;
60+
61+
// Handle negative indices: relative to the number of points parsed so far
62+
if (raw_idx < 0) {
63+
actual_idx = points.size() + raw_idx;
64+
} else {
65+
// Standard positive indices: OBJ is 1-based, convert to 0-based for C++
66+
actual_idx = raw_idx - 1;
67+
}
68+
69+
// Strict bounds checking to prevent segfaults
70+
if (actual_idx < 0 || actual_idx >= points.size()) {
71+
OPENVDB_THROW(openvdb::ValueError,
72+
"OBJ parse error on line " + std::to_string(lineNumber) +
73+
": Face index out of bounds (Raw: " + std::to_string(raw_idx) +
74+
", Computed: " + std::to_string(actual_idx) + ", Total Points: " +
75+
std::to_string(points.size()) + ")");
76+
}
77+
78+
faceIndices.push_back(actual_idx);
79+
}
80+
81+
// Add to the appropriate OpenVDB list
82+
if (faceIndices.size() == 3) {
83+
triangles.push_back(openvdb::Vec3I(faceIndices[0], faceIndices[1], faceIndices[2]));
84+
} else if (faceIndices.size() == 4) {
85+
quads.push_back(openvdb::Vec4I(faceIndices[0], faceIndices[1], faceIndices[2], faceIndices[3]));
86+
} else if (faceIndices.size() > 4) {
87+
std::cerr << "Warning on line " << lineNumber << ": Skipping face with "
88+
<< faceIndices.size() << " vertices. Triangulate your mesh!" << std::endl;
89+
}
90+
}
91+
}
92+
}
93+
94+
/// @brief This example depends on OpenVDB, NanoVDB, and CUDA
95+
int main(int argc, char *argv[])
96+
{
97+
using GridT = openvdb::FloatGrid;
98+
using BuildT = nanovdb::ValueOnIndex;
99+
100+
// Select the type of dilation here. The NN_EDGE case supports leaf dilation too (currently)
101+
// openvdb::tools::NearestNeighbors nnType = openvdb::tools::NN_FACE_EDGE_VERTEX;
102+
openvdb::tools::NearestNeighbors nnType = openvdb::tools::NN_FACE;
103+
104+
openvdb::util::CpuTimer cpuTimer;
105+
106+
try {
107+
108+
if (argc<2) OPENVDB_THROW(openvdb::ValueError, "usage: "+std::string(argv[0])+" input.obj [output.vdb]\n");
109+
std::string inputFile = argv[1];
110+
std::string outputFile = "output.vdb";
111+
if (argc > 2)
112+
outputFile = argv[2];
113+
float voxelSize = 0.001f;
114+
if (argc > 3)
115+
voxelSize = atof(argv[3]);
116+
117+
std::vector<openvdb::Vec3s> openvdb_points;
118+
std::vector<openvdb::Vec3I> openvdb_triangles;
119+
std::vector<openvdb::Vec4I> quads;
120+
121+
// Read the OBJ file
122+
std::cout << "Reading " << inputFile << "..." << std::endl;
123+
readOBJ(inputFile, openvdb_points, openvdb_triangles, quads);
124+
std::cout << "Loaded " << openvdb_points.size() << " vertices, "
125+
<< openvdb_triangles.size() << " openvdb_triangles, and "
126+
<< quads.size() << " quads." << std::endl;
127+
128+
// Initialize OpenVDB
129+
openvdb::initialize();
130+
131+
// Setup Grid Transform (Voxel Size)
132+
openvdb::math::Transform::Ptr transform =
133+
openvdb::math::Transform::createLinearTransform(voxelSize);
134+
135+
// Convert Mesh to Level Set (SDF)
136+
// halfband specifies the half-width of the narrow band in voxel units
137+
float halfband = 3.0f;
138+
cpuTimer.start("Converting mesh to OpenVDB level set");
139+
openvdb::FloatGrid::Ptr grid = openvdb::tools::meshToLevelSet<openvdb::FloatGrid>(
140+
*transform, openvdb_points, openvdb_triangles, quads, halfband);
141+
cpuTimer.stop();
142+
143+
144+
// Write the Grid to a VDB File
145+
grid->setName("LevelSet");
146+
grid->print(std::cout, 2);
147+
std::cout << "Writing to " << outputFile << "..." << std::endl;
148+
openvdb::GridPtrVec grids;
149+
grids.push_back(grid);
150+
openvdb::io::File file(outputFile);
151+
file.write(grids);
152+
file.close();
153+
154+
// Cast the raw pointers from the std::vector data
155+
const auto* nano_pts_data = reinterpret_cast<const nanovdb::Vec3f*>(openvdb_points.data());
156+
const auto* nano_tris_data = reinterpret_cast<const nanovdb::Vec3i*>(openvdb_triangles.data());
157+
158+
// Initialize the thrust vectors using the casted pointer ranges
159+
thrust::universal_vector<nanovdb::Vec3f> nanovdb_points(nano_pts_data, nano_pts_data + openvdb_points.size());
160+
thrust::universal_vector<nanovdb::Vec3i> nanovdb_triangles(nano_tris_data, nano_tris_data + openvdb_triangles.size());
161+
162+
// Convert OpenVDB transform to nanovdb::Map
163+
164+
const auto openvdb_mat4 = transform->baseMap()->getAffineMap()->getMat4();
165+
nanovdb::Map map;
166+
map.set(openvdb_mat4, openvdb_mat4.inverse());
167+
168+
mainMeshToGrid<BuildT>(
169+
nanovdb_points.data().get(),
170+
nanovdb_points.size(),
171+
nanovdb_triangles.data().get(),
172+
nanovdb_triangles.size(),
173+
map);
174+
175+
return 0;
176+
}
177+
catch (const std::exception& e) {
178+
std::cerr << "An exception occurred: \"" << e.what() << "\"" << std::endl;
179+
}
180+
return 0;
181+
}

‎nanovdb/nanovdb/examples/ex_mesh_to_grid_cuda/mesh_to_grid_cuda_kernels.cu‎

Lines changed: 155 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,155 @@
1+
// Copyright Contributors to the OpenVDB Project
2+
// SPDX-License-Identifier: Apache-2.0
3+
4+
5+
#include <nanovdb/NanoVDB.h>
6+
7+
#include <nanovdb/tools/cuda/MeshToGrid.cuh>
8+
9+
#if 0
10+
#include <nanovdb/tools/cuda/DilateGrid.cuh>
11+
#include <nanovdb/tools/cuda/PruneGrid.cuh>
12+
#include <nanovdb/util/cuda/Injection.cuh>
13+
14+
template<typename T>
15+
bool bufferCheck(const T* deviceBuffer, const T* hostBuffer, size_t elem_count) {
16+
T* tmpBuffer = new T[elem_count];
17+
cudaCheck(cudaMemcpy(tmpBuffer, deviceBuffer, elem_count * sizeof(T), cudaMemcpyDeviceToHost));
18+
bool same = true;
19+
for (int i=0; same && i< elem_count; ++i) { same = (tmpBuffer[i] == hostBuffer[i]); }
20+
delete [] tmpBuffer;
21+
return same;
22+
}
23+
#endif
24+
25+
template<typename BuildT>
26+
void mainMeshToGrid(
27+
const nanovdb::Vec3f *devicePoints,
28+
const int pointCount,
29+
const nanovdb::Vec3i *deviceTriangles,
30+
const int triangleCount,
31+
const nanovdb::Map map)
32+
{
33+
nanovdb::util::cuda::Timer gpuTimer;
34+
35+
// Initialize mesh-to-grid converter
36+
nanovdb::tools::cuda::MeshToGrid<BuildT> converter( devicePoints, pointCount, deviceTriangles, triangleCount, map );
37+
converter.setVerbose(1);
38+
converter.getHandle();
39+
40+
41+
// --- DIAGNOSTIC CHECK 1: The Modulus & Bounds Test ---
42+
uint64_t pairCount = converter.getPairCount();
43+
if (pairCount > 0) {
44+
std::cout << "\n--- Running GPU Diagnostics ---" << std::endl;
45+
46+
// Allocate host memory and download the pairs
47+
std::vector<typename nanovdb::tools::cuda::MeshToGrid<BuildT>::BoxTrianglePair> hostPairs(pairCount);
48+
cudaMemcpy(hostPairs.data(), converter.getDevicePairs(),
49+
pairCount * sizeof(typename nanovdb::tools::cuda::MeshToGrid<BuildT>::BoxTrianglePair),
50+
cudaMemcpyDeviceToHost);
51+
cudaDeviceSynchronize();
52+
53+
bool passed = true;
54+
for (uint64_t i = 0; i < pairCount; ++i) {
55+
const auto& pair = hostPairs[i];
56+
57+
// 1. Verify Modulo 4096 (Strict Root Tile Alignment)
58+
if (pair.origin[0] % 4096 != 0 ||
59+
pair.origin[1] % 4096 != 0 ||
60+
pair.origin[2] % 4096 != 0) {
61+
std::cerr << "FAIL: Misaligned Root Origin at index " << i
62+
<< " (" << pair.origin[0] << ", " << pair.origin[1] << ", " << pair.origin[2] << ")\n";
63+
passed = false;
64+
break;
65+
}
66+
67+
// 2. Verify Triangle ID bounds
68+
if (pair.triangleID >= triangleCount) {
69+
std::cerr << "FAIL: Out-of-bounds TriangleID " << pair.triangleID << " at index " << i << "\n";
70+
passed = false;
71+
break;
72+
}
73+
}
74+
75+
if (passed) {
76+
std::cout << "SUCCESS: All " << pairCount << " pairs are perfectly 4096-aligned and bounded!" << std::endl;
77+
// Print a sample to visually inspect
78+
std::cout << "Sample Pair [0]: Origin(" << hostPairs[0].origin[0] << ", "
79+
<< hostPairs[0].origin[1] << ", " << hostPairs[0].origin[2] << ") - TriID: "
80+
<< hostPairs[0].triangleID << std::endl;
81+
}
82+
}
83+
84+
85+
86+
#if 0
87+
dilator.setOperation(nanovdb::tools::morphology::NearestNeighbors(nnType));
88+
dilator.setChecksum(nanovdb::CheckMode::Default);
89+
90+
auto handle = dilator.getHandle();
91+
auto dstGrid = handle.template deviceGrid<BuildT>();
92+
93+
// Check for correctness
94+
if (bufferCheck((char*)dstGrid, (char*)indexGridDilated->data(), indexGridDilated->gridSize()))
95+
std::cout << "Result of DilateGrid check out CORRECT against reference" << std::endl;
96+
else
97+
std::cout << "Result of DilateGrid compares INCORRECT against reference" << std::endl;
98+
99+
// Re-run warm-started iterations
100+
dilator.setVerbose(0);
101+
for (int i = 0; i < benchmark_iters; i++) {
102+
gpuTimer.start("Re-running entire dilation after warmstart");
103+
auto dummyHandle = dilator.getHandle();
104+
gpuTimer.stop();
105+
}
106+
107+
uint32_t dstLeafCount = nanovdb::util::cuda::DeviceGridTraits<BuildT>::getTreeData(dstGrid).mNodeCount[0];
108+
nanovdb::cuda::DeviceBuffer dstLeafMaskBuffer;
109+
nanovdb::Mask<3>* dstLeafMasks = nullptr;
110+
if (dstLeafCount) {
111+
dstLeafMaskBuffer = nanovdb::cuda::DeviceBuffer::create( std::size_t(dstLeafCount) * sizeof(nanovdb::Mask<3>), nullptr, false );
112+
dstLeafMasks = static_cast<nanovdb::Mask<3>*>(dstLeafMaskBuffer.deviceData());
113+
if (!dstLeafMasks) throw std::runtime_error("No GPU buffer for dstLeafMask");
114+
}
115+
116+
const unsigned int numThreads = 128;
117+
auto numBlocks = [numThreads] (unsigned int n) {return (n + numThreads - 1) / numThreads;};
118+
gpuTimer.start("Injecting un-dilated topology as a pruning mask");
119+
if (dstLeafCount)
120+
nanovdb::util::cuda::lambdaKernel<<<numBlocks(dstLeafCount), numThreads>>>(dstLeafCount,
121+
nanovdb::util::cuda::InjectGridMaskFunctor<BuildT>(),
122+
deviceGridOriginal, dstGrid, dstLeafMasks );
123+
gpuTimer.stop();
124+
125+
// Initialize pruner
126+
nanovdb::tools::cuda::PruneGrid<BuildT> pruner( dstGrid, dstLeafMasks );
127+
pruner.setChecksum(nanovdb::CheckMode::Default);
128+
pruner.setVerbose(1);
129+
130+
auto prunedHandle = pruner.getHandle();
131+
auto prunedGrid = prunedHandle.template deviceGrid<BuildT>();
132+
133+
// Check for correctness
134+
if (bufferCheck((char*)prunedGrid, (char*)indexGridOriginal->data(), indexGridOriginal->gridSize()))
135+
std::cout << "Result of PruneGrid check out CORRECT against reference" << std::endl;
136+
else
137+
std::cout << "Result of PruneGrid compares INCORRECT against reference" << std::endl;
138+
139+
// Re-run warm-started iterations
140+
pruner.setVerbose(0);
141+
for (int i = 0; i < benchmark_iters; i++) {
142+
gpuTimer.start("Re-running entire pruning after warmstart");
143+
auto dummyHandle = pruner.getHandle();
144+
gpuTimer.stop();
145+
}
146+
#endif
147+
}
148+
149+
template
150+
void mainMeshToGrid<nanovdb::ValueOnIndex>(
151+
const nanovdb::Vec3f *devicePoints,
152+
const int pointCount,
153+
const nanovdb::Vec3i *deviceTriangles,
154+
const int triangleCount,
155+
const nanovdb::Map map);

0 commit comments

Comments
 (0)