From 905a8329058935a93b77ab525a7bfc771351f2b7 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Thu, 10 Sep 2026 06:27:21 +0000 Subject: [PATCH 1/2] NanoVDB: type the Mask scratch buffers as cuda::Buffer> (CUDA) Mask became trivially copyable in #2310, so the scratch that exists only to hold arrays of masks no longer needs to be a byte buffer reinterpreted at every use. TopologyBuilder's upper and lower mask members become Buffer> and Buffer>, and deviceUpperMasks/deviceLowerMasks return typed pointers, which the seven morphology functors and MeshToGrid now take directly - the only cast left is the array-shape view of the flat lower masks as one Mask<5>::SIZE row per upper node. MeshToGrid's two retain-mask buffers and its per-pass hit-mask buffer, the dilation example's leaf masks, and the dilation test's mask buffer (previously on the deprecated dual DeviceBuffer) are Buffer>. Allocation sizes are element counts and memsets use size_bytes(). Closes #2312. Signed-off-by: Mark Harris --- .../dilate_nanovdb_cuda_kernels.cu | 6 +- nanovdb/nanovdb/tools/cuda/CoarsenGrid.cuh | 2 +- nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh | 28 ++++----- nanovdb/nanovdb/tools/cuda/PruneGrid.cuh | 2 +- nanovdb/nanovdb/tools/cuda/RefineGrid.cuh | 2 +- .../nanovdb/tools/cuda/TopologyBuilder.cuh | 31 +++++----- nanovdb/nanovdb/unittest/TestNanoVDB.cu | 7 ++- nanovdb/nanovdb/util/cuda/Morphology.cuh | 57 +++++-------------- pendingchanges/nanovdb.txt | 1 + 9 files changed, 59 insertions(+), 77 deletions(-) diff --git a/nanovdb/nanovdb/examples/ex_dilate_nanovdb_cuda/dilate_nanovdb_cuda_kernels.cu b/nanovdb/nanovdb/examples/ex_dilate_nanovdb_cuda/dilate_nanovdb_cuda_kernels.cu index f9cf3bb282..6b024bcc18 100644 --- a/nanovdb/nanovdb/examples/ex_dilate_nanovdb_cuda/dilate_nanovdb_cuda_kernels.cu +++ b/nanovdb/nanovdb/examples/ex_dilate_nanovdb_cuda/dilate_nanovdb_cuda_kernels.cu @@ -51,11 +51,11 @@ void mainDilateGrid( } uint32_t dstLeafCount = nanovdb::util::cuda::DeviceGridTraits::getTreeData(dstGrid).mNodeCount[0]; - nanovdb::cuda::Buffer dstLeafMaskBuffer; + nanovdb::cuda::Buffer> dstLeafMaskBuffer; nanovdb::Mask<3>* dstLeafMasks = nullptr; if (dstLeafCount) { - dstLeafMaskBuffer = nanovdb::cuda::Buffer(cudaStream_t(0), std::size_t(dstLeafCount) * sizeof(nanovdb::Mask<3>), nanovdb::cuda::noInit); - dstLeafMasks = reinterpret_cast*>(dstLeafMaskBuffer.data()); + dstLeafMaskBuffer = nanovdb::cuda::Buffer>(cudaStream_t(0), std::size_t(dstLeafCount), nanovdb::cuda::noInit); + dstLeafMasks = dstLeafMaskBuffer.data(); if (!dstLeafMasks) throw std::runtime_error("No GPU buffer for dstLeafMask"); } diff --git a/nanovdb/nanovdb/tools/cuda/CoarsenGrid.cuh b/nanovdb/nanovdb/tools/cuda/CoarsenGrid.cuh index cca0ac00ac..5fbd7c7de6 100644 --- a/nanovdb/nanovdb/tools/cuda/CoarsenGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/CoarsenGrid.cuh @@ -215,7 +215,7 @@ void CoarsenGrid::coarsenInternalNodes() if (auto srcLeafCount = mSrcTreeData.mNodeCount[0]) { // Unless it's an empty grid util::cuda::lambdaKernel<<>>( srcLeafCount, util::morphology::cuda::CoarsenInternalNodesFunctor(), - mDeviceSrcGrid, mBuilder.deviceProcessedRoot(), mBuilder.mUpperMasks.data(), mBuilder.mLowerMasks.data() ); + mDeviceSrcGrid, mBuilder.deviceProcessedRoot(), mBuilder.deviceUpperMasks(), mBuilder.deviceLowerMasks() ); } }// CoarsenGrid::coarsenInternalNodes diff --git a/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh b/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh index 4962392b7b..2edc234f59 100644 --- a/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh @@ -18,6 +18,8 @@ #include #include +#include // for std::byte +#include // for std::runtime_error #include #include @@ -65,6 +67,7 @@ class MeshToGrid using PointT = nanovdb::Vec3f; using ScratchT = nanovdb::cuda::Buffer>; + using MaskBufT = nanovdb::cuda::Buffer, nanovdb::cuda::ResourceRef>; nanovdb::cuda::ResourceRef ref() { return mBuilder.ref(); } @@ -315,12 +318,12 @@ GridHandle MeshToGrid::getHandle(const BufferT &buff const uint32_t leafCount = mBuilder.data()->nodeCount[0]; auto handle = GridHandle(std::move(gridBuffer)); if (leafCount) { - ScratchT retainMaskBuffer = ScratchT(mStream, this->ref(), uint64_t(leafCount) * sizeof(nanovdb::Mask<3>), nanovdb::cuda::noInit); - cudaCheck(cudaMemsetAsync(retainMaskBuffer.data(), 0xFF, - uint64_t(leafCount) * sizeof(nanovdb::Mask<3>), mStream)); + MaskBufT retainMaskBuffer = MaskBufT(mStream, this->ref(), leafCount, nanovdb::cuda::noInit); + if (retainMaskBuffer.data() == nullptr) throw std::runtime_error("Failed to allocate retain mask buffer"); + cudaCheck(cudaMemsetAsync(retainMaskBuffer.data(), 0xFF, retainMaskBuffer.size_bytes(), mStream)); tools::cuda::PruneGrid pruner( static_cast(handle.deviceData()), - reinterpret_cast*>(retainMaskBuffer.data()), + retainMaskBuffer.data(), mStream); handle = pruner.template getHandle(buffer); } @@ -869,8 +872,8 @@ void MeshToGrid::rasterizeInternalNodes() using RasterizerT = util::rasterization::cuda::RasterizeInternalNodesFunctor; - auto *dUpperMasks = static_cast*>(mBuilder.deviceUpperMasks()); - auto *dLowerMasks = static_cast(*)[Mask<5>::SIZE]>(mBuilder.deviceLowerMasks()); + auto *dUpperMasks = mBuilder.deviceUpperMasks(); + auto *dLowerMasks = mBuilder.deviceLowerMasks(); util::cuda::lambdaKernel<<>>( mBoxTrianglePairCount, @@ -935,12 +938,11 @@ void MeshToGrid::processLeafTrianglePairs() for (int pass = 0; pass < 3; ++pass) { // Allocate Mask<3> buffer for the CTA hit results - // Size: mBoxTrianglePairCount * sizeof(nanovdb::Mask<3>) - ScratchT maskBuffer = ScratchT(mStream, this->ref(), mBoxTrianglePairCount * sizeof(nanovdb::Mask<3>), nanovdb::cuda::noInit); + MaskBufT maskBuffer = MaskBufT(mStream, this->ref(), mBoxTrianglePairCount, nanovdb::cuda::noInit); if (maskBuffer.data() == nullptr) { throw std::runtime_error("Failed to allocate mask buffer for subdivision pass"); } - auto* dMasks = reinterpret_cast*>(maskBuffer.data()); + auto* dMasks = maskBuffer.data(); // Allocate Counts buffer for Prefix Sum // Size: mBoxTrianglePairCount * sizeof(uint64_t) @@ -1125,12 +1127,12 @@ MeshToGrid::getHandleAndUDF(const GridBufferT& buffer, const const uint32_t leafCount = mBuilder.data()->nodeCount[0]; auto handle = GridHandle(std::move(gridBuffer)); if (leafCount) { - ScratchT retainMaskBuffer = ScratchT(mStream, this->ref(), uint64_t(leafCount) * sizeof(nanovdb::Mask<3>), nanovdb::cuda::noInit); - cudaCheck(cudaMemsetAsync(retainMaskBuffer.data(), 0xFF, - uint64_t(leafCount) * sizeof(nanovdb::Mask<3>), mStream)); + MaskBufT retainMaskBuffer = MaskBufT(mStream, this->ref(), leafCount, nanovdb::cuda::noInit); + if (retainMaskBuffer.data() == nullptr) throw std::runtime_error("Failed to allocate retain mask buffer"); + cudaCheck(cudaMemsetAsync(retainMaskBuffer.data(), 0xFF, retainMaskBuffer.size_bytes(), mStream)); tools::cuda::PruneGrid pruner( static_cast(handle.deviceData()), - reinterpret_cast*>(retainMaskBuffer.data()), + retainMaskBuffer.data(), mStream); handle = pruner.template getHandle(buffer); } diff --git a/nanovdb/nanovdb/tools/cuda/PruneGrid.cuh b/nanovdb/nanovdb/tools/cuda/PruneGrid.cuh index 7c2f0c9a8d..8a24dada0a 100644 --- a/nanovdb/nanovdb/tools/cuda/PruneGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/PruneGrid.cuh @@ -218,7 +218,7 @@ void PruneGrid::pruneInternalNodes() if (auto srcLeafCount = mSrcTreeData.mNodeCount[0]) { // Unless it's an empty grid util::cuda::lambdaKernel<<>>( srcLeafCount, util::morphology::cuda::PruneInternalNodesFunctor(), - mDeviceSrcGrid, mBuilder.deviceProcessedRoot(), mDeviceSrcLeafMask, mBuilder.mUpperMasks.data(), mBuilder.mLowerMasks.data() ); + mDeviceSrcGrid, mBuilder.deviceProcessedRoot(), mDeviceSrcLeafMask, mBuilder.deviceUpperMasks(), mBuilder.deviceLowerMasks() ); } }// PruneGrid::pruneInternalNodes diff --git a/nanovdb/nanovdb/tools/cuda/RefineGrid.cuh b/nanovdb/nanovdb/tools/cuda/RefineGrid.cuh index 0140e10ebe..a38629ee28 100644 --- a/nanovdb/nanovdb/tools/cuda/RefineGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/RefineGrid.cuh @@ -230,7 +230,7 @@ void RefineGrid::refineInternalNodes() if (auto srcLeafCount = mSrcTreeData.mNodeCount[0]) { // Unless it's an empty grid util::cuda::lambdaKernel<<>>( srcLeafCount, util::morphology::cuda::RefineInternalNodesFunctor(), - mDeviceSrcGrid, mBuilder.deviceProcessedRoot(), mBuilder.mUpperMasks.data(), mBuilder.mLowerMasks.data() ); + mDeviceSrcGrid, mBuilder.deviceProcessedRoot(), mBuilder.deviceUpperMasks(), mBuilder.deviceLowerMasks() ); } }// RefineGrid::refineInternalNodes diff --git a/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh b/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh index 4f82f1e0f5..e8040e8905 100644 --- a/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh +++ b/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh @@ -70,6 +70,8 @@ class TopologyBuilder /// Buffer rather than the dual DeviceBuffer, whose host pointer and /// per-device array they would leave unused. using ScratchT = nanovdb::cuda::Buffer>; + using UpperMaskBufT = nanovdb::cuda::Buffer, nanovdb::cuda::ResourceRef>; + using LowerMaskBufT = nanovdb::cuda::Buffer, nanovdb::cuda::ResourceRef>; using HostStagingT = nanovdb::cuda::Buffer; public: @@ -114,8 +116,8 @@ public: HostStagingT mHostRoot; // host staging for the processed root (pinned, so the upload is asynchronous) ScratchT mDeviceRoot; // device copy, made by uploadProcessedRoot - ScratchT mUpperMasks; - ScratchT mLowerMasks; + UpperMaskBufT mUpperMasks; + LowerMaskBufT mLowerMasks; ScratchT mUpperOffsets; ScratchT mLowerOffsets; ScratchT mLeafOffsets; @@ -155,8 +157,11 @@ public: mDeviceData = ScratchT(stream, nanovdb::cuda::ResourceRef(*mResource), sizeof(Data), nanovdb::cuda::noInit); cudaCheck(cudaMemcpyAsync(mDeviceData.data(), &mHostData, sizeof(Data), cudaMemcpyHostToDevice, stream)); } - void* deviceUpperMasks() { return mUpperMasks.data(); } - void* deviceLowerMasks() { return mLowerMasks.data(); } + Mask<5>* deviceUpperMasks() { return mUpperMasks.data(); } + /// @brief The densified lower masks: one row of Mask<5>::SIZE Mask<4> per upper node, + /// indexed [upper node][lower node offset]. The row shape is fixed here, beside the + /// allocation that defines it, so consumers never re-derive the stride. + Mask<4> (*deviceLowerMasks())[Mask<5>::SIZE] { return reinterpret_cast(*)[Mask<5>::SIZE]>(mLowerMasks.data()); } /// @brief A borrowing reference to the builder's resource, for consumers /// allocating sibling scratch from the same instance. nanovdb::cuda::ResourceRef ref() { return nanovdb::cuda::ResourceRef(*mResource); } @@ -196,17 +201,17 @@ void TopologyBuilder::allocateInternalMaskBuffers(cudaStream_ { if (hostProcessedRoot()->tileCount() == 0) return; // Processing empty grid(s); nothing to allocate - // Allocate (and zero-fill) buffers large enough to hold: - // (a) The serialized masks of all upper nodes, for all tiles in the updated root node, and - // (b) The serialized masks of all densified lower nodes, as if every upper node had a full set of 32^3 lower children - uint64_t upperSize = hostProcessedRoot()->tileCount() * sizeof(Mask<5>); - uint64_t lowerSize = hostProcessedRoot()->tileCount() * Mask<5>::SIZE * sizeof(Mask<4>); - mUpperMasks = ScratchT(stream, *mResource, upperSize, nanovdb::cuda::noInit); + // Allocate (and zero-fill) the mask arrays: + // (a) one Mask<5> per tile of the updated root node, and + // (b) Mask<5>::SIZE Mask<4> per tile, as if every upper node had a full set of 32^3 lower children + const uint64_t upperMaskCount = hostProcessedRoot()->tileCount(); + const uint64_t lowerMaskCount = upperMaskCount * Mask<5>::SIZE; + mUpperMasks = UpperMaskBufT(stream, *mResource, upperMaskCount, nanovdb::cuda::noInit); if (mUpperMasks.data() == nullptr) throw std::runtime_error("Failed to allocate upper mask buffer on device"); - cudaCheck(cudaMemsetAsync(mUpperMasks.data(), 0, upperSize, stream)); - mLowerMasks = ScratchT(stream, *mResource, lowerSize, nanovdb::cuda::noInit); + cudaCheck(cudaMemsetAsync(mUpperMasks.data(), 0, mUpperMasks.size_bytes(), stream)); + mLowerMasks = LowerMaskBufT(stream, *mResource, lowerMaskCount, nanovdb::cuda::noInit); if (mLowerMasks.data() == nullptr) throw std::runtime_error("Failed to allocate lower mask buffer on device"); - cudaCheck(cudaMemsetAsync(mLowerMasks.data(), 0, lowerSize, stream)); + cudaCheck(cudaMemsetAsync(mLowerMasks.data(), 0, mLowerMasks.size_bytes(), stream)); }// TopologyBuilder::allocateInternalMaskBuffers //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- diff --git a/nanovdb/nanovdb/unittest/TestNanoVDB.cu b/nanovdb/nanovdb/unittest/TestNanoVDB.cu index 1331c95974..694b8ebc72 100644 --- a/nanovdb/nanovdb/unittest/TestNanoVDB.cu +++ b/nanovdb/nanovdb/unittest/TestNanoVDB.cu @@ -28,6 +28,7 @@ #include #include #include +#include #include #include #include @@ -3765,9 +3766,9 @@ TEST(TestNanoVDBCUDA, DilateInjectPrune_ValueOnIndex) EXPECT_EQ(dilatedTreeData.mVoxelCount, 73); // Create a prune mask (set bits correspond to retained voxels) from the occupancy of the original grid - auto maskBuffer = nanovdb::cuda::DeviceBuffer::create( dilatedTreeData.mNodeCount[0] * sizeof(nanovdb::Mask<3>), nullptr, false); - EXPECT_TRUE(maskBuffer.deviceData()); - auto leafMasks = static_cast*>(maskBuffer.deviceData()); + nanovdb::cuda::Buffer> maskBuffer(cudaStream_t(0), dilatedTreeData.mNodeCount[0], nanovdb::cuda::noInit); + EXPECT_TRUE(maskBuffer.data()); + auto leafMasks = maskBuffer.data(); constexpr unsigned int num_threads = 128; unsigned int num_blocks = (static_cast(dilatedTreeData.mNodeCount[0]) + num_threads - 1) / num_threads; nanovdb::util::cuda::lambdaKernel<<>>(dilatedTreeData.mNodeCount[0], diff --git a/nanovdb/nanovdb/util/cuda/Morphology.cuh b/nanovdb/nanovdb/util/cuda/Morphology.cuh index 5b075f4838..68e47ef6b6 100644 --- a/nanovdb/nanovdb/util/cuda/Morphology.cuh +++ b/nanovdb/nanovdb/util/cuda/Morphology.cuh @@ -12,6 +12,7 @@ */ #include +#include // for std::size_t #include @@ -38,8 +39,8 @@ struct DilateInternalNodesFunctor operator()( const NanoGrid *srcGrid, const NanoRoot *dilatedRoot, - void *upperMasks_, - void *lowerMasks_) + Mask<5> *upperMasks, + Mask<4> (*lowerMasks)[Mask<5>::SIZE]) { int tID = threadIdx.x; int lowerID = blockIdx.x; @@ -47,10 +48,6 @@ struct DilateInternalNodesFunctor int threadInWarpID = threadIdx.x & 0x1f; int warpID = threadIdx.x >> 5; - using UpperMaskArrayT = Mask<5>*; - using LowerMaskArrayT = Mask<4>(*)[Mask<5>::SIZE]; - auto upperMasks = static_cast(upperMasks_); - auto lowerMasks = static_cast(lowerMasks_); using LowerMaskT = Mask<4>; using LowerMaskStencilT = LowerMaskT (&)[3][3][3]; @@ -329,16 +326,12 @@ struct MergeInternalNodesFunctor operator()( const NanoGrid *srcGrid, const NanoRoot *mergedRoot, - void *upperMasks_, - void *lowerMasks_) + Mask<5> *upperMasks, + Mask<4> (*lowerMasks)[Mask<5>::SIZE]) { int tID = threadIdx.x; int lowerID = blockIdx.x; - using UpperMaskArrayT = Mask<5>*; - using LowerMaskArrayT = Mask<4>(*)[Mask<5>::SIZE]; - auto upperMasks = static_cast(upperMasks_); - auto lowerMasks = static_cast(lowerMasks_); using LowerMaskT = Mask<4>; const auto& srcTree = srcGrid->tree(); @@ -371,13 +364,9 @@ struct PruneInternalNodesFunctor const NanoGrid* srcGrid, const NanoRoot* prunedRoot, const Mask<3>* srcLeafMask, - void *upperMasks_, - void *lowerMasks_) + Mask<5> *upperMasks, + Mask<4> (*lowerMasks)[Mask<5>::SIZE]) { - using UpperMaskArrayT = Mask<5>*; - using LowerMaskArrayT = Mask<4>(*)[Mask<5>::SIZE]; - auto upperMasks = static_cast(upperMasks_); - auto lowerMasks = static_cast(lowerMasks_); const auto& srcLeaf = srcGrid->tree().template getFirstNode<0>()[srcLeafID]; const auto& leafMask = srcLeafMask[srcLeafID]; @@ -408,13 +397,9 @@ struct RefineInternalNodesFunctor size_t srcLeafID, const NanoGrid* srcGrid, const NanoRoot* prunedRoot, - void *upperMasks_, - void *lowerMasks_) + Mask<5> *upperMasks, + Mask<4> (*lowerMasks)[Mask<5>::SIZE]) { - using UpperMaskArrayT = Mask<5>*; - using LowerMaskArrayT = Mask<4>(*)[Mask<5>::SIZE]; - auto upperMasks = static_cast(upperMasks_); - auto lowerMasks = static_cast(lowerMasks_); const auto& srcLeaf = srcGrid->tree().template getFirstNode<0>()[srcLeafID]; uint64_t octantPresent[2][2][2] = {}; @@ -454,13 +439,9 @@ struct CoarsenInternalNodesFunctor size_t srcLeafID, const NanoGrid* srcGrid, const NanoRoot* prunedRoot, - void *upperMasks_, - void *lowerMasks_) + Mask<5> *upperMasks, + Mask<4> (*lowerMasks)[Mask<5>::SIZE]) { - using UpperMaskArrayT = Mask<5>*; - using LowerMaskArrayT = Mask<4>(*)[Mask<5>::SIZE]; - auto upperMasks = static_cast(upperMasks_); - auto lowerMasks = static_cast(lowerMasks_); const auto& srcLeaf = srcGrid->tree().template getFirstNode<0>()[srcLeafID]; if (!srcLeaf.valueMask().isOff()) { // Gratuitous check; leaf should have at least one active voxel @@ -490,8 +471,8 @@ struct EnumerateNodesFunctor void __device__ operator()( - const void *upperMasks_, - const void *lowerMasks_, + const Mask<5> *upperMasks, + const Mask<4> (*lowerMasks)[Mask<5>::SIZE], uint32_t (*lowerCounts)[Mask<5>::SIZE], uint32_t (*leafCounts)[Mask<5>::SIZE] ) { @@ -500,10 +481,6 @@ struct EnumerateNodesFunctor int threadInWarpID = threadIdx.x & 0x1f; int warpID = threadIdx.x >> 5; - using UpperMaskArrayT = const Mask<5>*; - using LowerMaskArrayT = const Mask<4>(*)[Mask<5>::SIZE]; - auto upperMasks = static_cast(upperMasks_); - auto lowerMasks = static_cast(lowerMasks_); using WarpReduce = cub::WarpReduce; __shared__ typename WarpReduce::TempStorage temp_storage[WarpsPerBlock]; @@ -534,8 +511,8 @@ struct ProcessLowerNodesFunctor void __device__ operator()( - const void *upperMasks_, - const void *lowerMasks_, + const Mask<5> *upperMasks, + const Mask<4> (*lowerMasks)[Mask<5>::SIZE], const uint32_t *upperOffsets, const uint32_t (*lowerOffsets)[Mask<5>::SIZE], const uint32_t (*leafOffsets)[Mask<5>::SIZE], @@ -549,10 +526,6 @@ struct ProcessLowerNodesFunctor int threadInWarpID = threadIdx.x & 0x1f; int warpID = threadIdx.x >> 5; - using UpperMaskArrayT = const Mask<5>*; - using LowerMaskArrayT = const Mask<4>(*)[Mask<5>::SIZE]; - auto upperMasks = static_cast(upperMasks_); - auto lowerMasks = static_cast(lowerMasks_); using WarpScan = cub::WarpScan; __shared__ typename WarpScan::TempStorage temp_storage[WarpsPerBlock]; diff --git a/pendingchanges/nanovdb.txt b/pendingchanges/nanovdb.txt index fbd24e63ae..56845b81b7 100644 --- a/pendingchanges/nanovdb.txt +++ b/pendingchanges/nanovdb.txt @@ -16,6 +16,7 @@ NanoVDB: - nanovdb::cuda::createNodeManager gained a single-space overload: passing a memory resource instead of a pool buffer returns a NodeManagerHandle over cuda::Buffer> whose storage (and size scratch) allocates through that resource, which must outlive the handle. tools::cuda::GridStats now builds its temporary NodeManager through its injected resource, so every device allocation it makes honors ResourceT. NodeManagerHandle supports single-space buffers (deviceMgr maps onto the buffer, host accessors are compile-time errors) and no longer requires a default-constructible buffer type. - GridHandle::reset, NodeManagerHandle::reset and tools::VoxelBlockManager::reset now release storage through the buffer's destroy() when it provides one, and cuda::Buffer::clear/cuda::BufferView::clear are deprecated in favor of destroy() (BufferView::destroy detaches the non-owning view). The legacy HostBuffer/DeviceBuffer clear() methods are unaffected. - The bug-fix to the nanovdb::ReadAccessor (see below) improves random-access performance in some use-cases (especially on the CPU). + - The device-side mask scratch of tools::cuda::TopologyBuilder and tools::cuda::MeshToGrid, the dilation example and its test are now typed cuda::Buffer> instead of byte buffers reinterpreted as masks (possible since Mask became trivially copyable); TopologyBuilder::deviceUpperMasks/deviceLowerMasks return typed pointers and the morphology functors take them. - tools::cuda::DistributedPointsToGrid gained a defaulted ResourceT template parameter and constructor overloads taking one memory-resource instance per device in the mesh, indexed by device id: the CUB scratch pools and the per-device sort scratch route through the injected instances (memory pools are per-device, hence one instance each). The sort scratch is now a stream-ordered cuda::Buffer, replacing a synchronous cudaFree that stalled the device mid-pipeline, and the shared managed metadata and pinned merge intervals are owned by cuda::Buffer members over ManagedResource and PinnedResource instead of hand-paired allocate/free calls. The grid buffer can be a managed single-space handle: getHandle accepts cuda::Buffer alongside the legacy default. The multi-GPU tests and examples build managed cuda::Buffer handles now. Deprecations: From 6b3ce0b9d919ae69f09d20fc50c59e02fb92d29e Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Thu, 10 Sep 2026 22:34:03 +0000 Subject: [PATCH 2/2] NanoVDB: type the remaining scratch buffers in TopologyBuilder, MeshToGrid and SignedFloodFill (CUDA) The same pattern as the mask buffers, applied to every other scratch allocation that holds a single element type for its whole life: TopologyBuilder's upper/lower/leaf/voxel offsets, lower/leaf parents, the three count arrays of its enumerate pass and its device Data staging; MeshToGrid's transformed triangles, box-triangle pairs, unique root origins, and its nine local key/count/offset/pair buffers; SignedFloodFill's root-child array. Allocations are sized in elements, the reinterpret_casts at every use are gone, and the accessors return typed pointers. The lower and leaf offsets, which the node-processing kernels read as one row of Mask<5>::SIZE per upper node, expose that row shape from accessors beside the allocation, as deviceLowerMasks does. TopologyBuilder's root staging stays a byte buffer (its size is RootT::memUsage(tileCount), not an element count), which is now the only reinterpretation its alignment static_assert guards. Addresses the sweep portion of #2327. Signed-off-by: Mark Harris --- nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh | 80 ++++++------- .../nanovdb/tools/cuda/SignedFloodFill.cuh | 4 +- .../nanovdb/tools/cuda/TopologyBuilder.cuh | 107 ++++++++++-------- pendingchanges/nanovdb.txt | 2 +- 4 files changed, 103 insertions(+), 90 deletions(-) diff --git a/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh b/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh index 2edc234f59..b5ffc135fa 100644 --- a/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh @@ -18,7 +18,9 @@ #include #include -#include // for std::byte +#include // for std::size_t +#include // for std::string +#include // for std::vector #include // for std::runtime_error #include @@ -64,10 +66,12 @@ class MeshToGrid { static_assert(nanovdb::cuda::is_async_resource::value, "MeshToGrid allocates stream-ordered scratch and requires an AsyncResource"); + static_assert(ResourceT::DEFAULT_ALIGNMENT >= alignof(MeshToGridBoxTrianglePair), + "MeshToGrid's box-triangle pairs are 16-byte aligned and require allocations at least that aligned"); using PointT = nanovdb::Vec3f; - using ScratchT = nanovdb::cuda::Buffer>; - using MaskBufT = nanovdb::cuda::Buffer, nanovdb::cuda::ResourceRef>; + template + using BufT = nanovdb::cuda::Buffer>; nanovdb::cuda::ResourceRef ref() { return mBuilder.ref(); } @@ -188,16 +192,16 @@ private: const uint32_t mTriangleCount; const nanovdb::Map mMap; - ScratchT mXformedTriangles; - ScratchT mBoxTrianglePairsBuffer; + BufT mXformedTriangles; + BufT mBoxTrianglePairsBuffer; uint64_t mBoxTrianglePairCount{0}; - ScratchT mUniqueRootOriginsBuffer; + BufT mUniqueRootOriginsBuffer; uint64_t mUniqueRootTileCount{0}; - auto deviceXformedTriangles() { return reinterpret_cast(mXformedTriangles.data()); } - auto deviceBoxTrianglePairs() { return reinterpret_cast(mBoxTrianglePairsBuffer.data()); } - auto deviceUniqueRootOrigins() { return reinterpret_cast(mUniqueRootOriginsBuffer.data()); } - auto deviceUniqueRootOrigins() const { return reinterpret_cast(mUniqueRootOriginsBuffer.data()); } + auto deviceXformedTriangles() { return mXformedTriangles.data(); } + auto deviceBoxTrianglePairs() { return mBoxTrianglePairsBuffer.data(); } + auto deviceUniqueRootOrigins() { return mUniqueRootOriginsBuffer.data(); } + auto deviceUniqueRootOrigins() const { return mUniqueRootOriginsBuffer.data(); } nanovdb::cuda::TempPool mTempDevicePool; }; // tools::cuda::MeshToGrid @@ -318,7 +322,7 @@ GridHandle MeshToGrid::getHandle(const BufferT &buff const uint32_t leafCount = mBuilder.data()->nodeCount[0]; auto handle = GridHandle(std::move(gridBuffer)); if (leafCount) { - MaskBufT retainMaskBuffer = MaskBufT(mStream, this->ref(), leafCount, nanovdb::cuda::noInit); + BufT> retainMaskBuffer = BufT>(mStream, this->ref(), leafCount, nanovdb::cuda::noInit); if (retainMaskBuffer.data() == nullptr) throw std::runtime_error("Failed to allocate retain mask buffer"); cudaCheck(cudaMemsetAsync(retainMaskBuffer.data(), 0xFF, retainMaskBuffer.size_bytes(), mStream)); tools::cuda::PruneGrid pruner( @@ -362,7 +366,7 @@ void MeshToGrid::transformTriangles() int device = 0; cudaGetDevice(&device); - mXformedTriangles = ScratchT(mStream, this->ref(), mTriangleCount*sizeof(TriangleT), nanovdb::cuda::noInit); + mXformedTriangles = BufT(mStream, this->ref(), mTriangleCount, nanovdb::cuda::noInit); if (mXformedTriangles.data() == nullptr) throw std::runtime_error("Failed to allocate transofmed upper mask buffer on device"); util::cuda::lambdaKernel<<>>( @@ -500,15 +504,14 @@ void MeshToGrid::processRootTrianglePairs() // Pass 1: Count intersecting root boxes per triangle - ScratchT - rootBoxCounts = ScratchT(mStream, this->ref(), mTriangleCount * sizeof(uint64_t), nanovdb::cuda::noInit); + BufT rootBoxCounts = BufT(mStream, this->ref(), mTriangleCount, nanovdb::cuda::noInit); if (rootBoxCounts.data() == nullptr) throw std::runtime_error("Failed to allocate root box counts buffer"); util::cuda::lambdaKernel<<>>( mTriangleCount, topology::detail::CountRootBoxesFunctor{ deviceXformedTriangles(), - reinterpret_cast(rootBoxCounts.data()), + rootBoxCounts.data(), mBandWidth } ); @@ -516,27 +519,27 @@ void MeshToGrid::processRootTrianglePairs() // Pass 2: InclusiveSum Scan to compute offsets and total allocations - ScratchT rootBoxOffsets = ScratchT(mStream, this->ref(), (mTriangleCount+1)*sizeof(uint64_t), nanovdb::cuda::noInit); + BufT rootBoxOffsets = BufT(mStream, this->ref(), mTriangleCount+1, nanovdb::cuda::noInit); if (rootBoxOffsets.data() == nullptr) throw std::runtime_error("Failed to allocate root box offsets buffer"); cudaCheck(cudaMemsetAsync(rootBoxOffsets.data(), 0, sizeof(uint64_t), mStream)); CALL_CUBS(DeviceScan::InclusiveSum, - reinterpret_cast(rootBoxCounts.data()), - reinterpret_cast(rootBoxOffsets.data())+1, + rootBoxCounts.data(), + rootBoxOffsets.data()+1, mTriangleCount); - cudaCheck(cudaMemcpyAsync(&mBoxTrianglePairCount, reinterpret_cast(rootBoxOffsets.data())+mTriangleCount, sizeof(uint64_t), cudaMemcpyDeviceToHost, mStream)); + cudaCheck(cudaMemcpyAsync(&mBoxTrianglePairCount, rootBoxOffsets.data()+mTriangleCount, sizeof(uint64_t), cudaMemcpyDeviceToHost, mStream)); cudaStreamSynchronize(mStream); // Pass 3: Re-enumerate intersections of (padded) root boxes and triangles, and scatter to allocated list - mBoxTrianglePairsBuffer = ScratchT(mStream, this->ref(), mBoxTrianglePairCount * sizeof(MeshToGridBoxTrianglePair), nanovdb::cuda::noInit); + mBoxTrianglePairsBuffer = BufT(mStream, this->ref(), mBoxTrianglePairCount, nanovdb::cuda::noInit); if (mBoxTrianglePairsBuffer.data() == nullptr) throw std::runtime_error("Failed to allocate pairs buffer"); util::cuda::lambdaKernel<<>>( mTriangleCount, topology::detail::ScatterRootTrianglePairsFunctor{ deviceXformedTriangles(), - reinterpret_cast(rootBoxOffsets.data()), + rootBoxOffsets.data(), deviceBoxTrianglePairs(), mBandWidth } @@ -789,8 +792,8 @@ void MeshToGrid::enumerateRootTiles() cudaGetDevice(&device); // Step 1: Encode each pair's root origin as a sortable uint64_t key - ScratchT keysBuffer = ScratchT(mStream, this->ref(), mBoxTrianglePairCount * sizeof(uint64_t), nanovdb::cuda::noInit); - auto *dKeys = reinterpret_cast(keysBuffer.data()); + BufT keysBuffer = BufT(mStream, this->ref(), mBoxTrianglePairCount, nanovdb::cuda::noInit); + auto *dKeys = keysBuffer.data(); util::cuda::lambdaKernel<<>>( mBoxTrianglePairCount, @@ -799,17 +802,17 @@ void MeshToGrid::enumerateRootTiles() cudaCheckError(); // Step 2: Sort keys (SortKeys requires separate in/out buffers) - ScratchT sortedKeysBuffer = ScratchT(mStream, this->ref(), mBoxTrianglePairCount * sizeof(uint64_t), nanovdb::cuda::noInit); - auto *dSortedKeys = reinterpret_cast(sortedKeysBuffer.data()); + BufT sortedKeysBuffer = BufT(mStream, this->ref(), mBoxTrianglePairCount, nanovdb::cuda::noInit); + auto *dSortedKeys = sortedKeysBuffer.data(); CALL_CUBS(DeviceRadixSort::SortKeys, dKeys, dSortedKeys, (int)mBoxTrianglePairCount, 0, 64); // Step 3: Select unique keys - ScratchT uniqueKeysBuffer = ScratchT(mStream, this->ref(), mBoxTrianglePairCount * sizeof(uint64_t), nanovdb::cuda::noInit); - auto *dUniqueKeys = reinterpret_cast(uniqueKeysBuffer.data()); + BufT uniqueKeysBuffer = BufT(mStream, this->ref(), mBoxTrianglePairCount, nanovdb::cuda::noInit); + auto *dUniqueKeys = uniqueKeysBuffer.data(); - ScratchT numSelectedBuffer = ScratchT(mStream, this->ref(), sizeof(int32_t), nanovdb::cuda::noInit); - auto *dNumSelected = reinterpret_cast(numSelectedBuffer.data()); + BufT numSelectedBuffer = BufT(mStream, this->ref(), 1, nanovdb::cuda::noInit); + auto *dNumSelected = numSelectedBuffer.data(); CALL_CUBS(DeviceSelect::Unique, dSortedKeys, dUniqueKeys, dNumSelected, (int)mBoxTrianglePairCount); @@ -819,7 +822,7 @@ void MeshToGrid::enumerateRootTiles() mUniqueRootTileCount = static_cast(uniqueCount); // Step 4: Decode unique keys back to Coord origins - mUniqueRootOriginsBuffer = ScratchT(mStream, this->ref(), mUniqueRootTileCount * sizeof(nanovdb::Coord), nanovdb::cuda::noInit); + mUniqueRootOriginsBuffer = BufT(mStream, this->ref(), mUniqueRootTileCount, nanovdb::cuda::noInit); auto *dOrigins = deviceUniqueRootOrigins(); util::cuda::lambdaKernel<<>>( @@ -938,19 +941,18 @@ void MeshToGrid::processLeafTrianglePairs() for (int pass = 0; pass < 3; ++pass) { // Allocate Mask<3> buffer for the CTA hit results - MaskBufT maskBuffer = MaskBufT(mStream, this->ref(), mBoxTrianglePairCount, nanovdb::cuda::noInit); + BufT> maskBuffer = BufT>(mStream, this->ref(), mBoxTrianglePairCount, nanovdb::cuda::noInit); if (maskBuffer.data() == nullptr) { throw std::runtime_error("Failed to allocate mask buffer for subdivision pass"); } auto* dMasks = maskBuffer.data(); // Allocate Counts buffer for Prefix Sum - // Size: mBoxTrianglePairCount * sizeof(uint64_t) - ScratchT countsBuffer = ScratchT(mStream, this->ref(), mBoxTrianglePairCount * sizeof(uint64_t), nanovdb::cuda::noInit); + BufT countsBuffer = BufT(mStream, this->ref(), mBoxTrianglePairCount, nanovdb::cuda::noInit); if (countsBuffer.data() == nullptr) { throw std::runtime_error("Failed to allocate counts buffer for subdivision pass"); } - auto* dCounts = reinterpret_cast(countsBuffer.data()); + auto* dCounts = countsBuffer.data(); // Evaluate & Count: 1 CTA per parent pair, 512 threads per CTA. // Uses AABB-only test for large child scales (>= mSATThreshold), full SAT below. @@ -970,10 +972,10 @@ void MeshToGrid::processLeafTrianglePairs() // Prefix Sum: element [i+1] = exclusive write offset for parent i's children, // element [0] = 0, element [mBoxTrianglePairCount] = total child pair count. - ScratchT offsetsBuffer = ScratchT(mStream, this->ref(), (mBoxTrianglePairCount + 1) * sizeof(uint64_t), nanovdb::cuda::noInit); + BufT offsetsBuffer = BufT(mStream, this->ref(), mBoxTrianglePairCount + 1, nanovdb::cuda::noInit); if (offsetsBuffer.data() == nullptr) throw std::runtime_error("Failed to allocate offsets buffer for subdivision pass"); - auto* dOffsets = reinterpret_cast(offsetsBuffer.data()); + auto* dOffsets = offsetsBuffer.data(); cudaCheck(cudaMemsetAsync(dOffsets, 0, sizeof(uint64_t), mStream)); CALL_CUBS(DeviceScan::InclusiveSum, @@ -987,10 +989,10 @@ void MeshToGrid::processLeafTrianglePairs() cudaStreamSynchronize(mStream); // Allocate new child pair buffer - ScratchT newPairsBuffer = ScratchT(mStream, this->ref(), newPairCount * sizeof(BoxTrianglePair), nanovdb::cuda::noInit); + BufT newPairsBuffer = BufT(mStream, this->ref(), newPairCount, nanovdb::cuda::noInit); if (newPairsBuffer.data() == nullptr) throw std::runtime_error("Failed to allocate child pairs buffer for subdivision pass"); - auto* dNewPairs = reinterpret_cast(newPairsBuffer.data()); + auto* dNewPairs = newPairsBuffer.data(); // Scatter surviving child pairs into the new buffer util::cuda::lambdaKernel<<>>( @@ -1127,7 +1129,7 @@ MeshToGrid::getHandleAndUDF(const GridBufferT& buffer, const const uint32_t leafCount = mBuilder.data()->nodeCount[0]; auto handle = GridHandle(std::move(gridBuffer)); if (leafCount) { - MaskBufT retainMaskBuffer = MaskBufT(mStream, this->ref(), leafCount, nanovdb::cuda::noInit); + BufT> retainMaskBuffer = BufT>(mStream, this->ref(), leafCount, nanovdb::cuda::noInit); if (retainMaskBuffer.data() == nullptr) throw std::runtime_error("Failed to allocate retain mask buffer"); cudaCheck(cudaMemsetAsync(retainMaskBuffer.data(), 0xFF, retainMaskBuffer.size_bytes(), mStream)); tools::cuda::PruneGrid pruner( diff --git a/nanovdb/nanovdb/tools/cuda/SignedFloodFill.cuh b/nanovdb/nanovdb/tools/cuda/SignedFloodFill.cuh index a6a15a3eb6..91ecca422c 100644 --- a/nanovdb/nanovdb/tools/cuda/SignedFloodFill.cuh +++ b/nanovdb/nanovdb/tools/cuda/SignedFloodFill.cuh @@ -127,8 +127,8 @@ void processRoot(NanoTree *d_tree, cudaStream_t stream = 0) cudaCheck(cudaMemcpy(root + 1, (char*)(d_tree + 1) + sizeof(RootT), root->tileCount()*sizeof(TileT), cudaMemcpyDeviceToHost));// copy tiles // Sort the child nodes of the root in lexicographic order - ManagedBufT nodeBuffer(root->tileCount()*sizeof(ChildT), nanovdb::cuda::noInit); // potential over-allocation - auto *first = reinterpret_cast(nodeBuffer.data()), *last = first; + nanovdb::cuda::Buffer nodeBuffer(root->tileCount(), nanovdb::cuda::noInit); // potential over-allocation + auto *first = nodeBuffer.data(), *last = first; for (auto it=root->beginChild(); it; ++it) *last++ = ChildT(it.getCoord(), it.pos()); if (last - first < 2) return;// zero or one child node so nothing to do! std::sort(first, last, ChildT());// lexicographic ordering diff --git a/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh b/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh index e8040e8905..692d4100fc 100644 --- a/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh +++ b/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh @@ -24,6 +24,9 @@ #include #include // for the pinned host staging of the processed root +#include // for std::byte, std::size_t +#include // for std::runtime_error + namespace nanovdb { namespace tools::cuda { @@ -60,8 +63,8 @@ class TopologyBuilder static_assert(nanovdb::cuda::is_async_resource::value, "TopologyBuilder allocates stream-ordered scratch and requires an AsyncResource"); - static_assert(ResourceT::DEFAULT_ALIGNMENT >= alignof(uint64_t), - "TopologyBuilder reinterprets byte scratch as word-sized types and requires word-aligned allocations"); + static_assert(ResourceT::DEFAULT_ALIGNMENT >= NANOVDB_DATA_ALIGNMENT, + "TopologyBuilder stages the processed root as bytes and reinterprets it as a root node, which requires NANOVDB_DATA_ALIGNMENT-aligned allocations"); /// @brief Device-only scratch storage, borrowing the injected resource /// through a ResourceRef so all traffic reaches the caller's @@ -70,8 +73,10 @@ class TopologyBuilder /// Buffer rather than the dual DeviceBuffer, whose host pointer and /// per-device array they would leave unused. using ScratchT = nanovdb::cuda::Buffer>; - using UpperMaskBufT = nanovdb::cuda::Buffer, nanovdb::cuda::ResourceRef>; - using LowerMaskBufT = nanovdb::cuda::Buffer, nanovdb::cuda::ResourceRef>; + template + using BufT = nanovdb::cuda::Buffer>; + using UpperMaskBufT = BufT>; + using LowerMaskBufT = BufT>; using HostStagingT = nanovdb::cuda::Buffer; public: @@ -118,14 +123,14 @@ public: ScratchT mDeviceRoot; // device copy, made by uploadProcessedRoot UpperMaskBufT mUpperMasks; LowerMaskBufT mLowerMasks; - ScratchT mUpperOffsets; - ScratchT mLowerOffsets; - ScratchT mLeafOffsets; - ScratchT mVoxelOffsets; - ScratchT mLowerParents; - ScratchT mLeafParents; + BufT mUpperOffsets; + BufT mLowerOffsets; + BufT mLeafOffsets; + BufT mVoxelOffsets; + BufT mLowerParents; + BufT mLeafParents; Data mHostData{}; // host side of the builder parameters - ScratchT mDeviceData; // device copy, made by uploadData + BufT mDeviceData; // device copy, made by uploadData CheckMode mChecksum{CheckMode::Disable}; auto deviceProcessedRoot() { return reinterpret_cast(mDeviceRoot.data()); } @@ -154,20 +159,27 @@ public: void uploadData(cudaStream_t stream) { if (mDeviceData.empty()) - mDeviceData = ScratchT(stream, nanovdb::cuda::ResourceRef(*mResource), sizeof(Data), nanovdb::cuda::noInit); - cudaCheck(cudaMemcpyAsync(mDeviceData.data(), &mHostData, sizeof(Data), cudaMemcpyHostToDevice, stream)); + mDeviceData = BufT(stream, nanovdb::cuda::ResourceRef(*mResource), 1, nanovdb::cuda::noInit); + cudaCheck(cudaMemcpyAsync(mDeviceData.data(), &mHostData, mDeviceData.size_bytes(), cudaMemcpyHostToDevice, stream)); } Mask<5>* deviceUpperMasks() { return mUpperMasks.data(); } /// @brief The densified lower masks: one row of Mask<5>::SIZE Mask<4> per upper node, /// indexed [upper node][lower node offset]. The row shape is fixed here, beside the /// allocation that defines it, so consumers never re-derive the stride. Mask<4> (*deviceLowerMasks())[Mask<5>::SIZE] { return reinterpret_cast(*)[Mask<5>::SIZE]>(mLowerMasks.data()); } + //@{ + /// @brief The lower and leaf node offsets viewed one row of Mask<5>::SIZE per upper node, + /// indexed [upper node][lower node offset]; the row shape is fixed here, beside the + /// allocation that defines it, so consumers never re-derive the stride + uint32_t (*lowerOffsetRows())[Mask<5>::SIZE] { return reinterpret_cast::SIZE]>(mLowerOffsets.data()); } + uint32_t (*leafOffsetRows())[Mask<5>::SIZE] { return reinterpret_cast::SIZE]>(mLeafOffsets.data()); } + //@} /// @brief A borrowing reference to the builder's resource, for consumers /// allocating sibling scratch from the same instance. nanovdb::cuda::ResourceRef ref() { return nanovdb::cuda::ResourceRef(*mResource); } Data* data() { return &mHostData; } - Data* deviceData() { return reinterpret_cast(mDeviceData.data()); } + Data* deviceData() { return mDeviceData.data(); } private: static constexpr unsigned int mNumThreads = 128;// for kernels spawned via lambdaKernel (others may specialize) @@ -231,9 +243,9 @@ void TopologyBuilder::countNodes(cudaStream_t stream) // as well as the tile table at the root. std::size_t size = processedTileCount*Mask<5>::SIZE; - ScratchT upperCountsBuffer = ScratchT(stream, *mResource, processedTileCount*sizeof(uint32_t), nanovdb::cuda::noInit); - ScratchT lowerCountsBuffer = ScratchT(stream, *mResource, size*sizeof(uint32_t), nanovdb::cuda::noInit); - ScratchT leafCountsBuffer = ScratchT(stream, *mResource, size*sizeof(uint32_t), nanovdb::cuda::noInit); + BufT upperCountsBuffer = BufT(stream, *mResource, processedTileCount, nanovdb::cuda::noInit); + BufT lowerCountsBuffer = BufT(stream, *mResource, size, nanovdb::cuda::noInit); + BufT leafCountsBuffer = BufT(stream, *mResource, size, nanovdb::cuda::noInit); using CountType = uint32_t (*)[Mask<5>::SIZE]; auto lowerCounts = reinterpret_cast(lowerCountsBuffer.data()); @@ -244,37 +256,37 @@ void TopologyBuilder::countNodes(cudaStream_t stream) <<>> (deviceUpperMasks(), deviceLowerMasks(), lowerCounts, leafCounts); - mUpperOffsets = ScratchT(stream, *mResource, (processedTileCount+1)*sizeof(uint32_t), nanovdb::cuda::noInit); - mLowerOffsets = ScratchT(stream, *mResource, (size+1)*sizeof(uint32_t), nanovdb::cuda::noInit); - mLeafOffsets = ScratchT(stream, *mResource, (size+1)*sizeof(uint32_t), nanovdb::cuda::noInit); + mUpperOffsets = BufT(stream, *mResource, processedTileCount+1, nanovdb::cuda::noInit); + mLowerOffsets = BufT(stream, *mResource, size+1, nanovdb::cuda::noInit); + mLeafOffsets = BufT(stream, *mResource, size+1, nanovdb::cuda::noInit); cudaCheck(cudaMemsetAsync(mLowerOffsets.data(), 0, sizeof(uint32_t), stream)); CALL_CUBS(DeviceScan::InclusiveSum, - reinterpret_cast(lowerCountsBuffer.data()), - reinterpret_cast(mLowerOffsets.data())+1, + lowerCountsBuffer.data(), + mLowerOffsets.data()+1, size); - cudaCheck(cudaMemcpyAsync(&data()->nodeCount[1], reinterpret_cast(mLowerOffsets.data())+size, sizeof(uint32_t), cudaMemcpyDeviceToHost, stream)); + cudaCheck(cudaMemcpyAsync(&data()->nodeCount[1], mLowerOffsets.data()+size, sizeof(uint32_t), cudaMemcpyDeviceToHost, stream)); cudaCheck(cudaMemsetAsync(mLeafOffsets.data(), 0, sizeof(uint32_t), stream)); CALL_CUBS(DeviceScan::InclusiveSum, - reinterpret_cast(leafCountsBuffer.data()), - reinterpret_cast(mLeafOffsets.data())+1, + leafCountsBuffer.data(), + mLeafOffsets.data()+1, size); - cudaCheck(cudaMemcpyAsync(&data()->nodeCount[0], reinterpret_cast(mLeafOffsets.data())+size, sizeof(uint32_t), cudaMemcpyDeviceToHost, stream)); + cudaCheck(cudaMemcpyAsync(&data()->nodeCount[0], mLeafOffsets.data()+size, sizeof(uint32_t), cudaMemcpyDeviceToHost, stream)); util::cuda::lambdaKernel<<>>( processedTileCount, [] __device__(size_t tileID, CountType lowerOffsets, uint32_t* upperCounts) { upperCounts[tileID] = (lowerOffsets[tileID+1][0] > lowerOffsets[tileID][0]) ? 1 : 0; }, - reinterpret_cast(mLowerOffsets.data()), - reinterpret_cast(upperCountsBuffer.data())); + lowerOffsetRows(), + upperCountsBuffer.data()); cudaCheck(cudaMemsetAsync( mUpperOffsets.data(), 0, sizeof(uint32_t), stream)); CALL_CUBS(DeviceScan::InclusiveSum, - reinterpret_cast(upperCountsBuffer.data()), - reinterpret_cast(mUpperOffsets.data())+1, + upperCountsBuffer.data(), + mUpperOffsets.data()+1, processedTileCount); - cudaCheck(cudaMemcpyAsync(&data()->nodeCount[2], reinterpret_cast(mUpperOffsets.data())+processedTileCount, sizeof(uint32_t), cudaMemcpyDeviceToHost, stream)); + cudaCheck(cudaMemcpyAsync(&data()->nodeCount[2], mUpperOffsets.data()+processedTileCount, sizeof(uint32_t), cudaMemcpyDeviceToHost, stream)); }// TopologyBuilder::countNodes //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- @@ -300,7 +312,7 @@ BufferT TopologyBuilder::getBuffer(const BufferT &pool, cudaS data()->d_bufferPtr = nanovdb::cuda::detail::deviceStorageData(buffer); if (data()->d_bufferPtr == nullptr) throw std::runtime_error("Failed to allocate grid buffer on the device"); if (data()->nodeCount[2] != 0) // Unless the result is an empty grid - data()->d_upperOffsets = reinterpret_cast(mUpperOffsets.data()); + data()->d_upperOffsets = mUpperOffsets.data(); this->uploadData(stream); return buffer; @@ -502,25 +514,24 @@ inline void TopologyBuilder::processLowerNodes(cudaStream_t s // Fill out the contents of all newly allocated lower nodes (using the densified upper/lower mask arrays) // Also fill in the preamble (most of LeafData) for their leaf children auto processedTileCount = hostProcessedRoot()->tileCount(); - using CountType = uint32_t (*)[Mask<5>::SIZE]; if (processedTileCount) { // Unless output grid is empty std::size_t lowerCount = data()->nodeCount[1]; - mLowerParents = ScratchT(stream, *mResource, lowerCount*sizeof(uint32_t), nanovdb::cuda::noInit); + mLowerParents = BufT(stream, *mResource, lowerCount, nanovdb::cuda::noInit); std::size_t leafCount = data()->nodeCount[0]; - mLeafParents = ScratchT(stream, *mResource, leafCount*sizeof(uint32_t), nanovdb::cuda::noInit); + mLeafParents = BufT(stream, *mResource, leafCount, nanovdb::cuda::noInit); using Op = util::morphology::cuda::ProcessLowerNodesFunctor; util::cuda::operatorKernel <<>>( deviceUpperMasks(), deviceLowerMasks(), - reinterpret_cast(mUpperOffsets.data()), - reinterpret_cast(mLowerOffsets.data()), - reinterpret_cast(mLeafOffsets.data()), + mUpperOffsets.data(), + lowerOffsetRows(), + leafOffsetRows(), static_cast(data()->d_bufferPtr), - reinterpret_cast(mLowerParents.data()), - reinterpret_cast(mLeafParents.data()) + mLowerParents.data(), + mLeafParents.data() ); cudaCheckError(); } @@ -570,16 +581,16 @@ inline void TopologyBuilder::processLeafOffsets(cudaStream_t { std::size_t leafCount = data()->nodeCount[0]; if (leafCount) { // Unless output grid is empty - mVoxelOffsets = ScratchT(stream, *mResource, (leafCount+1)*sizeof(uint64_t), nanovdb::cuda::noInit); + mVoxelOffsets = BufT(stream, *mResource, leafCount+1, nanovdb::cuda::noInit); cudaCheck(cudaMemsetAsync(mVoxelOffsets.data(), 0, sizeof(uint64_t), stream)); util::cuda::lambdaKernel<<>>( - leafCount, topology::detail::UpdateLeafVoxelCountsAndPrefixSumFunctor(), deviceData(), reinterpret_cast(mVoxelOffsets.data())+1); + leafCount, topology::detail::UpdateLeafVoxelCountsAndPrefixSumFunctor(), deviceData(), mVoxelOffsets.data()+1); CALL_CUBS(DeviceScan::InclusiveSum, - reinterpret_cast(mVoxelOffsets.data())+1, - reinterpret_cast(mVoxelOffsets.data())+1, + mVoxelOffsets.data()+1, + mVoxelOffsets.data()+1, leafCount); util::cuda::lambdaKernel<<>>( - leafCount, topology::detail::UpdateLeafVoxelOffsetsFunctor(), deviceData(), reinterpret_cast(mVoxelOffsets.data())); + leafCount, topology::detail::UpdateLeafVoxelOffsetsFunctor(), deviceData(), mVoxelOffsets.data()); } }// TopologyBuilder::processLeafOffsets @@ -649,13 +660,13 @@ inline void TopologyBuilder::processBBox(cudaStream_t stream) // update and propagate bbox from leaf -> lower/parent nodes util::cuda::lambdaKernel<<nodeCount[0]), mNumThreads, 0, stream>>>( - data()->nodeCount[0], topology::detail::UpdateAndPropagateLeafBBoxFunctor(), deviceData(), reinterpret_cast(mLeafParents.data())); + data()->nodeCount[0], topology::detail::UpdateAndPropagateLeafBBoxFunctor(), deviceData(), mLeafParents.data()); mLeafParents.destroy(stream); cudaCheckError(); // propagate bbox from lower -> upper/parent node util::cuda::lambdaKernel<<nodeCount[1]), mNumThreads, 0, stream>>>( - data()->nodeCount[1], topology::detail::PropagateLowerBBoxFunctor(), deviceData(), reinterpret_cast(mLowerParents.data())); + data()->nodeCount[1], topology::detail::PropagateLowerBBoxFunctor(), deviceData(), mLowerParents.data()); mLowerParents.destroy(stream); cudaCheckError(); @@ -692,7 +703,7 @@ inline void TopologyBuilder::postProcessGridTree(cudaStream_t { // Finish updates to GridData/TreeData and (optionally) update checksum if (data()->nodeCount[0]) // if grid is empty, the default values are correct - util::cuda::lambdaKernel<<<1, 1, 0, stream>>>(1, topology::detail::PostProcessGridTreeFunctor(), deviceData(), reinterpret_cast(mVoxelOffsets.data())); + util::cuda::lambdaKernel<<<1, 1, 0, stream>>>(1, topology::detail::PostProcessGridTreeFunctor(), deviceData(), mVoxelOffsets.data()); cudaCheckError(); mVoxelOffsets.destroy(stream); diff --git a/pendingchanges/nanovdb.txt b/pendingchanges/nanovdb.txt index 56845b81b7..0ab66077ce 100644 --- a/pendingchanges/nanovdb.txt +++ b/pendingchanges/nanovdb.txt @@ -16,7 +16,7 @@ NanoVDB: - nanovdb::cuda::createNodeManager gained a single-space overload: passing a memory resource instead of a pool buffer returns a NodeManagerHandle over cuda::Buffer> whose storage (and size scratch) allocates through that resource, which must outlive the handle. tools::cuda::GridStats now builds its temporary NodeManager through its injected resource, so every device allocation it makes honors ResourceT. NodeManagerHandle supports single-space buffers (deviceMgr maps onto the buffer, host accessors are compile-time errors) and no longer requires a default-constructible buffer type. - GridHandle::reset, NodeManagerHandle::reset and tools::VoxelBlockManager::reset now release storage through the buffer's destroy() when it provides one, and cuda::Buffer::clear/cuda::BufferView::clear are deprecated in favor of destroy() (BufferView::destroy detaches the non-owning view). The legacy HostBuffer/DeviceBuffer clear() methods are unaffected. - The bug-fix to the nanovdb::ReadAccessor (see below) improves random-access performance in some use-cases (especially on the CPU). - - The device-side mask scratch of tools::cuda::TopologyBuilder and tools::cuda::MeshToGrid, the dilation example and its test are now typed cuda::Buffer> instead of byte buffers reinterpreted as masks (possible since Mask became trivially copyable); TopologyBuilder::deviceUpperMasks/deviceLowerMasks return typed pointers and the morphology functors take them. + - The device scratch of tools::cuda::TopologyBuilder, tools::cuda::MeshToGrid and tools::cuda::SignedFloodFill is now typed cuda::Buffer (masks, offsets, parents, counts, keys, triangles, box-triangle pairs, root origins, the builder's Data struct) instead of byte buffers reinterpreted at every use, with allocations sized in elements; TopologyBuilder's accessors return typed pointers, the lower/leaf offsets and masks expose their per-upper-node row shape once beside the allocation, and the morphology functors take the typed pointers. The dilation example and test use cuda::Buffer> likewise (possible since Mask became trivially copyable). TopologyBuilder's variable-length root staging stays a byte buffer, being a serialized node image. - tools::cuda::DistributedPointsToGrid gained a defaulted ResourceT template parameter and constructor overloads taking one memory-resource instance per device in the mesh, indexed by device id: the CUB scratch pools and the per-device sort scratch route through the injected instances (memory pools are per-device, hence one instance each). The sort scratch is now a stream-ordered cuda::Buffer, replacing a synchronous cudaFree that stalled the device mid-pipeline, and the shared managed metadata and pinned merge intervals are owned by cuda::Buffer members over ManagedResource and PinnedResource instead of hand-paired allocate/free calls. The grid buffer can be a managed single-space handle: getHandle accepts cuda::Buffer alongside the legacy default. The multi-GPU tests and examples build managed cuda::Buffer handles now. Deprecations: