Skip to content

Commit 3b2ad8f

Browse files
sifakisclaude
sifakis
and
claude
committed
Handle empty triangle list gracefully in MeshToGrid
Replace throws-on-zero with early returns throughout the pipeline so that getHandle() and getHandleAndUDF() return a valid empty grid when triangleCount == 0. Kernel launches over null sets are explicitly guarded before any buffer pointer is dereferenced. Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com> Signed-off-by: Efty Sifakis <esifakis@nvidia.com>
1 parent 79e9bab commit 3b2ad8f

2 files changed

Lines changed: 74 additions & 58 deletions

File tree

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

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -23,6 +23,14 @@ void mainMeshToGrid(
2323
const nanovdb::Map map,
2424
const openvdb::FloatGrid::Ptr refGrid)
2525
{
26+
// Test topology-only path (getHandle)
27+
{
28+
nanovdb::tools::cuda::MeshToGrid<BuildT> conv(devicePoints, pointCount, deviceTriangles, triangleCount, map);
29+
conv.setVerbose(1);
30+
auto handle = conv.getHandle();
31+
std::cout << "[getHandle] completed, buffer size: " << handle.bufferSize() << " bytes\n";
32+
}
33+
2634
// Initialize mesh-to-grid converter
2735
nanovdb::tools::cuda::MeshToGrid<BuildT> converter( devicePoints, pointCount, deviceTriangles, triangleCount, map );
2836
converter.setVerbose(1);

‎nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh‎

Lines changed: 66 additions & 58 deletions
Original file line numberDiff line numberDiff line change
@@ -6,7 +6,8 @@
66
77
\authors Efty Sifakis
88
9-
\brief Rasterization of triangle mesh into a sparse NanoVDB indexGrid on the device
9+
\brief Rasterization of triangle mesh into a sparse NanoVDB indexGrid on the device.
10+
Optionally an Unsigned Distance Field can be returned in a newly allocated sidecar.
1011
1112
\warning The header file contains cuda device code so be sure
1213
to only include it in .cu files (or other .cuh files)
@@ -50,15 +51,15 @@ struct Triangle {
5051
template <typename BuildT>
5152
class MeshToGrid
5253
{
53-
using PointT = nanovdb::Vec3f;
54+
using PointT = nanovdb::Vec3f;
5455
using TriangleIndexT = nanovdb::Vec3i;
5556
using TriangleT = Triangle;
56-
using GridT = NanoGrid<BuildT>;
57-
using TreeT = NanoTree<BuildT>;
58-
using RootT = NanoRoot<BuildT>;
59-
using UpperT = NanoUpper<BuildT>;
60-
using LowerT = NanoLower<BuildT>;
61-
using LeafT = NanoLeaf<BuildT>;
57+
using GridT = NanoGrid<BuildT>;
58+
using TreeT = NanoTree<BuildT>;
59+
using RootT = NanoRoot<BuildT>;
60+
using UpperT = NanoUpper<BuildT>;
61+
using LowerT = NanoLower<BuildT>;
62+
using LeafT = NanoLeaf<BuildT>;
6263

6364
public:
6465
struct alignas(16) BoxTrianglePair { // sizeof(BoxTrianglePair) = 16B
@@ -259,7 +260,7 @@ MeshToGrid<BuildT>::getHandle(const BufferT &pool)
259260

260261
// Allocate output grid buffer
261262
if (mVerbose==1) mTimer.start("Allocating grid buffer");
262-
auto gridBuffer = mBuilder.getBuffer(pool, mStream);
263+
auto buffer = mBuilder.getBuffer(pool, mStream);
263264
if (mVerbose==1) mTimer.stop();
264265

265266
// Initialize grid/tree/root metadata
@@ -281,8 +282,10 @@ MeshToGrid<BuildT>::getHandle(const BufferT &pool)
281282
if (mVerbose==1) mTimer.start("Rasterizing leaf nodes");
282283
rasterizeLeafNodes();
283284
if (mVerbose==1) mTimer.stop();
284-
mXformedTriangles.clear(mStream);
285-
mBoxTrianglePairsBuffer.clear(mStream);
285+
if (mBoxTrianglePairCount) {
286+
mXformedTriangles.clear(mStream);
287+
mBoxTrianglePairsBuffer.clear(mStream);
288+
}
286289

287290
// Update leaf value offsets (prefix sums of per-leaf active voxel counts)
288291
if (mVerbose==1) mTimer.start("Processing leaf offsets");
@@ -308,17 +311,20 @@ MeshToGrid<BuildT>::getHandle(const BufferT &pool)
308311
if (mVerbose==1) mTimer.start("Pruning empty leaves");
309312
int device = 0; cudaGetDevice(&device);
310313
const uint32_t leafCount = mBuilder.data()->nodeCount[0];
311-
nanovdb::cuda::DeviceBuffer retainMaskBuffer = nanovdb::cuda::DeviceBuffer::create(
312-
uint64_t(leafCount) * sizeof(nanovdb::Mask<3>), nullptr, device, mStream);
313-
cudaCheck(cudaMemsetAsync(retainMaskBuffer.deviceData(), 0xFF,
314-
uint64_t(leafCount) * sizeof(nanovdb::Mask<3>), mStream));
315-
tools::cuda::PruneGrid<BuildT> pruner(
316-
static_cast<const GridT*>(gridBuffer.deviceData()),
317-
static_cast<nanovdb::Mask<3>*>(retainMaskBuffer.deviceData()),
318-
mStream);
319-
auto prunedHandle = pruner.template getHandle<BufferT>(pool);
314+
auto handle = GridHandle<BufferT>(std::move(buffer));
315+
if (leafCount) {
316+
nanovdb::cuda::DeviceBuffer retainMaskBuffer = nanovdb::cuda::DeviceBuffer::create(
317+
uint64_t(leafCount) * sizeof(nanovdb::Mask<3>), nullptr, device, mStream);
318+
cudaCheck(cudaMemsetAsync(retainMaskBuffer.deviceData(), 0xFF,
319+
uint64_t(leafCount) * sizeof(nanovdb::Mask<3>), mStream));
320+
tools::cuda::PruneGrid<BuildT> pruner(
321+
static_cast<const GridT*>(handle.deviceData()),
322+
static_cast<nanovdb::Mask<3>*>(retainMaskBuffer.deviceData()),
323+
mStream);
324+
handle = pruner.template getHandle<BufferT>(pool);
325+
}
320326
if (mVerbose==1) mTimer.stop();
321-
return prunedHandle;
327+
return handle;
322328
} // MeshToGrid<BuildT>::getHandle
323329

324330
//-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------
@@ -347,9 +353,7 @@ struct TransformTrianglesFunctor
347353
template<typename BuildT>
348354
void MeshToGrid<BuildT>::transformTriangles()
349355
{
350-
// TODO: Handle null input case
351-
if (mTriangleCount == 0)
352-
throw std::runtime_error("MeshToGrid currently requires mTriangleCount > 0 (Holistic zero-handling pending).");
356+
if (mTriangleCount == 0) return;
353357

354358
int device = 0;
355359
cudaGetDevice(&device);
@@ -486,9 +490,7 @@ struct ScatterRootTrianglePairsFunctor
486490
template<typename BuildT>
487491
void MeshToGrid<BuildT>::processRootTrianglePairs()
488492
{
489-
// TODO: Handle null input case
490-
if (mTriangleCount == 0)
491-
throw std::runtime_error("MeshToGrid currently requires mTriangleCount > 0 (Holistic zero-handling pending).");
493+
if (mTriangleCount == 0) { mBoxTrianglePairCount = 0; return; }
492494

493495
int device = 0;
494496
cudaGetDevice(&device);
@@ -867,10 +869,6 @@ void MeshToGrid<BuildT>::buildRasterizedRoot()
867869
// Origins are already in coordToKey-sorted order from enumerateRootTiles(),
868870
// so no further sorting or deduplication is required here.
869871
uint32_t tileCount = static_cast<uint32_t>(mUniqueRootTileCount);
870-
std::vector<nanovdb::Coord> hostOrigins(tileCount);
871-
cudaCheck(cudaMemcpy(hostOrigins.data(),
872-
deviceUniqueRootOrigins(),
873-
tileCount * sizeof(nanovdb::Coord), cudaMemcpyDeviceToHost));
874872

875873
// Build the root node on CPU: one tile per unique root origin.
876874
// Only the NanoVDB tile key is set here; child pointers and values are
@@ -880,19 +878,26 @@ void MeshToGrid<BuildT>::buildRasterizedRoot()
880878
auto *rootPtr = static_cast<RootT*>(mBuilder.mProcessedRoot.data());
881879
rootPtr->mTableSize = tileCount;
882880
rootPtr->mBackground = typename RootT::ValueType{};
883-
for (uint32_t t = 0; t < tileCount; ++t)
884-
*rootPtr->tile(t) = typename RootT::DataType::Tile{RootT::CoordToKey(hostOrigins[t])};
885-
886-
mBuilder.mProcessedRoot.deviceUpload(device, mStream, false);
887881

888-
mUniqueRootOriginsBuffer.clear(mStream);
882+
if (tileCount) {
883+
std::vector<nanovdb::Coord> hostOrigins(tileCount);
884+
cudaCheck(cudaMemcpy(hostOrigins.data(),
885+
deviceUniqueRootOrigins(),
886+
tileCount * sizeof(nanovdb::Coord), cudaMemcpyDeviceToHost));
887+
for (uint32_t t = 0; t < tileCount; ++t)
888+
*rootPtr->tile(t) = typename RootT::DataType::Tile{RootT::CoordToKey(hostOrigins[t])};
889+
mBuilder.mProcessedRoot.deviceUpload(device, mStream, false);
890+
mUniqueRootOriginsBuffer.clear(mStream);
891+
}
889892
} // MeshToGrid<BuildT>::buildRasterizedRoot
890893

891894
//-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------
892895

893896
template<typename BuildT>
894897
void MeshToGrid<BuildT>::rasterizeInternalNodes()
895898
{
899+
if (mBoxTrianglePairCount == 0) return;
900+
896901
using RasterizerT = util::rasterization::cuda::RasterizeInternalNodesFunctor<BuildT, BoxTrianglePair>;
897902

898903
auto *dUpperMasks = static_cast<Mask<5>*>(mBuilder.deviceUpperMasks());
@@ -949,9 +954,7 @@ void MeshToGrid<BuildT>::rasterizeLeafNodes()
949954
template<typename BuildT>
950955
void MeshToGrid<BuildT>::processLeafTrianglePairs()
951956
{
952-
// TODO: Handle null input case
953-
if (mTriangleCount == 0)
954-
throw std::runtime_error("MeshToGrid currently requires mTriangleCount > 0 (Holistic zero-handling pending).");
957+
if (mBoxTrianglePairCount == 0) return;
955958

956959
int device = 0;
957960
cudaGetDevice(&device);
@@ -1132,7 +1135,7 @@ MeshToGrid<BuildT>::getHandleAndUDF(const GridBufferT& gridPool, const SidecarBu
11321135
mBuilder.data()->nodeCount[0]);
11331136

11341137
if (mVerbose==1) mTimer.start("Allocating grid buffer");
1135-
auto gridBuffer = mBuilder.getBuffer(gridPool, mStream);
1138+
auto buffer = mBuilder.getBuffer(gridPool, mStream);
11361139
if (mVerbose==1) mTimer.stop();
11371140

11381141
if (mVerbose==1) mTimer.start("Processing grid/tree/root");
@@ -1167,15 +1170,18 @@ MeshToGrid<BuildT>::getHandleAndUDF(const GridBufferT& gridPool, const SidecarBu
11671170
if (mVerbose==1) mTimer.start("Pruning empty leaves");
11681171
int device = 0; cudaGetDevice(&device);
11691172
const uint32_t leafCount = mBuilder.data()->nodeCount[0];
1170-
nanovdb::cuda::DeviceBuffer retainMaskBuffer = nanovdb::cuda::DeviceBuffer::create(
1171-
uint64_t(leafCount) * sizeof(nanovdb::Mask<3>), nullptr, device, mStream);
1172-
cudaCheck(cudaMemsetAsync(retainMaskBuffer.deviceData(), 0xFF,
1173-
uint64_t(leafCount) * sizeof(nanovdb::Mask<3>), mStream));
1174-
tools::cuda::PruneGrid<BuildT> pruner(
1175-
static_cast<const GridT*>(gridBuffer.deviceData()),
1176-
static_cast<nanovdb::Mask<3>*>(retainMaskBuffer.deviceData()),
1177-
mStream);
1178-
auto handle = pruner.template getHandle<GridBufferT>(gridPool);
1173+
auto handle = GridHandle<GridBufferT>(std::move(buffer));
1174+
if (leafCount) {
1175+
nanovdb::cuda::DeviceBuffer retainMaskBuffer = nanovdb::cuda::DeviceBuffer::create(
1176+
uint64_t(leafCount) * sizeof(nanovdb::Mask<3>), nullptr, device, mStream);
1177+
cudaCheck(cudaMemsetAsync(retainMaskBuffer.deviceData(), 0xFF,
1178+
uint64_t(leafCount) * sizeof(nanovdb::Mask<3>), mStream));
1179+
tools::cuda::PruneGrid<BuildT> pruner(
1180+
static_cast<const GridT*>(handle.deviceData()),
1181+
static_cast<nanovdb::Mask<3>*>(retainMaskBuffer.deviceData()),
1182+
mStream);
1183+
handle = pruner.template getHandle<GridBufferT>(gridPool);
1184+
}
11791185
if (mVerbose==1) mTimer.stop();
11801186

11811187
// ---- UDF sidecar ----
@@ -1198,16 +1204,18 @@ MeshToGrid<BuildT>::getHandleAndUDF(const GridBufferT& gridPool, const SidecarBu
11981204
if (mVerbose==1) mTimer.stop();
11991205

12001206
if (mVerbose==1) mTimer.start("Computing UDF via leaf/triangle pairs");
1201-
using UDFFunctorT = util::rasterization::cuda::ComputeUDFFunctor<BuildT, BoxTrianglePair, Triangle>;
1202-
util::cuda::operatorKernelInstance<UDFFunctorT>
1203-
<<<mBoxTrianglePairCount, UDFFunctorT::MaxThreadsPerBlock, 0, mStream>>>(
1204-
UDFFunctorT{ deviceBoxTrianglePairs(), deviceXformedTriangles(),
1205-
handle.template deviceGrid<BuildT>(), dSidecar,
1206-
mBandWidth * mBandWidth });
1207-
cudaCheckError();
1207+
if (mBoxTrianglePairCount) {
1208+
using UDFFunctorT = util::rasterization::cuda::ComputeUDFFunctor<BuildT, BoxTrianglePair, Triangle>;
1209+
util::cuda::operatorKernelInstance<UDFFunctorT>
1210+
<<<mBoxTrianglePairCount, UDFFunctorT::MaxThreadsPerBlock, 0, mStream>>>(
1211+
UDFFunctorT{ deviceBoxTrianglePairs(), deviceXformedTriangles(),
1212+
handle.template deviceGrid<BuildT>(), dSidecar,
1213+
mBandWidth * mBandWidth });
1214+
cudaCheckError();
1215+
mXformedTriangles.clear(mStream);
1216+
mBoxTrianglePairsBuffer.clear(mStream);
1217+
}
12081218
if (mVerbose==1) mTimer.stop();
1209-
mXformedTriangles.clear(mStream);
1210-
mBoxTrianglePairsBuffer.clear(mStream);
12111219

12121220
if (mVerbose==1) mTimer.start("Finalizing UDF sidecar (sqrt + clamp)");
12131221
util::cuda::lambdaKernel<<<numBlocks(activeVoxelCount + 1), mNumThreads, 0, mStream>>>(

0 commit comments

Comments
 (0)