From 04fd4bc3174f189b4a624cef8abed30a56884f7a Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 5 Aug 2026 01:18:58 +0000 Subject: [PATCH 01/17] NanoVDB: align cuda::Buffer member names with cuda::buffer Name the free operation destroy, as cuda::buffer does, and add the stream-taking overload it also provides. clear stays but only as a transitional delegate, marked as such: it exists because GridHandle::reset still calls it, and goes away with the legacy dual buffers that cuda::Buffer replaces. Rename setStream to set_stream. Member names cannot be aliased, so matching the standard spelling is the whole reason the resource concept kept allocate_async; the same argument applies here. Note in passing that ours deliberately does not synchronize, matching cuda::buffer's set_stream_unsynchronized rather than its set_stream, whose own documentation and implementation disagree (NVIDIA/cccl#10649). Add swap. The generic std::swap already does the right thing through the move operations, but both std::vector and rmm::device_buffer provide one. Co-Authored-By: Claude Opus 5 (1M context) Signed-off-by: Mark Harris --- nanovdb/nanovdb/cuda/Buffer.h | 49 ++++++++++++++++--- nanovdb/nanovdb/unittest/TestBuffer.cu | 66 ++++++++++++++++++++++++-- 2 files changed, 103 insertions(+), 12 deletions(-) diff --git a/nanovdb/nanovdb/cuda/Buffer.h b/nanovdb/nanovdb/cuda/Buffer.h index 41d15d511a..d9144591ca 100644 --- a/nanovdb/nanovdb/cuda/Buffer.h +++ b/nanovdb/nanovdb/cuda/Buffer.h @@ -50,7 +50,7 @@ struct StreamHolder { cudaStream_t mStream = 0; }; /// is_resource). When @c R provides both interfaces the stream-ordered /// one is used. /// @details With a stream-ordered resource the Buffer retains the stream of -/// the most recent allocation (or the one supplied via setStream) +/// the most recent allocation (or the one supplied via set_stream) /// and orders its deallocation on that stream. Buffer is move-only. template class Buffer : private detail::StreamHolder::value> @@ -133,7 +133,7 @@ class Buffer : private detail::StreamHolder::value> Buffer& operator=(Buffer&& other) noexcept { if (this != &other) { - this->clear(); + this->destroy(); static_cast&>(*this) = other; mResource = std::move(other.mResource); mData = other.mData; @@ -168,7 +168,7 @@ class Buffer : private detail::StreamHolder::value> /// @brief D-tor. A stream-ordered resource frees on the retained stream; /// a synchronous resource frees immediately. - ~Buffer() { this->clear(); } + ~Buffer() { this->destroy(); } /// @brief Returns the retained stream, i.e. the stream the buffer's memory /// will be freed on. @@ -179,9 +179,11 @@ class Buffer : private detail::StreamHolder::value> /// deallocation (and destruction) is ordered on @c stream instead. /// @param stream cuda stream subsequent deallocation is ordered on /// @warning The caller is responsible for ordering @c stream after any - /// in-flight work that uses the buffer's memory. + /// in-flight work that uses the buffer's memory. This deliberately + /// does not synchronize, matching cuda::buffer's + /// set_stream_unsynchronized rather than its set_stream. template::value, int> = 0> - void setStream(cudaStream_t stream) { this->mStream = stream; } + void set_stream(cudaStream_t stream) { this->mStream = stream; } /// @brief Resizes the buffer to @c count elements, preserving the leading /// min(old, new) elements. Every operation — the new allocation, the @@ -219,7 +221,7 @@ class Buffer : private detail::StreamHolder::value> mSize = count; } else { - this->mStream = stream; // no reallocation: setStream semantics + this->mStream = stream; // no reallocation: set_stream semantics } } @@ -263,13 +265,46 @@ class Buffer : private detail::StreamHolder::value> bool empty() const { return mSize == 0; } /// @brief Frees the buffer memory (if any) and resets to the empty state. - void clear() + /// A stream-ordered resource frees on the retained stream. + /// @note Spelled destroy to match cuda::buffer. This is the name to use. + void destroy() { this->deallocate(mData, mSize); mData = nullptr; mSize = 0; } + /// @brief Frees the buffer memory (if any) and resets to the empty state. + /// @note Transitional, and not the name to use: it exists only because + /// GridHandle::reset still calls clear() on its buffer. It goes away + /// when the legacy dual buffers do and GridHandle moves to destroy(). + void clear() { this->destroy(); } + + /// @brief Frees the buffer memory (if any) on @c stream and resets to the + /// empty state. @c stream becomes the retained stream. + /// @param stream cuda stream the deallocation is ordered on + /// @warning The caller is responsible for ordering @c stream after any + /// in-flight work that uses the buffer's memory. + template::value, int> = 0> + void destroy(cudaStream_t stream) + { + this->mStream = stream; + this->destroy(); + } + + /// @brief Exchanges the contents of this buffer with @c other. Neither + /// buffer allocates, frees, or copies element data. + /// @param other buffer to exchange contents with + void swap(Buffer& other) noexcept + { + auto& lhs = static_cast&>(*this); + auto& rhs = static_cast&>(other); + std::swap(lhs, rhs); + std::swap(mResource, other.mResource); + std::swap(mData, other.mData); + std::swap(mSize, other.mSize); + } + private: /// @brief Returns @c count * sizeof(T), throwing std::runtime_error if the /// byte size would overflow size_t instead of silently wrapping into diff --git a/nanovdb/nanovdb/unittest/TestBuffer.cu b/nanovdb/nanovdb/unittest/TestBuffer.cu index 6a871802c6..40be6878a6 100644 --- a/nanovdb/nanovdb/unittest/TestBuffer.cu +++ b/nanovdb/nanovdb/unittest/TestBuffer.cu @@ -258,9 +258,65 @@ TEST(TestBuffer, ClearFreesAndEmpties) ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess); } +TEST(TestBuffer, DestroyIsTheSpellingClearDelegatesTo) +{ + Counters c; + nanovdb::cuda::Buffer buf(0, CountingResource{&c}, 64, nanovdb::cuda::noInit); + ASSERT_EQ(c.allocs, 1); + buf.destroy(); // cuda::buffer's spelling + EXPECT_EQ(buf.data(), nullptr); + EXPECT_EQ(buf.size(), 0u); + EXPECT_EQ(c.deallocs, 1); + buf.destroy(); // idempotent + EXPECT_EQ(c.deallocs, 1); + ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess); +} + +TEST(TestBuffer, DestroyOnStreamRetargetsTheFree) +{ + cudaStream_t a, b; + ASSERT_EQ(cudaStreamCreate(&a), cudaSuccess); + ASSERT_EQ(cudaStreamCreate(&b), cudaSuccess); + { + StreamLog log; + nanovdb::cuda::Buffer buf(a, StreamRecordingResource{&log}, 32, nanovdb::cuda::noInit); + ASSERT_EQ(log.allocStreams.size(), 1u); + EXPECT_EQ(log.allocStreams[0], a); // allocated on a + buf.destroy(b); // explicit stream overload + ASSERT_EQ(log.deallocStreams.size(), 1u); + EXPECT_EQ(log.deallocStreams[0], b); // freed on b, not the retained stream a + EXPECT_EQ(buf.stream(), b); // b is retained afterwards + } + ASSERT_EQ(cudaStreamSynchronize(a), cudaSuccess); + ASSERT_EQ(cudaStreamSynchronize(b), cudaSuccess); + cudaStreamDestroy(a); + cudaStreamDestroy(b); +} + +TEST(TestBuffer, SwapExchangesWithoutAllocatingOrFreeing) +{ + Counters c; + nanovdb::cuda::Buffer x(0, CountingResource{&c}, 128, nanovdb::cuda::noInit); + nanovdb::cuda::Buffer y(0, CountingResource{&c}, 64, nanovdb::cuda::noInit); + ASSERT_EQ(c.allocs, 2); + auto* px = x.data(); + auto* py = y.data(); + const int allocs = c.allocs, deallocs = c.deallocs; + + x.swap(y); + + EXPECT_EQ(c.allocs, allocs); // no allocation + EXPECT_EQ(c.deallocs, deallocs);// no free + EXPECT_EQ(x.data(), py); + EXPECT_EQ(y.data(), px); + EXPECT_EQ(x.size(), 64u); + EXPECT_EQ(y.size(), 128u); + ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess); +} + //====================================================================== // Stream retention: the destructor and resize free on the retained stream -// (stream of the most recent allocation, or the one supplied via setStream) +// (stream of the most recent allocation, or the one supplied via set_stream) //====================================================================== TEST(TestBuffer, DestructorFreesOnAllocationStream) @@ -288,7 +344,7 @@ TEST(TestBuffer, SetStreamRedirectsTheFree) StreamLog log; { nanovdb::cuda::Buffer buf(a, StreamRecordingResource{&log}, 32, nanovdb::cuda::noInit); - buf.setStream(b); // member update only, no synchronization + buf.set_stream(b); // member update only, no synchronization EXPECT_EQ(buf.stream(), b); } ASSERT_EQ(log.allocStreams.size(), 1u); @@ -473,7 +529,7 @@ TEST(TestBuffer, AsyncPathIsGraphCapturable) template struct HasSetStream : std::false_type {}; template -struct HasSetStream().setStream(cudaStream_t{0}))>> : std::true_type {}; +struct HasSetStream().set_stream(cudaStream_t{0}))>> : std::true_type {}; template struct HasStreamGetter : std::false_type {}; @@ -484,9 +540,9 @@ using PinnedBufferF = nanovdb::cuda::Buffer; // A Buffer over a synchronous resource exposes no stream API at all. -static_assert(!HasSetStream::value, "sync-resource Buffer must not expose setStream"); +static_assert(!HasSetStream::value, "sync-resource Buffer must not expose set_stream"); static_assert(!HasStreamGetter::value, "sync-resource Buffer must not expose stream()"); -static_assert(HasSetStream::value, "async-resource Buffer exposes setStream"); +static_assert(HasSetStream::value, "async-resource Buffer exposes set_stream"); static_assert(HasStreamGetter::value, "async-resource Buffer exposes stream()"); TEST(TestBuffer, PinnedBufferIsPageLocked) From a5ad1fcca6de50e088b42afcf9da0957819c5ef2 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 5 Aug 2026 01:19:08 +0000 Subject: [PATCH 02/17] NanoVDB: allocate TopologyBuilder scratch from an injected resource Eleven of the builder's buffers are device-only -- nothing reads them on the host -- yet they were cuda::DeviceBuffer, whose host pointer and per-device array they never use. Move them to the single-space cuda::Buffer and give the builder a resource parameter, so its scratch is injectable like PointsToGrid's already is, and freed by scope rather than by hand. mProcessedRoot and mData are read on the host and stay dual. This is not a speedup: DeviceBuffer::init allocates host or device memory but never both, and these buffers always passed a real device id, so no cudaMallocHost was on this path to begin with. Measured dilate on 1k/20k/ 200k points, and the difference is within run-to-run noise. Hoist the Data struct out of the class. It does not depend on the resource, and leaving it nested would give every ResourceT its own incompatible type for the device functors to name. Co-Authored-By: Claude Opus 5 (1M context) Signed-off-by: Mark Harris --- nanovdb/nanovdb/tools/cuda/CoarsenGrid.cuh | 2 +- nanovdb/nanovdb/tools/cuda/PruneGrid.cuh | 2 +- nanovdb/nanovdb/tools/cuda/RefineGrid.cuh | 2 +- .../nanovdb/tools/cuda/TopologyBuilder.cuh | 243 ++++++++++-------- 4 files changed, 135 insertions(+), 114 deletions(-) diff --git a/nanovdb/nanovdb/tools/cuda/CoarsenGrid.cuh b/nanovdb/nanovdb/tools/cuda/CoarsenGrid.cuh index d22e5462b4..f0686eac84 100644 --- a/nanovdb/nanovdb/tools/cuda/CoarsenGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/CoarsenGrid.cuh @@ -210,7 +210,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.deviceData(), mBuilder.mLowerMasks.deviceData() ); + mDeviceSrcGrid, mBuilder.deviceProcessedRoot(), mBuilder.mUpperMasks.data(), mBuilder.mLowerMasks.data() ); } }// CoarsenGrid::coarsenInternalNodes diff --git a/nanovdb/nanovdb/tools/cuda/PruneGrid.cuh b/nanovdb/nanovdb/tools/cuda/PruneGrid.cuh index 113acb6837..cb445c23ec 100644 --- a/nanovdb/nanovdb/tools/cuda/PruneGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/PruneGrid.cuh @@ -213,7 +213,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.deviceData(), mBuilder.mLowerMasks.deviceData() ); + mDeviceSrcGrid, mBuilder.deviceProcessedRoot(), mDeviceSrcLeafMask, mBuilder.mUpperMasks.data(), mBuilder.mLowerMasks.data() ); } }// PruneGrid::pruneInternalNodes diff --git a/nanovdb/nanovdb/tools/cuda/RefineGrid.cuh b/nanovdb/nanovdb/tools/cuda/RefineGrid.cuh index d611629abd..07704aa16d 100644 --- a/nanovdb/nanovdb/tools/cuda/RefineGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/RefineGrid.cuh @@ -225,7 +225,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.deviceData(), mBuilder.mLowerMasks.deviceData() ); + mDeviceSrcGrid, mBuilder.deviceProcessedRoot(), mBuilder.mUpperMasks.data(), mBuilder.mLowerMasks.data() ); } }// RefineGrid::refineInternalNodes diff --git a/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh b/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh index 6ccbb25e6e..1aa5a68c28 100644 --- a/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh +++ b/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh @@ -17,6 +17,8 @@ #include #include +#include +#include #include #include @@ -24,7 +26,25 @@ namespace nanovdb { namespace tools::cuda { +/// @brief Shared grid/tree offsets and node counts handed to the device +/// functors. Independent of the resource the builder allocates from, +/// so it lives outside TopologyBuilder and stays one type across every +/// ResourceT instantiation. template +struct TopologyBuilderData { + void *d_bufferPtr; + uint64_t grid, tree, root, upper, lower, leaf, size;// byte offsets to nodes in buffer + uint32_t nodeCount[3];// 0=leaf,1=lower, 2=upper + uint32_t *d_upperOffsets; + __hostdev__ NanoGrid& getGrid() const {return *util::PtrAdd>(d_bufferPtr, grid);} + __hostdev__ NanoTree& getTree() const {return *util::PtrAdd>(d_bufferPtr, tree);} + __hostdev__ NanoRoot& getRoot() const {return *util::PtrAdd>(d_bufferPtr, root);} + __hostdev__ NanoUpper& getUpper(int i) const {return *(util::PtrAdd>(d_bufferPtr, upper)+i);} + __hostdev__ NanoLower& getLower(int i) const {return *(util::PtrAdd>(d_bufferPtr, lower)+i);} + __hostdev__ NanoLeaf& getLeaf(int i) const {return *(util::PtrAdd>(d_bufferPtr, leaf)+i);} +};// TopologyBuilderData + +template class TopologyBuilder { static_assert(nanovdb::BuildTraits::is_onindex);// For now, only OnIndexGrids supported @@ -36,25 +56,25 @@ class TopologyBuilder using LowerT = NanoLower; using LeafT = NanoLeaf; + /// @brief Device-only scratch storage, allocated from the injected + /// resource. These buffers are never read on the host, so they use + /// the single-space Buffer rather than the dual DeviceBuffer, whose + /// host pointer and per-device array they would leave unused. + using ScratchT = nanovdb::cuda::Buffer; + public: - TopologyBuilder(cudaStream_t stream) + /// @param stream cuda stream the scratch allocations are ordered on + /// @param resource resource instance all device scratch is allocated from; + /// must outlive this builder + TopologyBuilder(cudaStream_t stream, ResourceT& resource = nanovdb::cuda::default_resource()) + : mResource(&resource) + , mTempDevicePool(resource) { mData = nanovdb::cuda::DeviceBuffer::create(sizeof(Data)); } - struct Data { - void *d_bufferPtr; - uint64_t grid, tree, root, upper, lower, leaf, size;// byte offsets to nodes in buffer - uint32_t nodeCount[3];// 0=leaf,1=lower, 2=upper - uint32_t *d_upperOffsets; - __hostdev__ GridT& getGrid() const {return *util::PtrAdd(d_bufferPtr, grid);} - __hostdev__ TreeT& getTree() const {return *util::PtrAdd(d_bufferPtr, tree);} - __hostdev__ RootT& getRoot() const {return *util::PtrAdd(d_bufferPtr, root);} - __hostdev__ UpperT& getUpper(int i) const {return *(util::PtrAdd(d_bufferPtr, upper)+i);} - __hostdev__ LowerT& getLower(int i) const {return *(util::PtrAdd(d_bufferPtr, lower)+i);} - __hostdev__ LeafT& getLeaf(int i) const {return *(util::PtrAdd(d_bufferPtr, leaf)+i);} - };// Data + using Data = TopologyBuilderData; void allocateInternalMaskBuffers(cudaStream_t stream); @@ -74,21 +94,21 @@ public: void postProcessGridTree(cudaStream_t stream); nanovdb::cuda::DeviceBuffer mProcessedRoot; - nanovdb::cuda::DeviceBuffer mUpperMasks; - nanovdb::cuda::DeviceBuffer mLowerMasks; - nanovdb::cuda::DeviceBuffer mUpperOffsets; - nanovdb::cuda::DeviceBuffer mLowerOffsets; - nanovdb::cuda::DeviceBuffer mLeafOffsets; - nanovdb::cuda::DeviceBuffer mVoxelOffsets; - nanovdb::cuda::DeviceBuffer mLowerParents; - nanovdb::cuda::DeviceBuffer mLeafParents; + ScratchT mUpperMasks; + ScratchT mLowerMasks; + ScratchT mUpperOffsets; + ScratchT mLowerOffsets; + ScratchT mLeafOffsets; + ScratchT mVoxelOffsets; + ScratchT mLowerParents; + ScratchT mLeafParents; nanovdb::cuda::DeviceBuffer mData; CheckMode mChecksum{CheckMode::Disable}; auto deviceProcessedRoot() { return static_cast(mProcessedRoot.deviceData()); } auto hostProcessedRoot() { return static_cast(mProcessedRoot.data()); } - void* deviceUpperMasks() { return mUpperMasks.deviceData(); } - void* deviceLowerMasks() { return mLowerMasks.deviceData(); } + void* deviceUpperMasks() { return mUpperMasks.data(); } + void* deviceLowerMasks() { return mLowerMasks.data(); } Data* data() { return static_cast(mData.data()); } Data* deviceData() { return static_cast(mData.deviceData()); } @@ -96,8 +116,9 @@ private: static constexpr unsigned int mNumThreads = 128;// for kernels spawned via lambdaKernel (others may specialize) static unsigned int numBlocks(unsigned int n) {return (n + mNumThreads - 1) / mNumThreads;} - nanovdb::cuda::TempDevicePool mTempDevicePool; -};// tools::cuda::TopologyBuilder + ResourceT* mResource;// non-owning; all device scratch routes through this instance + nanovdb::cuda::TempPool mTempDevicePool; +};// tools::cuda::TopologyBuilder //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- @@ -118,8 +139,8 @@ private: //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template -void TopologyBuilder::allocateInternalMaskBuffers(cudaStream_t stream) +template +void TopologyBuilder::allocateInternalMaskBuffers(cudaStream_t stream) { if (hostProcessedRoot()->tileCount() == 0) return; // Processing empty grid(s); nothing to allocate @@ -130,18 +151,18 @@ void TopologyBuilder::allocateInternalMaskBuffers(cudaStream_t stream) cudaGetDevice(&device); uint64_t upperSize = hostProcessedRoot()->tileCount() * sizeof(Mask<5>); uint64_t lowerSize = hostProcessedRoot()->tileCount() * Mask<5>::SIZE * sizeof(Mask<4>); - mUpperMasks = nanovdb::cuda::DeviceBuffer::create(upperSize, nullptr, device, stream); - if (mUpperMasks.deviceData() == nullptr) throw std::runtime_error("Failed to allocate upper mask buffer on device"); - cudaCheck(cudaMemsetAsync(mUpperMasks.deviceData(), 0, upperSize, stream)); - mLowerMasks = nanovdb::cuda::DeviceBuffer::create( lowerSize, nullptr, device, stream ); - if (mLowerMasks.deviceData() == nullptr) throw std::runtime_error("Failed to allocate lower mask buffer on device"); - cudaCheck(cudaMemsetAsync(mLowerMasks.deviceData(), 0, lowerSize, stream)); -}// TopologyBuilder::allocateInternalMaskBuffers + mUpperMasks = ScratchT(stream, *mResource, upperSize, 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); + if (mLowerMasks.data() == nullptr) throw std::runtime_error("Failed to allocate lower mask buffer on device"); + cudaCheck(cudaMemsetAsync(mLowerMasks.data(), 0, lowerSize, stream)); +}// TopologyBuilder::allocateInternalMaskBuffers //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template -void TopologyBuilder::countNodes(cudaStream_t stream) +template +void TopologyBuilder::countNodes(cudaStream_t stream) { auto processedTileCount = hostProcessedRoot()->tileCount(); if (processedTileCount == 0) { // Processing empty grid(s); zero nodes at all levels @@ -157,57 +178,57 @@ void TopologyBuilder::countNodes(cudaStream_t stream) int device = 0; cudaGetDevice(&device); - nanovdb::cuda::DeviceBuffer upperCountsBuffer = nanovdb::cuda::DeviceBuffer::create(processedTileCount*sizeof(uint32_t), nullptr, device, stream); - nanovdb::cuda::DeviceBuffer lowerCountsBuffer = nanovdb::cuda::DeviceBuffer::create(size*sizeof(uint32_t), nullptr, device, stream); - nanovdb::cuda::DeviceBuffer leafCountsBuffer = nanovdb::cuda::DeviceBuffer::create(size*sizeof(uint32_t), nullptr, device, stream); + 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); using CountType = uint32_t (*)[Mask<5>::SIZE]; - auto lowerCounts = reinterpret_cast( lowerCountsBuffer.deviceData() ); - auto leafCounts = reinterpret_cast( leafCountsBuffer.deviceData() ); + auto lowerCounts = reinterpret_cast(lowerCountsBuffer.data()); + auto leafCounts = reinterpret_cast(leafCountsBuffer.data()); using Op = util::morphology::cuda::EnumerateNodesFunctor; util::cuda::operatorKernel <<>> (deviceUpperMasks(), deviceLowerMasks(), lowerCounts, leafCounts); - mUpperOffsets = nanovdb::cuda::DeviceBuffer::create((processedTileCount+1)*sizeof(uint32_t), nullptr, device, stream); - mLowerOffsets = nanovdb::cuda::DeviceBuffer::create((size+1)*sizeof(uint32_t), nullptr, device, stream); - mLeafOffsets = nanovdb::cuda::DeviceBuffer::create((size+1)*sizeof(uint32_t), nullptr, device, stream); + 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); - cudaCheck(cudaMemsetAsync(mLowerOffsets.deviceData(), 0, sizeof(uint32_t), stream)); + cudaCheck(cudaMemsetAsync(mLowerOffsets.data(), 0, sizeof(uint32_t), stream)); CALL_CUBS(DeviceScan::InclusiveSum, - static_cast(lowerCountsBuffer.deviceData()), - static_cast(mLowerOffsets.deviceData())+1, + reinterpret_cast(lowerCountsBuffer.data()), + reinterpret_cast(mLowerOffsets.data())+1, size); - cudaCheck(cudaMemcpyAsync(&data()->nodeCount[1], static_cast(mLowerOffsets.deviceData())+size, sizeof(uint32_t), cudaMemcpyDeviceToHost, stream)); + cudaCheck(cudaMemcpyAsync(&data()->nodeCount[1], reinterpret_cast(mLowerOffsets.data())+size, sizeof(uint32_t), cudaMemcpyDeviceToHost, stream)); - cudaCheck(cudaMemsetAsync(mLeafOffsets.deviceData(), 0, sizeof(uint32_t), stream)); + cudaCheck(cudaMemsetAsync(mLeafOffsets.data(), 0, sizeof(uint32_t), stream)); CALL_CUBS(DeviceScan::InclusiveSum, - static_cast(leafCountsBuffer.deviceData()), - static_cast(mLeafOffsets.deviceData())+1, + reinterpret_cast(leafCountsBuffer.data()), + reinterpret_cast(mLeafOffsets.data())+1, size); - cudaCheck(cudaMemcpyAsync(&data()->nodeCount[0], static_cast(mLeafOffsets.deviceData())+size, sizeof(uint32_t), cudaMemcpyDeviceToHost, stream)); + cudaCheck(cudaMemcpyAsync(&data()->nodeCount[0], reinterpret_cast(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; }, - static_cast(mLowerOffsets.deviceData()), - static_cast(upperCountsBuffer.deviceData())); + reinterpret_cast(mLowerOffsets.data()), + reinterpret_cast(upperCountsBuffer.data())); - cudaCheck(cudaMemsetAsync( mUpperOffsets.deviceData(), 0, sizeof(uint32_t), stream)); + cudaCheck(cudaMemsetAsync( mUpperOffsets.data(), 0, sizeof(uint32_t), stream)); CALL_CUBS(DeviceScan::InclusiveSum, - static_cast(upperCountsBuffer.deviceData()), - static_cast(mUpperOffsets.deviceData())+1, + reinterpret_cast(upperCountsBuffer.data()), + reinterpret_cast(mUpperOffsets.data())+1, processedTileCount); - cudaCheck(cudaMemcpyAsync(&data()->nodeCount[2], static_cast(mUpperOffsets.deviceData())+processedTileCount, sizeof(uint32_t), cudaMemcpyDeviceToHost, stream)); -}// TopologyBuilder::countNodes + cudaCheck(cudaMemcpyAsync(&data()->nodeCount[2], reinterpret_cast(mUpperOffsets.data())+processedTileCount, sizeof(uint32_t), cudaMemcpyDeviceToHost, stream)); +}// TopologyBuilder::countNodes //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template +template template -BufferT TopologyBuilder::getBuffer(const BufferT &pool, cudaStream_t stream) +BufferT TopologyBuilder::getBuffer(const BufferT &pool, cudaStream_t stream) { // Allocates a device buffer for the destination grid, once the topology/size of the tree is known data()->grid = 0;// grid is always stored at the start of the buffer! @@ -226,11 +247,11 @@ BufferT TopologyBuilder::getBuffer(const BufferT &pool, cudaStream_t str data()->d_bufferPtr = buffer.deviceData(); 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 = static_cast(mUpperOffsets.deviceData()); + data()->d_upperOffsets = reinterpret_cast(mUpperOffsets.data()); mData.deviceUpload(device, stream, false); return buffer; -}// TopologyBuilder::getBuffer +}// TopologyBuilder::getBuffer //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- @@ -240,7 +261,7 @@ template struct BuildGridTreeRootFunctor { __device__ - void operator()(size_t, typename TopologyBuilder::Data *d_data) { + void operator()(size_t, TopologyBuilderData *d_data) { // process Root auto &root = d_data->getRoot(); @@ -320,7 +341,7 @@ struct InitGridTreeRootFunctor Map map; // transform to embed in the output grid __device__ - void operator()(size_t, typename TopologyBuilder::Data *d_data) { + void operator()(size_t, TopologyBuilderData *d_data) { // process Root (identical to BuildGridTreeRootFunctor) auto &root = d_data->getRoot(); @@ -389,7 +410,7 @@ template struct BuildUpperNodesFunctor { __device__ - void operator()(size_t processedTileID, typename TopologyBuilder::Data *d_data, NanoRoot *d_processedRoot) { + void operator()(size_t processedTileID, TopologyBuilderData *d_data, NanoRoot *d_processedRoot) { uint32_t tileID = d_data->d_upperOffsets[processedTileID]; if (tileID != d_data->d_upperOffsets[processedTileID+1]) // if the offsets are the same, this was a speculatively introduced tile which was not necessary { @@ -406,8 +427,8 @@ struct BuildUpperNodesFunctor }// namespace topology::detail -template -inline void TopologyBuilder::processUpperNodes(cudaStream_t stream) +template +inline void TopologyBuilder::processUpperNodes(cudaStream_t stream) { // Connect all newly allocated upper nodes to their respective tiles // Also fill in any necessary part of the preamble (in InternalData) of upper nodes @@ -418,12 +439,12 @@ inline void TopologyBuilder::processUpperNodes(cudaStream_t stream) processedTileCount, topology::detail::BuildUpperNodesFunctor(), deviceData(), deviceProcessedRoot()); cudaCheckError(); } -}// TopologyBuilder::processUpperNodes +}// TopologyBuilder::processUpperNodes //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template -inline void TopologyBuilder::processLowerNodes(cudaStream_t stream) +template +inline void TopologyBuilder::processLowerNodes(cudaStream_t stream) { // 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 @@ -434,31 +455,31 @@ inline void TopologyBuilder::processLowerNodes(cudaStream_t stream) int device = 0; cudaGetDevice(&device); std::size_t lowerCount = data()->nodeCount[1]; - mLowerParents = nanovdb::cuda::DeviceBuffer::create(lowerCount*sizeof(uint32_t), nullptr, device, stream); + mLowerParents = ScratchT(stream, *mResource, lowerCount*sizeof(uint32_t), nanovdb::cuda::noInit); std::size_t leafCount = data()->nodeCount[0]; - mLeafParents = nanovdb::cuda::DeviceBuffer::create(leafCount*sizeof(uint32_t), nullptr, device, stream); + mLeafParents = ScratchT(stream, *mResource, leafCount*sizeof(uint32_t), nanovdb::cuda::noInit); using Op = util::morphology::cuda::ProcessLowerNodesFunctor; util::cuda::operatorKernel <<>>( deviceUpperMasks(), deviceLowerMasks(), - static_cast(mUpperOffsets.deviceData()), - static_cast(mLowerOffsets.deviceData()), - static_cast(mLeafOffsets.deviceData()), + reinterpret_cast(mUpperOffsets.data()), + reinterpret_cast(mLowerOffsets.data()), + reinterpret_cast(mLeafOffsets.data()), static_cast(data()->d_bufferPtr), - static_cast(mLowerParents.deviceData()), - static_cast(mLeafParents.deviceData()) + reinterpret_cast(mLowerParents.data()), + reinterpret_cast(mLeafParents.data()) ); cudaCheckError(); } mProcessedRoot.clear(stream); - mUpperMasks.clear(stream); - mLowerMasks.clear(stream); - mLowerOffsets.clear(stream); - mLeafOffsets.clear(stream); -}// TopologyBuilder::processLowerNodes + mUpperMasks.destroy(); + mLowerMasks.destroy(); + mLowerOffsets.destroy(); + mLeafOffsets.destroy(); +}// TopologyBuilder::processLowerNodes //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- @@ -468,7 +489,7 @@ template struct UpdateLeafVoxelCountsAndPrefixSumFunctor { __device__ - void operator()(size_t leafID, typename TopologyBuilder::Data *d_data, uint64_t *d_voxelCounts) { + void operator()(size_t leafID, TopologyBuilderData *d_data, uint64_t *d_voxelCounts) { auto &leaf = d_data->getGrid().tree().template getFirstNode<0>()[leafID]; const uint64_t *w = leaf.mValueMask.words(); uint64_t prefixSum = 0, sum = util::countOn(*w++); @@ -485,32 +506,32 @@ template struct UpdateLeafVoxelOffsetsFunctor { __device__ - void operator()(size_t leafID, typename TopologyBuilder::Data *d_data, uint64_t *d_voxelOffsets) { + void operator()(size_t leafID, TopologyBuilderData *d_data, uint64_t *d_voxelOffsets) { auto &leaf = d_data->getGrid().tree().template getFirstNode<0>()[leafID]; leaf.mOffset = d_voxelOffsets[leafID]+1; } }; }// namespace topology::detail -template -inline void TopologyBuilder::processLeafOffsets(cudaStream_t stream) +template +inline void TopologyBuilder::processLeafOffsets(cudaStream_t stream) { int device = 0; cudaGetDevice(&device); std::size_t leafCount = data()->nodeCount[0]; if (leafCount) { // Unless output grid is empty - mVoxelOffsets = nanovdb::cuda::DeviceBuffer::create((leafCount+1)*sizeof(uint64_t), nullptr, device, stream); - cudaCheck(cudaMemsetAsync(mVoxelOffsets.deviceData(), 0, sizeof(uint64_t), stream)); + mVoxelOffsets = ScratchT(stream, *mResource, (leafCount+1)*sizeof(uint64_t), nanovdb::cuda::noInit); + cudaCheck(cudaMemsetAsync(mVoxelOffsets.data(), 0, sizeof(uint64_t), stream)); util::cuda::lambdaKernel<<>>( - leafCount, topology::detail::UpdateLeafVoxelCountsAndPrefixSumFunctor(), deviceData(), static_cast(mVoxelOffsets.deviceData())+1); + leafCount, topology::detail::UpdateLeafVoxelCountsAndPrefixSumFunctor(), deviceData(), reinterpret_cast(mVoxelOffsets.data())+1); CALL_CUBS(DeviceScan::InclusiveSum, - static_cast(mVoxelOffsets.deviceData())+1, - static_cast(mVoxelOffsets.deviceData())+1, + reinterpret_cast(mVoxelOffsets.data())+1, + reinterpret_cast(mVoxelOffsets.data())+1, leafCount); util::cuda::lambdaKernel<<>>( - leafCount, topology::detail::UpdateLeafVoxelOffsetsFunctor(), deviceData(), static_cast(mVoxelOffsets.deviceData())); + leafCount, topology::detail::UpdateLeafVoxelOffsetsFunctor(), deviceData(), reinterpret_cast(mVoxelOffsets.data())); } -}// TopologyBuilder::processLeafOffsets +}// TopologyBuilder::processLeafOffsets //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- @@ -527,7 +548,7 @@ template struct UpdateAndPropagateLeafBBoxFunctor { __device__ - void operator()(size_t tid, typename TopologyBuilder::Data *d_data, const uint32_t* leafParents) { + void operator()(size_t tid, TopologyBuilderData *d_data, const uint32_t* leafParents) { auto &lower = d_data->getLower(leafParents[tid]); auto &leaf = d_data->getLeaf(tid); leaf.updateBBox(); @@ -539,7 +560,7 @@ template struct PropagateLowerBBoxFunctor { __device__ - void operator()(size_t tid, typename TopologyBuilder::Data *d_data, const uint32_t* lowerParents) { + void operator()(size_t tid, TopologyBuilderData *d_data, const uint32_t* lowerParents) { auto &upper = d_data->getUpper(lowerParents[tid]); auto &lower = d_data->getLower(tid); upper.mBBox.expandAtomic(lower.bbox()); } @@ -549,7 +570,7 @@ template struct PropagateUpperBBoxFunctor { __device__ - void operator()(size_t tid, typename TopologyBuilder::Data *d_data) { + void operator()(size_t tid, TopologyBuilderData *d_data) { d_data->getRoot().mBBox.expandAtomic(d_data->getUpper(tid).bbox()); } }; @@ -558,7 +579,7 @@ template struct UpdateRootWorldBBoxFunctor { __device__ - void operator()(size_t tid, typename TopologyBuilder::Data *d_data) { + void operator()(size_t tid, TopologyBuilderData *d_data) { // TODO: check that the correct semantics are followed in this transformation auto BBox = d_data->getRoot().mBBox; BBox.max() += 1; @@ -569,8 +590,8 @@ struct UpdateRootWorldBBoxFunctor }// namespace topology::detail -template -inline void TopologyBuilder::processBBox(cudaStream_t stream) +template +inline void TopologyBuilder::processBBox(cudaStream_t stream) { if (data()->nodeCount[0] == 0) return; // Output grid is empty; retain empty bounding box @@ -578,14 +599,14 @@ 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(), static_cast(mLeafParents.deviceData())); - mLeafParents.clear(stream); + data()->nodeCount[0], topology::detail::UpdateAndPropagateLeafBBoxFunctor(), deviceData(), reinterpret_cast(mLeafParents.data())); + mLeafParents.destroy(); cudaCheckError(); // propagate bbox from lower -> upper/parent node util::cuda::lambdaKernel<<nodeCount[1]), mNumThreads, 0, stream>>>( - data()->nodeCount[1], topology::detail::PropagateLowerBBoxFunctor(), deviceData(), static_cast(mLowerParents.deviceData())); - mLowerParents.clear(stream); + data()->nodeCount[1], topology::detail::PropagateLowerBBoxFunctor(), deviceData(), reinterpret_cast(mLowerParents.data())); + mLowerParents.destroy(); cudaCheckError(); // propagate bbox from upper -> root/parent node @@ -595,7 +616,7 @@ inline void TopologyBuilder::processBBox(cudaStream_t stream) // update the world-bbox in the root node util::cuda::lambdaKernel<<<1, 1, 0, stream>>>(1, topology::detail::UpdateRootWorldBBoxFunctor(), deviceData()); cudaCheckError(); -}// TopologyBuilder::processBBox +}// TopologyBuilder::processBBox //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- @@ -605,7 +626,7 @@ template struct PostProcessGridTreeFunctor { __device__ - void operator()(size_t tid, typename TopologyBuilder::Data *d_data, uint64_t* d_voxelOffsets) { + void operator()(size_t tid, TopologyBuilderData *d_data, uint64_t* d_voxelOffsets) { auto& grid = d_data->getGrid(); auto& tree = grid.tree(); auto leafCount = tree.mNodeCount[0]; @@ -616,17 +637,17 @@ struct PostProcessGridTreeFunctor }// namespace topology::detail -template -inline void TopologyBuilder::postProcessGridTree(cudaStream_t stream) +template +inline void TopologyBuilder::postProcessGridTree(cudaStream_t stream) { // 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(), static_cast(mVoxelOffsets.deviceData())); + util::cuda::lambdaKernel<<<1, 1, 0, stream>>>(1, topology::detail::PostProcessGridTreeFunctor(), deviceData(), reinterpret_cast(mVoxelOffsets.data())); cudaCheckError(); - mVoxelOffsets.clear(stream); + mVoxelOffsets.destroy(); tools::cuda::updateChecksum((GridData*)data()->d_bufferPtr, mChecksum, stream); -}// TopologyBuilder::postProcessGridTree +}// TopologyBuilder::postProcessGridTree //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- From 1e693339cbc99d9dd7cc6f16e8cf8f1f8a5ce01a Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 5 Aug 2026 01:52:07 +0000 Subject: [PATCH 03/17] NanoVDB: add SyncFromAsync, and give MeshToGrid a resource seam A stream-ordered resource must also model the synchronous concept, which means writing four methods where two would do. The synchronous pair is not a bare delegate -- memory from allocate must be usable on any stream when it returns, so the null-stream allocation has to be synchronized first -- and omitting that yields memory which satisfies the concept but is not actually synchronous. Put it in one place rather than leaving each author to rediscover it. The two resources in TestMemoryResource are the first users, and were already wrong in exactly that way: they provide only the async pair, so they never modelled is_async_resource. TempPool duck-typed and never checked, so nothing caught it. MeshToGrid was the last builder allocating from a hard-wired DeviceResource, through TempDevicePool. Give it a ResourceT parameter and thread it into both its TopologyBuilder and its pool. As with Data in the builder, BoxTrianglePair is hoisted out of the class: it does not depend on the resource, and leaving it nested would give every ResourceT its own incompatible type for the device functors to name. Co-Authored-By: Claude Opus 5 (1M context) Signed-off-by: Mark Harris --- nanovdb/nanovdb/cuda/DeviceResource.h | 45 +++++++++ nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh | 99 ++++++++++--------- nanovdb/nanovdb/unittest/TestBuffer.cu | 48 +++++++++ .../nanovdb/unittest/TestMemoryResource.cu | 4 +- 4 files changed, 147 insertions(+), 49 deletions(-) diff --git a/nanovdb/nanovdb/cuda/DeviceResource.h b/nanovdb/nanovdb/cuda/DeviceResource.h index a18ee789b6..fee39332bd 100644 --- a/nanovdb/nanovdb/cuda/DeviceResource.h +++ b/nanovdb/nanovdb/cuda/DeviceResource.h @@ -145,6 +145,51 @@ struct is_resource().deallocate(std::declval(), size_t{0}, size_t{0}))>> : std::true_type {}; +/// @brief CRTP base supplying the synchronous half of the resource concept in +/// terms of the stream-ordered half, so a custom stream-ordered resource +/// only has to write allocate_async and deallocate_async. +/// @tparam Derived the resource deriving from this base +/// @details A stream-ordered resource must also model the synchronous concept +/// (is_async_resource implies is_resource), which means writing four +/// methods where two would do. The synchronous pair is not a bare +/// delegate: memory from allocate must be usable immediately on any +/// stream, so the null-stream allocation has to be synchronized before +/// it is returned. Omitting that synchronization yields memory that +/// satisfies the concept but is not actually synchronous -- a race +/// rather than a compile error -- so it lives here rather than being +/// rewritten per resource. +/// @code +/// struct MyResource : nanovdb::cuda::SyncFromAsync { +/// static constexpr size_t DEFAULT_ALIGNMENT = 256; +/// void* allocate_async(size_t bytes, size_t alignment, cudaStream_t stream); +/// void deallocate_async(void* p, size_t bytes, size_t alignment, cudaStream_t stream); +/// }; +/// @endcode +template +struct SyncFromAsync +{ + /// @brief Allocates @c bytes usable on any stream when this returns. + /// @param bytes number of bytes to allocate + /// @param alignment requested alignment + void* allocate(size_t bytes, size_t alignment) + { + void* p = static_cast(*this).allocate_async(bytes, alignment, cudaStream_t{0}); + cudaCheck(cudaStreamSynchronize(cudaStream_t{0})); + return p; + } + + /// @brief Frees @c p on the null stream. + /// @param p pointer previously returned by allocate + /// @param bytes size passed to the matching allocate + /// @param alignment alignment passed to the matching allocate + /// @note No synchronization here: the synchronous concept's contract is + /// that the memory is already quiescent when deallocate is called. + void deallocate(void* p, size_t bytes, size_t alignment) + { + static_cast(*this).deallocate_async(p, bytes, alignment, cudaStream_t{0}); + } +}; + } } // namespace nanovdb::cuda diff --git a/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh b/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh index d143793dce..b93db43d9b 100644 --- a/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh @@ -48,7 +48,15 @@ struct Triangle { __hostdev__ nanovdb::Vec3f& operator[](int i) { return v[i]; } }; -template +/// @brief Pairing of a leaf-node origin with a triangle id. Independent of the +/// resource the converter allocates from, so it lives outside MeshToGrid +/// and stays one type across every ResourceT instantiation. +struct alignas(16) MeshToGridBoxTrianglePair { // sizeof = 16B + nanovdb::Coord origin; // 12B + uint32_t triangleID; // 4B +}; + +template class MeshToGrid { using PointT = nanovdb::Vec3f; @@ -62,10 +70,7 @@ class MeshToGrid using LeafT = NanoLeaf; public: - struct alignas(16) BoxTrianglePair { // sizeof(BoxTrianglePair) = 16B - nanovdb::Coord origin; // 12B - uint32_t triangleID; // 4B - }; + using BoxTrianglePair = MeshToGridBoxTrianglePair; /// @brief Constructor /// @param devicePoints Vertex list for input triangle surface @@ -79,10 +84,11 @@ public: const nanovdb::Vec3i *deviceTriangles, const uint32_t triangleCount, const nanovdb::Map map = nanovdb::Map(), - cudaStream_t stream = 0 + cudaStream_t stream = 0, + ResourceT& resource = nanovdb::cuda::default_resource() ) - : mStream(stream), mTimer(stream), mBuilder(stream), mDevicePoints(devicePoints), mPointCount(pointCount), - mDeviceTriangles(deviceTriangles), mTriangleCount(triangleCount), mMap(map) + : mStream(stream), mTimer(stream), mBuilder(stream, resource), mDevicePoints(devicePoints), mPointCount(pointCount), + mDeviceTriangles(deviceTriangles), mTriangleCount(triangleCount), mMap(map), mTempDevicePool(resource) {} /// @brief Toggle on and off verbose mode @@ -155,7 +161,7 @@ private: static constexpr unsigned int mNumThreads = 128;// for kernels spawned via lambdaKernel (others may specialize) static unsigned int numBlocks(unsigned int n) {return (n + mNumThreads - 1) / mNumThreads;} - TopologyBuilder mBuilder; + TopologyBuilder mBuilder; cudaStream_t mStream{0}; std::string mGridName; util::cuda::Timer mTimer; @@ -178,8 +184,8 @@ private: auto deviceBoxTrianglePairs() { return static_cast(mBoxTrianglePairsBuffer.deviceData()); } auto deviceUniqueRootOrigins() const { return static_cast(mUniqueRootOriginsBuffer.deviceData()); } - nanovdb::cuda::TempDevicePool mTempDevicePool; -}; // tools::cuda::MeshToGrid + nanovdb::cuda::TempPool mTempDevicePool; +}; // tools::cuda::MeshToGrid //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- @@ -200,10 +206,9 @@ private: //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template +template template -GridHandle -MeshToGrid::getHandle(const BufferT &buffer) +GridHandle MeshToGrid::getHandle(const BufferT &buffer) { cudaStreamSynchronize(mStream); @@ -310,7 +315,7 @@ MeshToGrid::getHandle(const BufferT &buffer) } if (mVerbose==1) mTimer.stop(); return handle; -} // MeshToGrid::getHandle +} // MeshToGrid::getHandle //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- @@ -335,8 +340,8 @@ struct TransformTrianglesFunctor } // namespace topology::detail -template -void MeshToGrid::transformTriangles() +template +void MeshToGrid::transformTriangles() { if (mTriangleCount == 0) return; @@ -355,7 +360,7 @@ void MeshToGrid::transformTriangles() cudaCheckError(); -} // MeshToGrid::transformTriangles +} // MeshToGrid::transformTriangles //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- @@ -416,7 +421,7 @@ struct CountRootBoxesFunctor template struct ScatterRootTrianglePairsFunctor { - using PairT = typename MeshToGrid::BoxTrianglePair; + using PairT = MeshToGridBoxTrianglePair; const Triangle* dXformedTriangles; const uint64_t* dOffsets; @@ -471,8 +476,8 @@ struct ScatterRootTrianglePairsFunctor } // namespace topology::detail -template -void MeshToGrid::processRootTrianglePairs() +template +void MeshToGrid::processRootTrianglePairs() { if (mTriangleCount == 0) { mBoxTrianglePairCount = 0; return; } @@ -512,7 +517,7 @@ void MeshToGrid::processRootTrianglePairs() // Pass 3: Re-enumerate intersections of (padded) root boxes and triangles, and scatter to allocated list mBoxTrianglePairsBuffer = nanovdb::cuda::DeviceBuffer::create( - mBoxTrianglePairCount * sizeof(typename MeshToGrid::BoxTrianglePair), nullptr, device, mStream); + mBoxTrianglePairCount * sizeof(MeshToGridBoxTrianglePair), nullptr, device, mStream); if (mBoxTrianglePairsBuffer.deviceData() == nullptr) throw std::runtime_error("Failed to allocate pairs buffer"); util::cuda::lambdaKernel<<>>( @@ -525,7 +530,7 @@ void MeshToGrid::processRootTrianglePairs() } ); -} // MeshToGrid::processRootTrianglePairs +} // MeshToGrid::processRootTrianglePairs //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- @@ -545,7 +550,7 @@ namespace topology::detail { template struct ScatterChildPairsFunctor { - using PairT = typename MeshToGrid::BoxTrianglePair; + using PairT = MeshToGridBoxTrianglePair; const PairT* dParents; const nanovdb::Mask<3>* dMasks; @@ -657,7 +662,7 @@ __device__ inline bool testTriangleAABB( template __global__ void evaluateAndCountSubBoxesKernel( - const typename MeshToGrid::BoxTrianglePair* dParents, + const MeshToGridBoxTrianglePair* dParents, const Triangle* dXformedTriangles, nanovdb::Mask<3>* dMasks, uint64_t* dCounts, @@ -746,7 +751,7 @@ __device__ inline nanovdb::Coord keyToCoord(uint64_t key) template struct EncodeRootOriginsFunctor { - const typename MeshToGrid::BoxTrianglePair* dPairs; + const MeshToGridBoxTrianglePair* dPairs; uint64_t* dKeys; __device__ void operator()(size_t i) const { dKeys[i] = coordToKey(dPairs[i].origin); } @@ -763,8 +768,8 @@ struct DecodeRootOriginsFunctor } // namespace topology::detail -template -void MeshToGrid::enumerateRootTiles() +template +void MeshToGrid::enumerateRootTiles() { if (mBoxTrianglePairCount == 0) return; @@ -816,12 +821,12 @@ void MeshToGrid::enumerateRootTiles() ); cudaCheckError(); -} // MeshToGrid::enumerateRootTiles +} // MeshToGrid::enumerateRootTiles //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template -void MeshToGrid::buildRasterizedRoot() +template +void MeshToGrid::buildRasterizedRoot() { int device = 0; cudaGetDevice(&device); @@ -850,12 +855,12 @@ void MeshToGrid::buildRasterizedRoot() mBuilder.mProcessedRoot.deviceUpload(device, mStream, false); mUniqueRootOriginsBuffer.clear(mStream); } -} // MeshToGrid::buildRasterizedRoot +} // MeshToGrid::buildRasterizedRoot //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template -void MeshToGrid::rasterizeInternalNodes() +template +void MeshToGrid::rasterizeInternalNodes() { if (mBoxTrianglePairCount == 0) return; @@ -870,12 +875,12 @@ void MeshToGrid::rasterizeInternalNodes() ); cudaCheckError(); -} // MeshToGrid::rasterizeInternalNodes +} // MeshToGrid::rasterizeInternalNodes //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template -void MeshToGrid::processGridTreeRoot() +template +void MeshToGrid::processGridTreeRoot() { // Initialize grid/tree/root metadata from scratch using the provided map. // InitGridTreeRootFunctor sets all GridData fields explicitly (magic, version, @@ -892,12 +897,12 @@ void MeshToGrid::processGridTreeRoot() cudaCheck(cudaMemsetAsync(dst, 0, GridData::MaxNameSize, mStream)); } -} // MeshToGrid::processGridTreeRoot +} // MeshToGrid::processGridTreeRoot //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template -void MeshToGrid::rasterizeLeafNodes() +template +void MeshToGrid::rasterizeLeafNodes() { if (mBoxTrianglePairCount == 0) return; @@ -908,12 +913,12 @@ void MeshToGrid::rasterizeLeafNodes() &mBuilder.data()->getGrid(), mBandWidth * mBandWidth }); cudaCheckError(); -} // MeshToGrid::rasterizeLeafNodes +} // MeshToGrid::rasterizeLeafNodes //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template -void MeshToGrid::processLeafTrianglePairs() +template +void MeshToGrid::processLeafTrianglePairs() { if (mBoxTrianglePairCount == 0) return; @@ -999,7 +1004,7 @@ void MeshToGrid::processLeafTrianglePairs() scale /= 8; } -} // MeshToGrid::processLeafTrianglePairs +} // MeshToGrid::processLeafTrianglePairs //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- @@ -1038,10 +1043,10 @@ struct FinalizeSidecarFunctor //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template +template template std::pair, SidecarBufferT> -MeshToGrid::getHandleAndUDF(const GridBufferT& buffer, const SidecarBufferT&) +MeshToGrid::getHandleAndUDF(const GridBufferT& buffer, const SidecarBufferT&) { cudaStreamSynchronize(mStream); @@ -1172,7 +1177,7 @@ MeshToGrid::getHandleAndUDF(const GridBufferT& buffer, const SidecarBuff cudaStreamSynchronize(mStream); return { std::move(handle), std::move(sidecarBuffer) }; -} // MeshToGrid::getHandleAndUDF +} // MeshToGrid::getHandleAndUDF //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- diff --git a/nanovdb/nanovdb/unittest/TestBuffer.cu b/nanovdb/nanovdb/unittest/TestBuffer.cu index 40be6878a6..f6200eb73b 100644 --- a/nanovdb/nanovdb/unittest/TestBuffer.cu +++ b/nanovdb/nanovdb/unittest/TestBuffer.cu @@ -258,6 +258,54 @@ TEST(TestBuffer, ClearFreesAndEmpties) ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess); } +// A custom stream-ordered resource written the short way: two methods plus the +// mixin, rather than four. +struct MixinResource : nanovdb::cuda::SyncFromAsync +{ + static constexpr size_t DEFAULT_ALIGNMENT = nanovdb::cuda::DeviceResource::DEFAULT_ALIGNMENT; + Counters* counters = nullptr; + void* allocate_async(size_t bytes, size_t alignment, cudaStream_t stream) { + void* p = nanovdb::cuda::DeviceResource{}.allocate_async(bytes, alignment, stream); + if (p) { ++counters->allocs; counters->allocBytes = bytes; } + return p; + } + void deallocate_async(void* p, size_t bytes, size_t alignment, cudaStream_t stream) { + if (p) { ++counters->deallocs; counters->deallocBytes = bytes; } + nanovdb::cuda::DeviceResource{}.deallocate_async(p, bytes, alignment, stream); + } +}; + +// The mixin supplies the synchronous half, so both concepts are satisfied. +static_assert(nanovdb::cuda::is_async_resource::value, + "SyncFromAsync user must still model AsyncResource"); +static_assert(nanovdb::cuda::is_resource::value, + "SyncFromAsync must supply the synchronous half of the concept"); + +TEST(TestBuffer, SyncFromAsyncSuppliesTheSynchronousPair) +{ + Counters c; + MixinResource r{{}, &c}; + // the inherited synchronous pair routes through the derived async methods + void* p = r.allocate(1024, MixinResource::DEFAULT_ALIGNMENT); + ASSERT_NE(p, nullptr); + EXPECT_EQ(c.allocs, 1); + r.deallocate(p, 1024, MixinResource::DEFAULT_ALIGNMENT); + EXPECT_EQ(c.deallocs, 1); + ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess); +} + +TEST(TestBuffer, BufferWorksOverAMixinResource) +{ + Counters c; + { + nanovdb::cuda::Buffer buf(0, MixinResource{{}, &c}, 64, nanovdb::cuda::noInit); + EXPECT_EQ(c.allocs, 1); + EXPECT_NE(buf.data(), nullptr); + } + EXPECT_EQ(c.deallocs, 1); + ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess); +} + TEST(TestBuffer, DestroyIsTheSpellingClearDelegatesTo) { Counters c; diff --git a/nanovdb/nanovdb/unittest/TestMemoryResource.cu b/nanovdb/nanovdb/unittest/TestMemoryResource.cu index 136e409323..ed5591055b 100644 --- a/nanovdb/nanovdb/unittest/TestMemoryResource.cu +++ b/nanovdb/nanovdb/unittest/TestMemoryResource.cu @@ -25,7 +25,7 @@ namespace { /// @brief Resource that counts (non-null) allocations and deallocations so /// leaks can be asserted. Delegates the actual work to DeviceResource. -struct CountingResource +struct CountingResource : nanovdb::cuda::SyncFromAsync { static constexpr size_t DEFAULT_ALIGNMENT = nanovdb::cuda::DeviceResource::DEFAULT_ALIGNMENT; int allocs = 0; @@ -43,7 +43,7 @@ struct CountingResource /// @brief Resource that records the stream of every allocation/deallocation, /// to verify stream-ordered teardown. Delegates work to DeviceResource. -struct StreamRecordingResource +struct StreamRecordingResource : nanovdb::cuda::SyncFromAsync { static constexpr size_t DEFAULT_ALIGNMENT = nanovdb::cuda::DeviceResource::DEFAULT_ALIGNMENT; std::vector allocStreams; From 35d61dd52bca00fddc40a74931858433f6659677 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 5 Aug 2026 03:14:37 +0000 Subject: [PATCH 04/17] NanoVDB: add ResourceRef and route TempPool's scratch through cuda::Buffer Buffer holds its resource by value, matching cuda::buffer -- whose model this completes: in CCCL the ownership semantics are selected by what is placed in the by-value slot, an owning any_resource or a borrowing resource_ref. We adopted the slot without the borrowing type, so a container like TempPool, whose contract is a non-owning pointer to a possibly stateful resource, had no way to hold a Buffer without copying that resource and stranding its state. ResourceRef is the missing piece: a non-owning reference that is itself a resource, so copying the ref shares the underlying instance. Its async methods exist only when R models AsyncResource, so a ref over a synchronous resource does not misreport its tier, and two refs compare equal exactly when they reference the same resource. TempPool now keeps its bytes in a Buffer>: same resource contract, same stream retention, same discard-on-growth reallocation, but the block is freed by ownership rather than by hand. The TempPool unit tests, which assert traffic against the caller's own resource instance, pass unchanged. Co-Authored-By: Claude Opus 5 (1M context) Signed-off-by: Mark Harris --- nanovdb/nanovdb/cuda/DeviceResource.h | 53 ++++++++++++++++++++++++++ nanovdb/nanovdb/cuda/TempPool.h | 51 ++++++++++++++----------- nanovdb/nanovdb/unittest/TestBuffer.cu | 48 +++++++++++++++++++++++ pendingchanges/nanovdb.txt | 1 + 4 files changed, 130 insertions(+), 23 deletions(-) diff --git a/nanovdb/nanovdb/cuda/DeviceResource.h b/nanovdb/nanovdb/cuda/DeviceResource.h index fee39332bd..0fcd3d0328 100644 --- a/nanovdb/nanovdb/cuda/DeviceResource.h +++ b/nanovdb/nanovdb/cuda/DeviceResource.h @@ -190,6 +190,59 @@ struct SyncFromAsync } }; +/// @brief Non-owning reference to a memory resource that is itself a resource: +/// copying the ref shares the underlying resource rather than copying it. +/// @tparam R the referenced resource type +/// @details Types that hold their resource by value -- cuda::Buffer, matching +/// cuda::buffer -- select their ownership semantics by what is placed +/// in that slot: a concrete resource is owned as a copy, while a +/// ResourceRef borrows. This is the same division cuda::mr draws +/// between any_resource (owning) and resource_ref (borrowing), and the +/// same shape as std::pmr::polymorphic_allocator over memory_resource*. +/// Use it when a resource is stateful or long-lived and a container +/// must allocate through *that* instance rather than a copy of it. +/// @warning The referenced resource must outlive every use of this ref and of +/// all copies of it, including any container holding one. +template +struct ResourceRef +{ + static_assert(is_async_resource::value || is_resource::value, + "ResourceRef requires R to model the AsyncResource or the Resource concept"); + + static constexpr size_t DEFAULT_ALIGNMENT = R::DEFAULT_ALIGNMENT; + + /// @brief Constructs a ref borrowing @c resource. + /// @param resource resource to allocate from; must outlive this ref + ResourceRef(R& resource) : mResource(&resource) {} + + /// @{ + /// @brief Stream-ordered pair, present only when @c R models AsyncResource, + /// so a ref over a synchronous resource does not misreport its tier. + template::value, int> = 0> + void* allocate_async(size_t bytes, size_t alignment, cudaStream_t stream) + { + return mResource->allocate_async(bytes, alignment, stream); + } + template::value, int> = 0> + void deallocate_async(void* p, size_t bytes, size_t alignment, cudaStream_t stream) + { + mResource->deallocate_async(p, bytes, alignment, stream); + } + /// @} + + /// @brief Synchronous pair, forwarding to the referenced resource. + void* allocate(size_t bytes, size_t alignment) { return mResource->allocate(bytes, alignment); } + void deallocate(void* p, size_t bytes, size_t alignment) { mResource->deallocate(p, bytes, alignment); } + + /// @brief Two refs compare equal iff they reference the same resource, i.e. + /// memory allocated through one may be deallocated through the other. + friend bool operator==(ResourceRef lhs, ResourceRef rhs) { return lhs.mResource == rhs.mResource; } + friend bool operator!=(ResourceRef lhs, ResourceRef rhs) { return lhs.mResource != rhs.mResource; } + +private: + R* mResource; +};// ResourceRef + } } // namespace nanovdb::cuda diff --git a/nanovdb/nanovdb/cuda/TempPool.h b/nanovdb/nanovdb/cuda/TempPool.h index 63a97569d2..65317e5f3b 100644 --- a/nanovdb/nanovdb/cuda/TempPool.h +++ b/nanovdb/nanovdb/cuda/TempPool.h @@ -10,6 +10,7 @@ #ifndef NANOVDB_CUDA_TEMPPOOL_H_HAS_BEEN_INCLUDED #define NANOVDB_CUDA_TEMPPOOL_H_HAS_BEEN_INCLUDED +#include #include #include @@ -21,31 +22,33 @@ namespace cuda { template class TempPool { + static_assert(is_async_resource::value, + "TempPool allocates stream-ordered scratch and requires an AsyncResource"); + // The buffer borrows the pool's resource through a ResourceRef rather than + // copying it, preserving the pool's contract that all traffic reaches the + // caller's resource instance (which may be stateful). + using BufferT = Buffer>; public: /// @brief Default c-tor of an empty memory pool that uses the default /// instance of @c Resource for all allocations. - TempPool() : mResource(&default_resource()), mData(nullptr), mSize(0), mRequestedSize(0), mStream(nullptr) {} + TempPool() : TempPool(default_resource()) {} /// @brief C-tor of an empty memory pool that routes all allocations through /// the supplied @c Resource instance. /// @param resource resource instance to allocate from; must outlive this pool. - explicit TempPool(Resource& resource) : mResource(&resource), mData(nullptr), mSize(0), mRequestedSize(0), mStream(nullptr) {} - - /// @brief Destructor. Frees the managed memory on the stream of the most - /// recent reallocate(), so the stream-ordered free is ordered after - /// the work that used the memory (rather than on the null stream). - ~TempPool() { - mRequestedSize = 0; - mResource->deallocate_async(mData, mSize, Resource::DEFAULT_ALIGNMENT, mStream); - mData = nullptr; - mSize = 0; + explicit TempPool(Resource& resource) + : mResource(&resource) + , mBuffer(cudaStream_t{0}, ResourceRef(resource), 0, noInit) + { } /// @brief Returns a non-const void pointer to the data managed by this instance. - void* data() {return mData;} + void* data() {return mBuffer.data();} /// @brief Returns a non-const reference to the actual size of the data managed by this instance. + /// @note Returned by reference because cub's two-pass API takes the storage + /// size as a size_t&, so this cannot forward Buffer::size() by value. size_t& size() {return mSize;} /// @brief Returns a non-const reference to the requested size of the data managed by this instance. @@ -54,25 +57,27 @@ class TempPool { /// @brief Returns the stream that the managed memory was last (re)allocated on, /// i.e. the stream this pool will free on at destruction. - cudaStream_t stream() const {return mStream;} + cudaStream_t stream() const {return mBuffer.stream();} /// @brief Re-allocation of the data managed by this instance. Only has affect if the pool in empty or /// the requested memory is larger than the existing size. /// @param stream cuda stream used for asynchronous de-allocation and allocation. + /// @note Scratch is discarded, never resized: preserving a prefix of + /// temporary storage would be a wasted copy. void reallocate(cudaStream_t stream) { - if (!mData || mRequestedSize > mSize) { - mResource->deallocate_async(mData, mSize, Resource::DEFAULT_ALIGNMENT, stream); - mData = mResource->allocate_async(mRequestedSize, Resource::DEFAULT_ALIGNMENT, stream); - mSize = mRequestedSize; + if (mBuffer.empty() || mRequestedSize > mSize) { + mBuffer.destroy(stream);// free the outgrown block on this stream + mBuffer = BufferT(stream, ResourceRef(*mResource), mRequestedSize, noInit); + mSize = mBuffer.size(); + } else { + mBuffer.set_stream(stream);// retained so the d-tor frees on the most-recently-used stream } - mStream = stream;// retained so the destructor frees on the most-recently-used stream } private: - Resource *mResource; - void *mData; - size_t mSize; - size_t mRequestedSize; - cudaStream_t mStream; + Resource *mResource;// non-owning; must outlive this pool and its buffer + BufferT mBuffer; + size_t mSize{0}; + size_t mRequestedSize{0}; };// TempPool class using TempDevicePool = TempPool; diff --git a/nanovdb/nanovdb/unittest/TestBuffer.cu b/nanovdb/nanovdb/unittest/TestBuffer.cu index f6200eb73b..61fdcac549 100644 --- a/nanovdb/nanovdb/unittest/TestBuffer.cu +++ b/nanovdb/nanovdb/unittest/TestBuffer.cu @@ -306,6 +306,54 @@ TEST(TestBuffer, BufferWorksOverAMixinResource) ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess); } +// State held inline, not behind a pointer: copying such a resource strands its +// accounting, which is exactly what ResourceRef exists to avoid. +struct StatefulInlineResource : nanovdb::cuda::SyncFromAsync +{ + static constexpr size_t DEFAULT_ALIGNMENT = nanovdb::cuda::DeviceResource::DEFAULT_ALIGNMENT; + int allocs = 0, deallocs = 0; + void* allocate_async(size_t bytes, size_t alignment, cudaStream_t stream) { + ++allocs; return nanovdb::cuda::DeviceResource{}.allocate_async(bytes, alignment, stream); + } + void deallocate_async(void* p, size_t bytes, size_t alignment, cudaStream_t stream) { + ++deallocs; nanovdb::cuda::DeviceResource{}.deallocate_async(p, bytes, alignment, stream); + } +}; + +// A ref over an async resource models both tiers; over a synchronous-only +// resource it models only the synchronous one -- it must not misreport. +static_assert(nanovdb::cuda::is_async_resource>::value, + "ref over an async resource must model AsyncResource"); +static_assert(nanovdb::cuda::is_resource>::value, + "ref over an async resource must model Resource"); +static_assert(!nanovdb::cuda::is_async_resource>::value, + "ref over a synchronous resource must not claim AsyncResource"); +static_assert(nanovdb::cuda::is_resource>::value, + "ref over a synchronous resource must model Resource"); + +TEST(TestBuffer, ResourceRefSharesTheUnderlyingResource) +{ + StatefulInlineResource res; // the original; a by-value copy would strand these counters + { + nanovdb::cuda::Buffer> buf( + cudaStream_t{0}, nanovdb::cuda::ResourceRef(res), 1024, nanovdb::cuda::noInit); + EXPECT_NE(buf.data(), nullptr); + EXPECT_EQ(res.allocs, 1); // traffic reaches the original, not a copy + EXPECT_EQ(res.deallocs, 0); + } + ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess); + EXPECT_EQ(res.allocs, 1); + EXPECT_EQ(res.deallocs, 1); +} + +TEST(TestBuffer, ResourceRefEqualityIsIdentity) +{ + StatefulInlineResource a, b; + nanovdb::cuda::ResourceRef ra(a), raAgain(a), rb(b); + EXPECT_TRUE(ra == raAgain); // same underlying resource + EXPECT_TRUE(ra != rb); // different underlying resources +} + TEST(TestBuffer, DestroyIsTheSpellingClearDelegatesTo) { Counters c; diff --git a/pendingchanges/nanovdb.txt b/pendingchanges/nanovdb.txt index 8e3bdbc0ca..c676c3be27 100644 --- a/pendingchanges/nanovdb.txt +++ b/pendingchanges/nanovdb.txt @@ -4,6 +4,7 @@ NanoVDB: - Added new _hostdev_ function named nanovdb::math::isoCrossing, which intersects a ray against a user-defined iso-surface. Improvements: + - The GPU builders now allocate all scratch through an injectable memory resource: tools::cuda::TopologyBuilder and tools::cuda::MeshToGrid gained a ResourceT template parameter (defaulted, so existing code is unaffected), joining PointsToGrid. Added nanovdb::cuda::SyncFromAsync, a CRTP base that derives the synchronous half of the resource concept from the stream-ordered half, and nanovdb::cuda::ResourceRef, a non-owning reference to a resource that is itself a resource, for containers that hold their resource by value. - The bug-fix to the nanovdb::ReadAccessor (see below) improves random-access performance in some use-cases (especially on the CPU). Fixes: From fb05c71475e167e85dca6566582489c8945f2bef Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 5 Aug 2026 03:18:17 +0000 Subject: [PATCH 05/17] NanoVDB: note cuda::Buffer and cuda::BufferView in pendingchanges Co-Authored-By: Claude Opus 5 (1M context) Signed-off-by: Mark Harris --- pendingchanges/nanovdb.txt | 1 + 1 file changed, 1 insertion(+) diff --git a/pendingchanges/nanovdb.txt b/pendingchanges/nanovdb.txt index 8e3bdbc0ca..f342f5c47a 100644 --- a/pendingchanges/nanovdb.txt +++ b/pendingchanges/nanovdb.txt @@ -2,6 +2,7 @@ NanoVDB: New Features: - Added new _hostdev_ function named nanovdb::math::isoCrossing, which intersects a ray against a user-defined iso-surface. + - Added nanovdb::cuda::Buffer and nanovdb::cuda::BufferView (CUDA): a typed, resource-aware, stream-ordered container that allocates from an injectable memory resource and frees on its retained stream, and a non-owning view over externally managed memory that a GridHandle can wrap without copying. Member names follow cuda::buffer (destroy, set_stream, swap). Also added the synchronous resource concept nanovdb::cuda::is_resource alongside is_async_resource. Improvements: - The bug-fix to the nanovdb::ReadAccessor (see below) improves random-access performance in some use-cases (especially on the CPU). From 5b192c996776feb82cd983ba02c13c3bdb2b4d75 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 5 Aug 2026 03:53:01 +0000 Subject: [PATCH 06/17] NanoVDB: express PointsToGrid's density search as a loop The bisection search over voxel size was written as a backward goto, which made the lifetimes of the buffers it retries over non-lexical. Rewrite it as while(true) with continue on retry and break on convergence; the six hand-written frees before the jump are unchanged. The change is easiest to review with whitespace ignored, since the loop body re-indents: git diff -w shows 27 changed lines. d_keys and d_node_count carry results past the loop, so their declarations move above it, as does the copy event, which was created inside the retried region on every iteration but destroyed only once at the end -- each retry leaked the previous handle. Co-Authored-By: Claude Opus 5 (1M context) Signed-off-by: Mark Harris --- nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh | 199 +++++++++++--------- 1 file changed, 105 insertions(+), 94 deletions(-) diff --git a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh index 4ac4ed386d..3fb555851a 100644 --- a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh @@ -610,112 +610,123 @@ void PointsToGrid::countNodes(const PtrT points, size_t point bool operator<(const Foo &rhs) const {return density < rhs.density || (density == rhs.density && dx < rhs.dx);} } min{0.0, 1}, max{0.0, 0};// min: as dx -> 0 density -> 1 point per voxel, max: density is 0 i.e. undefined -jump:// this marks the beginning of the actual algorithm - - mData.d_keys = static_cast(mResource->allocate_async(pointCount*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); - mData.d_indx = static_cast(mResource->allocate_async(pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream));// uint32_t can index 4.29 billion Coords, corresponding to 48 GB - cudaCheck(cudaMemcpyAsync(mDeviceData, &mData, sizeof(PointsToGridData), cudaMemcpyHostToDevice, mStream));// copy mData from CPU -> GPU + // Declared ahead of the search loop below: d_keys and d_node_count carry + // results past it, and the event is recorded and re-used across iterations + // (previously it was re-created per iteration, leaking the prior handle). + uint64_t *d_keys = nullptr; + uint32_t *d_indx = nullptr, *d_points_per_tile = nullptr, *d_node_count = nullptr; + cudaEvent_t copyEvent; + cudaCheck(cudaEventCreate(©Event)); - if (mVerbose==2) mTimer.start("\nAllocating arrays for keys and indices"); - auto *d_keys = static_cast(mResource->allocate_async(pointCount*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); - auto *d_indx = static_cast(mResource->allocate_async(pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + // Bisection search for the voxel size dx that yields the target point + // density: each iteration builds tile and voxel keys at the current dx, + // then either converges or frees this iteration's buffers and retries. + while (true) { - if (mVerbose==2) mTimer.restart("Generate tile keys"); - util::cuda::lambdaKernel<<>>(pointCount, TileKeyFunctor(), mDeviceData, points, d_keys, d_indx); - cudaCheckError(); - if (mVerbose==2) mTimer.restart("DeviceRadixSort of "+std::to_string(pointCount)+" tile keys"); - CALL_CUBS(DeviceRadixSort::SortPairs, d_keys, mData.d_keys, d_indx, mData.d_indx, pointCount, 0, 63);// 21 bits per coord - std::swap(d_indx, mData.d_indx);// sorted indices are now in d_indx + mData.d_keys = static_cast(mResource->allocate_async(pointCount*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + mData.d_indx = static_cast(mResource->allocate_async(pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream));// uint32_t can index 4.29 billion Coords, corresponding to 48 GB + cudaCheck(cudaMemcpyAsync(mDeviceData, &mData, sizeof(PointsToGridData), cudaMemcpyHostToDevice, mStream));// copy mData from CPU -> GPU - if (mVerbose==2) mTimer.restart("Allocate runs"); - auto *d_points_per_tile = static_cast(mResource->allocate_async(pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); - uint32_t *d_node_count = static_cast(mResource->allocate_async(3*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + if (mVerbose==2) mTimer.start("\nAllocating arrays for keys and indices"); + d_keys = static_cast(mResource->allocate_async(pointCount*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + d_indx = static_cast(mResource->allocate_async(pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); - if (mVerbose==2) mTimer.restart("DeviceRunLengthEncode tile keys"); - CALL_CUBS(DeviceRunLengthEncode::Encode, mData.d_keys, d_keys, d_points_per_tile, d_node_count+2, pointCount); - cudaCheck(cudaMemcpyAsync(mData.nodeCount+2, d_node_count+2, sizeof(uint32_t), cudaMemcpyDeviceToHost, mStream)); - cudaCheck(cudaStreamSynchronize(mStream)); - mData.d_tile_keys = static_cast(mResource->allocate_async(mData.nodeCount[2]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); - cudaCheck(cudaMemcpyAsync(mData.d_tile_keys, d_keys, mData.nodeCount[2]*sizeof(uint64_t), cudaMemcpyDeviceToDevice, mStream)); - - static constexpr uint32_t SEGMENTED_SORT_TILE_THRESHOLD = 32; - if (mData.nodeCount[2] >= SEGMENTED_SORT_TILE_THRESHOLD) { - // Bulk segmented sort: one kernel launch + one segmented radix sort (faster for many tiles) - if (mVerbose==2) mTimer.restart("Segmented radix sort of " + std::to_string(pointCount) + " voxel keys in " + std::to_string(mData.nodeCount[2]) + " tiles"); - auto *d_tile_offsets = static_cast(mResource->allocate_async((mData.nodeCount[2]+1)*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); - cudaCheck(cudaMemsetAsync(d_tile_offsets, 0, sizeof(uint32_t), mStream)); - CALL_CUBS(DeviceScan::InclusiveSum, d_points_per_tile, d_tile_offsets + 1, mData.nodeCount[2]); - mResource->deallocate_async(d_points_per_tile, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - - util::cuda::lambdaKernel<<>>(pointCount, BulkVoxelKeyFunctor(), mDeviceData, points, d_tile_offsets, mData.nodeCount[2], d_keys, d_indx, uint32_t(0)); + if (mVerbose==2) mTimer.restart("Generate tile keys"); + util::cuda::lambdaKernel<<>>(pointCount, TileKeyFunctor(), mDeviceData, points, d_keys, d_indx); cudaCheckError(); - CALL_CUBS(DeviceSegmentedRadixSort::SortPairs, d_keys, mData.d_keys, d_indx, mData.d_indx, (int)pointCount, (int)mData.nodeCount[2], d_tile_offsets, d_tile_offsets + 1, 0, 36); - mResource->deallocate_async(d_tile_offsets, (mData.nodeCount[2]+1)*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - } else { - // Serial per-tile sort: individual kernel + sort per tile (lower overhead for few tiles) - if (mVerbose==2) mTimer.restart("DeviceRadixSort of " + std::to_string(pointCount) + " voxel keys in " + std::to_string(mData.nodeCount[2]) + " tiles"); - uint32_t *points_per_tile = new uint32_t[mData.nodeCount[2]]; - cudaCheck(cudaMemcpyAsync(points_per_tile, d_points_per_tile, mData.nodeCount[2]*sizeof(uint32_t), cudaMemcpyDeviceToHost, mStream)); - mResource->deallocate_async(d_points_per_tile, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - for (uint32_t id = 0, offset = 0; id < mData.nodeCount[2]; ++id) { - const uint32_t count = points_per_tile[id]; - util::cuda::offsetLambdaKernel<<>>(count, offset, VoxelKeyFunctor(), mDeviceData, points, id, d_keys, d_indx); + if (mVerbose==2) mTimer.restart("DeviceRadixSort of "+std::to_string(pointCount)+" tile keys"); + CALL_CUBS(DeviceRadixSort::SortPairs, d_keys, mData.d_keys, d_indx, mData.d_indx, pointCount, 0, 63);// 21 bits per coord + std::swap(d_indx, mData.d_indx);// sorted indices are now in d_indx + + if (mVerbose==2) mTimer.restart("Allocate runs"); + d_points_per_tile = static_cast(mResource->allocate_async(pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + d_node_count = static_cast(mResource->allocate_async(3*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + + if (mVerbose==2) mTimer.restart("DeviceRunLengthEncode tile keys"); + CALL_CUBS(DeviceRunLengthEncode::Encode, mData.d_keys, d_keys, d_points_per_tile, d_node_count+2, pointCount); + cudaCheck(cudaMemcpyAsync(mData.nodeCount+2, d_node_count+2, sizeof(uint32_t), cudaMemcpyDeviceToHost, mStream)); + cudaCheck(cudaStreamSynchronize(mStream)); + mData.d_tile_keys = static_cast(mResource->allocate_async(mData.nodeCount[2]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + cudaCheck(cudaMemcpyAsync(mData.d_tile_keys, d_keys, mData.nodeCount[2]*sizeof(uint64_t), cudaMemcpyDeviceToDevice, mStream)); + + static constexpr uint32_t SEGMENTED_SORT_TILE_THRESHOLD = 32; + if (mData.nodeCount[2] >= SEGMENTED_SORT_TILE_THRESHOLD) { + // Bulk segmented sort: one kernel launch + one segmented radix sort (faster for many tiles) + if (mVerbose==2) mTimer.restart("Segmented radix sort of " + std::to_string(pointCount) + " voxel keys in " + std::to_string(mData.nodeCount[2]) + " tiles"); + auto *d_tile_offsets = static_cast(mResource->allocate_async((mData.nodeCount[2]+1)*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + cudaCheck(cudaMemsetAsync(d_tile_offsets, 0, sizeof(uint32_t), mStream)); + CALL_CUBS(DeviceScan::InclusiveSum, d_points_per_tile, d_tile_offsets + 1, mData.nodeCount[2]); + mResource->deallocate_async(d_points_per_tile, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + + util::cuda::lambdaKernel<<>>(pointCount, BulkVoxelKeyFunctor(), mDeviceData, points, d_tile_offsets, mData.nodeCount[2], d_keys, d_indx, uint32_t(0)); cudaCheckError(); - CALL_CUBS(DeviceRadixSort::SortPairs, d_keys + offset, mData.d_keys + offset, d_indx + offset, mData.d_indx + offset, count, 0, 36); - offset += count; + CALL_CUBS(DeviceSegmentedRadixSort::SortPairs, d_keys, mData.d_keys, d_indx, mData.d_indx, (int)pointCount, (int)mData.nodeCount[2], d_tile_offsets, d_tile_offsets + 1, 0, 36); + mResource->deallocate_async(d_tile_offsets, (mData.nodeCount[2]+1)*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + } else { + // Serial per-tile sort: individual kernel + sort per tile (lower overhead for few tiles) + if (mVerbose==2) mTimer.restart("DeviceRadixSort of " + std::to_string(pointCount) + " voxel keys in " + std::to_string(mData.nodeCount[2]) + " tiles"); + uint32_t *points_per_tile = new uint32_t[mData.nodeCount[2]]; + cudaCheck(cudaMemcpyAsync(points_per_tile, d_points_per_tile, mData.nodeCount[2]*sizeof(uint32_t), cudaMemcpyDeviceToHost, mStream)); + mResource->deallocate_async(d_points_per_tile, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + for (uint32_t id = 0, offset = 0; id < mData.nodeCount[2]; ++id) { + const uint32_t count = points_per_tile[id]; + util::cuda::offsetLambdaKernel<<>>(count, offset, VoxelKeyFunctor(), mDeviceData, points, id, d_keys, d_indx); + cudaCheckError(); + CALL_CUBS(DeviceRadixSort::SortPairs, d_keys + offset, mData.d_keys + offset, d_indx + offset, mData.d_indx + offset, count, 0, 36); + offset += count; + } + delete [] points_per_tile; } - delete [] points_per_tile; - } - mResource->deallocate_async(d_indx, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + mResource->deallocate_async(d_indx, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - if (mVerbose==2) mTimer.restart("Count points per voxel"); + if (mVerbose==2) mTimer.restart("Count points per voxel"); - cudaEvent_t copyEvent; - cudaCheck(cudaEventCreate(©Event)); - mData.pointsPerVoxel = static_cast(mResource->allocate_async(pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); - uint32_t *d_voxel_count = static_cast(mResource->allocate_async(sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); - CALL_CUBS(DeviceRunLengthEncode::Encode, mData.d_keys, d_keys, mData.pointsPerVoxel, d_voxel_count, pointCount); - cudaCheck(cudaMemcpyAsync(&mData.voxelCount, d_voxel_count, sizeof(uint32_t), cudaMemcpyDeviceToHost, mStream)); - cudaCheck(cudaEventRecord(copyEvent, mStream)); - mResource->deallocate_async(d_voxel_count, sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - - if (util::is_same::value) { - if (mVerbose==2) mTimer.restart("Count max points per voxel"); - uint32_t *d_maxPointsPerVoxel = static_cast(mResource->allocate_async(sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)), maxPointsPerVoxel; - cudaCheck(cudaEventSynchronize(copyEvent)); - CALL_CUBS(DeviceReduce::Max, mData.pointsPerVoxel, d_maxPointsPerVoxel, mData.voxelCount); - cudaCheck(cudaMemcpyAsync(&maxPointsPerVoxel, d_maxPointsPerVoxel, sizeof(uint32_t), cudaMemcpyDeviceToHost, mStream)); + mData.pointsPerVoxel = static_cast(mResource->allocate_async(pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + uint32_t *d_voxel_count = static_cast(mResource->allocate_async(sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + CALL_CUBS(DeviceRunLengthEncode::Encode, mData.d_keys, d_keys, mData.pointsPerVoxel, d_voxel_count, pointCount); + cudaCheck(cudaMemcpyAsync(&mData.voxelCount, d_voxel_count, sizeof(uint32_t), cudaMemcpyDeviceToHost, mStream)); cudaCheck(cudaEventRecord(copyEvent, mStream)); - mResource->deallocate_async(d_maxPointsPerVoxel, sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - double dx = mData.map.getVoxelSize()[0]; - cudaCheck(cudaEventSynchronize(copyEvent)); - if (++iterCounter >= mMaxIterations || pointCount == 1u || math::Abs((int)maxPointsPerVoxel - (int)mMaxPointsPerVoxel) <= mTolerance) { - mMaxPointsPerVoxel = maxPointsPerVoxel; - } else { - const Foo tmp{dx, maxPointsPerVoxel}; - if (maxPointsPerVoxel < mMaxPointsPerVoxel) { - if (min < tmp) min = tmp; - } else if (max.density == 0 || tmp < max) { - max = tmp; - } - if (max.density) { - dx = (min.dx*(max.density - mMaxPointsPerVoxel) + max.dx*(mMaxPointsPerVoxel-min.density))/double(max.density-min.density); - } else if (maxPointsPerVoxel > 1u) { - dx *= (mMaxPointsPerVoxel-1.0)/(maxPointsPerVoxel-1.0); - } else {// maxPointsPerVoxel = 1 so increase dx significantly - dx *= 10.0; + mResource->deallocate_async(d_voxel_count, sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + + if (util::is_same::value) { + if (mVerbose==2) mTimer.restart("Count max points per voxel"); + uint32_t *d_maxPointsPerVoxel = static_cast(mResource->allocate_async(sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)), maxPointsPerVoxel; + cudaCheck(cudaEventSynchronize(copyEvent)); + CALL_CUBS(DeviceReduce::Max, mData.pointsPerVoxel, d_maxPointsPerVoxel, mData.voxelCount); + cudaCheck(cudaMemcpyAsync(&maxPointsPerVoxel, d_maxPointsPerVoxel, sizeof(uint32_t), cudaMemcpyDeviceToHost, mStream)); + cudaCheck(cudaEventRecord(copyEvent, mStream)); + mResource->deallocate_async(d_maxPointsPerVoxel, sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + double dx = mData.map.getVoxelSize()[0]; + cudaCheck(cudaEventSynchronize(copyEvent)); + if (++iterCounter >= mMaxIterations || pointCount == 1u || math::Abs((int)maxPointsPerVoxel - (int)mMaxPointsPerVoxel) <= mTolerance) { + mMaxPointsPerVoxel = maxPointsPerVoxel; + } else { + const Foo tmp{dx, maxPointsPerVoxel}; + if (maxPointsPerVoxel < mMaxPointsPerVoxel) { + if (min < tmp) min = tmp; + } else if (max.density == 0 || tmp < max) { + max = tmp; + } + if (max.density) { + dx = (min.dx*(max.density - mMaxPointsPerVoxel) + max.dx*(mMaxPointsPerVoxel-min.density))/double(max.density-min.density); + } else if (maxPointsPerVoxel > 1u) { + dx *= (mMaxPointsPerVoxel-1.0)/(maxPointsPerVoxel-1.0); + } else {// maxPointsPerVoxel = 1 so increase dx significantly + dx *= 10.0; + } + if (mVerbose==2) printf("\ntarget density = %" PRIu32 ", current density = %" PRIu32 ", current dx = %f, next dx = %f\n", mMaxPointsPerVoxel, maxPointsPerVoxel, tmp.dx, dx); + mData.map = Map(dx); + mResource->deallocate_async(mData.d_keys, pointCount*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + mResource->deallocate_async(mData.d_indx, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + mResource->deallocate_async(d_keys, pointCount*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + mResource->deallocate_async(mData.d_tile_keys, mData.nodeCount[2]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + mResource->deallocate_async(d_node_count, 3*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + mResource->deallocate_async(mData.pointsPerVoxel, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + continue; } - if (mVerbose==2) printf("\ntarget density = %" PRIu32 ", current density = %" PRIu32 ", current dx = %f, next dx = %f\n", mMaxPointsPerVoxel, maxPointsPerVoxel, tmp.dx, dx); - mData.map = Map(dx); - mResource->deallocate_async(mData.d_keys, pointCount*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - mResource->deallocate_async(mData.d_indx, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - mResource->deallocate_async(d_keys, pointCount*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - mResource->deallocate_async(mData.d_tile_keys, mData.nodeCount[2]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - mResource->deallocate_async(d_node_count, 3*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - mResource->deallocate_async(mData.pointsPerVoxel, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - goto jump; } - } + break; + }// while (true) if (iterCounter>1 && mVerbose) std::cerr << "Used " << iterCounter << " attempts to determine dx that produces a target dpoint denisty\n\n"; if (mVerbose==2) mTimer.restart("Compute prefix sum of points per voxel"); From a12201c9856e20ed81837b28f07aa765a43b1667 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 5 Aug 2026 04:02:36 +0000 Subject: [PATCH 07/17] NanoVDB: own PointsToGrid's device arrays with cuda::Buffer Every device array PointsToGrid allocates is now owned by a Buffer> borrowing the injected resource -- members where the pipeline frees them in a later member function than the one that allocated them (countNodes allocates; processUpperNodes, processLeafNodes, processPoints and processBBox release), locals where the lifetime is contained. The raw pointers survive only as views: the device-visible fields inside mData, and the working pointers the cub and kernel calls take. Every owner event -- assignment, swap, destroy -- immediately refreshes its view, and released views are nulled so a stale use faults instead of reading a freed block. The index ping-pong becomes a swap of owners across the member/local boundary, replacing the bare pointer swap whose safety depended on nothing reading the device copy of d_indx between the two uploads. The density-search retry keeps its free-before-reallocate order via explicit destroy calls, so peak device memory is unchanged. The hand-matched byte sizes at every free site disappear, and the arrays released one line before scope exit now just leave scope. One behavior change worth naming: the too-many-points-per-leaf throw previously leaked the reduction scratch; ownership now releases it during unwind. Verified against the previous commit with a counting resource over three shapes (bulk segmented-sort branch, serial per-tile branch, and the bisection retry engaged): allocation count, free count, and total bytes are identical. Full CUDA and memory-resource suites unchanged. Co-Authored-By: Claude Opus 5 (1M context) Signed-off-by: Mark Harris --- nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh | 141 +++++++++++++------- pendingchanges/nanovdb.txt | 2 +- 2 files changed, 93 insertions(+), 50 deletions(-) diff --git a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh index 3fb555851a..5fd5640d2f 100644 --- a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh @@ -303,10 +303,20 @@ public: , mResource(&resource) , mTimer(stream) , mPointType(util::is_same::value ? PointType::Default : PointType::Disable) + , mDeviceDataBuf(stream, nanovdb::cuda::ResourceRef(resource), 1, nanovdb::cuda::noInit) + , mKeysBuf(stream, nanovdb::cuda::ResourceRef(resource), 0, nanovdb::cuda::noInit) + , mTileKeysBuf(stream, nanovdb::cuda::ResourceRef(resource), 0, nanovdb::cuda::noInit) + , mLeafKeysBuf(stream, nanovdb::cuda::ResourceRef(resource), 0, nanovdb::cuda::noInit) + , mLowerKeysBuf(stream, nanovdb::cuda::ResourceRef(resource), 0, nanovdb::cuda::noInit) + , mIndxBuf(stream, nanovdb::cuda::ResourceRef(resource), 0, nanovdb::cuda::noInit) + , mPointsPerVoxelBuf(stream, nanovdb::cuda::ResourceRef(resource), 0, nanovdb::cuda::noInit) + , mPointsPerVoxelPrefixBuf(stream, nanovdb::cuda::ResourceRef(resource), 0, nanovdb::cuda::noInit) + , mPointsPerLeafBuf(stream, nanovdb::cuda::ResourceRef(resource), 0, nanovdb::cuda::noInit) + , mPointsPerLeafPrefixBuf(stream, nanovdb::cuda::ResourceRef(resource), 0, nanovdb::cuda::noInit) , mTempDevicePool(resource) { mData.map = map; - mDeviceData = static_cast*>(mResource->allocate_async(sizeof(PointsToGridData), ResourceT::DEFAULT_ALIGNMENT, mStream)); + mDeviceData = mDeviceDataBuf.data(); } /// @brief Default constructor that calls the Map constructor defined above @@ -327,7 +337,6 @@ public: mMaxIterations = maxIterations; } - ~PointsToGrid(){ mResource->deallocate_async(mDeviceData, sizeof(PointsToGridData), ResourceT::DEFAULT_ALIGNMENT, mStream); } /// @brief Toggle on and off verbose mode /// @param level Verbose level: 0=quiet, 1=timing, 2=benchmarking @@ -392,13 +401,26 @@ private: static constexpr unsigned int mNumThreads = 128;// seems faster than the old value of 256! static unsigned int numBlocks(unsigned int n) {return (n + mNumThreads - 1) / mNumThreads;} + template + using BufT = nanovdb::cuda::Buffer>; + nanovdb::cuda::ResourceRef ref() { return nanovdb::cuda::ResourceRef(*mResource); } + cudaStream_t mStream{0}; ResourceT* mResource;// non-owning; all device allocations (mDeviceData + scratch) route through this resource instance util::cuda::Timer mTimer; PointType mPointType; std::string mGridName; int mVerbose{0}; - PointsToGridData mData, *mDeviceData; + PointsToGridData mData, *mDeviceData;// mDeviceData views mDeviceDataBuf + // Owners of the device arrays that mData's raw pointers view. Raw views are + // refreshed immediately after every owner event (assign, swap, destroy), so + // the views -- including the device-visible fields inside mData -- are + // never stale. Members rather than locals because the pipeline frees them + // across member functions (countNodes allocates; processUpperNodes, + // processLeafNodes, processPoints and processBBox release). + BufT> mDeviceDataBuf; + BufT mKeysBuf, mTileKeysBuf, mLeafKeysBuf, mLowerKeysBuf; + BufT mIndxBuf, mPointsPerVoxelBuf, mPointsPerVoxelPrefixBuf, mPointsPerLeafBuf, mPointsPerLeafPrefixBuf; uint32_t mMaxPointsPerVoxel{0u}, mMaxPointsPerLeaf{0u}; int mTolerance{1}, mMaxIterations{1}; CheckMode mChecksum{CheckMode::Disable}; @@ -613,6 +635,9 @@ void PointsToGrid::countNodes(const PtrT points, size_t point // Declared ahead of the search loop below: d_keys and d_node_count carry // results past it, and the event is recorded and re-used across iterations // (previously it was re-created per iteration, leaking the prior handle). + BufT keysScratch(mStream, this->ref(), 0, nanovdb::cuda::noInit); + BufT indxScratch(mStream, this->ref(), 0, nanovdb::cuda::noInit); + BufT nodeCountScratch(mStream, this->ref(), 0, nanovdb::cuda::noInit); uint64_t *d_keys = nullptr; uint32_t *d_indx = nullptr, *d_points_per_tile = nullptr, *d_node_count = nullptr; cudaEvent_t copyEvent; @@ -623,51 +648,60 @@ void PointsToGrid::countNodes(const PtrT points, size_t point // then either converges or frees this iteration's buffers and retries. while (true) { - mData.d_keys = static_cast(mResource->allocate_async(pointCount*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); - mData.d_indx = static_cast(mResource->allocate_async(pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream));// uint32_t can index 4.29 billion Coords, corresponding to 48 GB + mKeysBuf = BufT(mStream, this->ref(), pointCount, nanovdb::cuda::noInit); + mIndxBuf = BufT(mStream, this->ref(), pointCount, nanovdb::cuda::noInit);// uint32_t can index 4.29 billion Coords, corresponding to 48 GB + mData.d_keys = mKeysBuf.data(); + mData.d_indx = mIndxBuf.data(); cudaCheck(cudaMemcpyAsync(mDeviceData, &mData, sizeof(PointsToGridData), cudaMemcpyHostToDevice, mStream));// copy mData from CPU -> GPU if (mVerbose==2) mTimer.start("\nAllocating arrays for keys and indices"); - d_keys = static_cast(mResource->allocate_async(pointCount*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); - d_indx = static_cast(mResource->allocate_async(pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + keysScratch = BufT(mStream, this->ref(), pointCount, nanovdb::cuda::noInit); + indxScratch = BufT(mStream, this->ref(), pointCount, nanovdb::cuda::noInit); + d_keys = keysScratch.data(); + d_indx = indxScratch.data(); if (mVerbose==2) mTimer.restart("Generate tile keys"); util::cuda::lambdaKernel<<>>(pointCount, TileKeyFunctor(), mDeviceData, points, d_keys, d_indx); cudaCheckError(); if (mVerbose==2) mTimer.restart("DeviceRadixSort of "+std::to_string(pointCount)+" tile keys"); CALL_CUBS(DeviceRadixSort::SortPairs, d_keys, mData.d_keys, d_indx, mData.d_indx, pointCount, 0, 63);// 21 bits per coord - std::swap(d_indx, mData.d_indx);// sorted indices are now in d_indx + mIndxBuf.swap(indxScratch);// the sorted indices' owner is now indxScratch + d_indx = indxScratch.data();// sorted indices + mData.d_indx = mIndxBuf.data();// receives the voxel-sorted indices below if (mVerbose==2) mTimer.restart("Allocate runs"); - d_points_per_tile = static_cast(mResource->allocate_async(pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); - d_node_count = static_cast(mResource->allocate_async(3*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + BufT pointsPerTileScratch(mStream, this->ref(), pointCount, nanovdb::cuda::noInit); + d_points_per_tile = pointsPerTileScratch.data(); + nodeCountScratch = BufT(mStream, this->ref(), 3, nanovdb::cuda::noInit); + d_node_count = nodeCountScratch.data(); if (mVerbose==2) mTimer.restart("DeviceRunLengthEncode tile keys"); CALL_CUBS(DeviceRunLengthEncode::Encode, mData.d_keys, d_keys, d_points_per_tile, d_node_count+2, pointCount); cudaCheck(cudaMemcpyAsync(mData.nodeCount+2, d_node_count+2, sizeof(uint32_t), cudaMemcpyDeviceToHost, mStream)); cudaCheck(cudaStreamSynchronize(mStream)); - mData.d_tile_keys = static_cast(mResource->allocate_async(mData.nodeCount[2]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + mTileKeysBuf = BufT(mStream, this->ref(), mData.nodeCount[2], nanovdb::cuda::noInit); + mData.d_tile_keys = mTileKeysBuf.data(); cudaCheck(cudaMemcpyAsync(mData.d_tile_keys, d_keys, mData.nodeCount[2]*sizeof(uint64_t), cudaMemcpyDeviceToDevice, mStream)); static constexpr uint32_t SEGMENTED_SORT_TILE_THRESHOLD = 32; if (mData.nodeCount[2] >= SEGMENTED_SORT_TILE_THRESHOLD) { // Bulk segmented sort: one kernel launch + one segmented radix sort (faster for many tiles) if (mVerbose==2) mTimer.restart("Segmented radix sort of " + std::to_string(pointCount) + " voxel keys in " + std::to_string(mData.nodeCount[2]) + " tiles"); - auto *d_tile_offsets = static_cast(mResource->allocate_async((mData.nodeCount[2]+1)*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + BufT tileOffsetsScratch(mStream, this->ref(), mData.nodeCount[2]+1, nanovdb::cuda::noInit); + auto *d_tile_offsets = tileOffsetsScratch.data(); cudaCheck(cudaMemsetAsync(d_tile_offsets, 0, sizeof(uint32_t), mStream)); CALL_CUBS(DeviceScan::InclusiveSum, d_points_per_tile, d_tile_offsets + 1, mData.nodeCount[2]); - mResource->deallocate_async(d_points_per_tile, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + pointsPerTileScratch.destroy(); d_points_per_tile = nullptr; util::cuda::lambdaKernel<<>>(pointCount, BulkVoxelKeyFunctor(), mDeviceData, points, d_tile_offsets, mData.nodeCount[2], d_keys, d_indx, uint32_t(0)); cudaCheckError(); CALL_CUBS(DeviceSegmentedRadixSort::SortPairs, d_keys, mData.d_keys, d_indx, mData.d_indx, (int)pointCount, (int)mData.nodeCount[2], d_tile_offsets, d_tile_offsets + 1, 0, 36); - mResource->deallocate_async(d_tile_offsets, (mData.nodeCount[2]+1)*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); } else { // Serial per-tile sort: individual kernel + sort per tile (lower overhead for few tiles) if (mVerbose==2) mTimer.restart("DeviceRadixSort of " + std::to_string(pointCount) + " voxel keys in " + std::to_string(mData.nodeCount[2]) + " tiles"); uint32_t *points_per_tile = new uint32_t[mData.nodeCount[2]]; cudaCheck(cudaMemcpyAsync(points_per_tile, d_points_per_tile, mData.nodeCount[2]*sizeof(uint32_t), cudaMemcpyDeviceToHost, mStream)); - mResource->deallocate_async(d_points_per_tile, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + pointsPerTileScratch.destroy(); d_points_per_tile = nullptr; for (uint32_t id = 0, offset = 0; id < mData.nodeCount[2]; ++id) { const uint32_t count = points_per_tile[id]; util::cuda::offsetLambdaKernel<<>>(count, offset, VoxelKeyFunctor(), mDeviceData, points, id, d_keys, d_indx); @@ -677,25 +711,28 @@ void PointsToGrid::countNodes(const PtrT points, size_t point } delete [] points_per_tile; } - mResource->deallocate_async(d_indx, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + indxScratch.destroy(); d_indx = nullptr;// tile-order copy, superseded by the voxel sort if (mVerbose==2) mTimer.restart("Count points per voxel"); - mData.pointsPerVoxel = static_cast(mResource->allocate_async(pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); - uint32_t *d_voxel_count = static_cast(mResource->allocate_async(sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + mPointsPerVoxelBuf = BufT(mStream, this->ref(), pointCount, nanovdb::cuda::noInit); + mData.pointsPerVoxel = mPointsPerVoxelBuf.data(); + BufT voxelCountScratch(mStream, this->ref(), 1, nanovdb::cuda::noInit); + uint32_t *d_voxel_count = voxelCountScratch.data(); CALL_CUBS(DeviceRunLengthEncode::Encode, mData.d_keys, d_keys, mData.pointsPerVoxel, d_voxel_count, pointCount); cudaCheck(cudaMemcpyAsync(&mData.voxelCount, d_voxel_count, sizeof(uint32_t), cudaMemcpyDeviceToHost, mStream)); cudaCheck(cudaEventRecord(copyEvent, mStream)); - mResource->deallocate_async(d_voxel_count, sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + voxelCountScratch.destroy(); d_voxel_count = nullptr; if (util::is_same::value) { if (mVerbose==2) mTimer.restart("Count max points per voxel"); - uint32_t *d_maxPointsPerVoxel = static_cast(mResource->allocate_async(sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)), maxPointsPerVoxel; + BufT maxPointsPerVoxelScratch(mStream, this->ref(), 1, nanovdb::cuda::noInit); + uint32_t *d_maxPointsPerVoxel = maxPointsPerVoxelScratch.data(), maxPointsPerVoxel; cudaCheck(cudaEventSynchronize(copyEvent)); CALL_CUBS(DeviceReduce::Max, mData.pointsPerVoxel, d_maxPointsPerVoxel, mData.voxelCount); cudaCheck(cudaMemcpyAsync(&maxPointsPerVoxel, d_maxPointsPerVoxel, sizeof(uint32_t), cudaMemcpyDeviceToHost, mStream)); cudaCheck(cudaEventRecord(copyEvent, mStream)); - mResource->deallocate_async(d_maxPointsPerVoxel, sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + maxPointsPerVoxelScratch.destroy(); d_maxPointsPerVoxel = nullptr; double dx = mData.map.getVoxelSize()[0]; cudaCheck(cudaEventSynchronize(copyEvent)); if (++iterCounter >= mMaxIterations || pointCount == 1u || math::Abs((int)maxPointsPerVoxel - (int)mMaxPointsPerVoxel) <= mTolerance) { @@ -716,12 +753,14 @@ void PointsToGrid::countNodes(const PtrT points, size_t point } if (mVerbose==2) printf("\ntarget density = %" PRIu32 ", current density = %" PRIu32 ", current dx = %f, next dx = %f\n", mMaxPointsPerVoxel, maxPointsPerVoxel, tmp.dx, dx); mData.map = Map(dx); - mResource->deallocate_async(mData.d_keys, pointCount*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - mResource->deallocate_async(mData.d_indx, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - mResource->deallocate_async(d_keys, pointCount*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - mResource->deallocate_async(mData.d_tile_keys, mData.nodeCount[2]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - mResource->deallocate_async(d_node_count, 3*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - mResource->deallocate_async(mData.pointsPerVoxel, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + // free before the next iteration reallocates, so peak device + // memory matches the pre-loop behavior + mKeysBuf.destroy(); mData.d_keys = nullptr; + mIndxBuf.destroy(); mData.d_indx = nullptr; + keysScratch.destroy(); d_keys = nullptr; + mTileKeysBuf.destroy(); mData.d_tile_keys = nullptr; + nodeCountScratch.destroy(); d_node_count = nullptr; + mPointsPerVoxelBuf.destroy();mData.pointsPerVoxel = nullptr; continue; } } @@ -731,16 +770,19 @@ void PointsToGrid::countNodes(const PtrT points, size_t point if (mVerbose==2) mTimer.restart("Compute prefix sum of points per voxel"); cudaCheck(cudaEventSynchronize(copyEvent)); - mData.pointsPerVoxelPrefix = static_cast(mResource->allocate_async(mData.voxelCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + mPointsPerVoxelPrefixBuf = BufT(mStream, this->ref(), mData.voxelCount, nanovdb::cuda::noInit); + mData.pointsPerVoxelPrefix = mPointsPerVoxelPrefixBuf.data(); CALL_CUBS(DeviceScan::ExclusiveSum, mData.pointsPerVoxel, mData.pointsPerVoxelPrefix, mData.voxelCount); - mData.pointsPerLeaf = static_cast(mResource->allocate_async(pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + mPointsPerLeafBuf = BufT(mStream, this->ref(), pointCount, nanovdb::cuda::noInit); + mData.pointsPerLeaf = mPointsPerLeafBuf.data(); CALL_CUBS(DeviceRunLengthEncode::Encode, thrust::make_transform_iterator(mData.d_keys, ShiftRight<9>()), d_keys, mData.pointsPerLeaf, d_node_count, pointCount); cudaCheck(cudaMemcpyAsync(mData.nodeCount, d_node_count, sizeof(uint32_t), cudaMemcpyDeviceToHost, mStream)); cudaCheck(cudaEventRecord(copyEvent, mStream)); if constexpr(util::is_same::value) { - uint32_t *d_maxPointsPerLeaf = static_cast(mResource->allocate_async(sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + BufT maxPointsPerLeafScratch(mStream, this->ref(), 1, nanovdb::cuda::noInit); + uint32_t *d_maxPointsPerLeaf = maxPointsPerLeafScratch.data(); cudaCheck(cudaEventSynchronize(copyEvent)); CALL_CUBS(DeviceReduce::Max, mData.pointsPerLeaf, d_maxPointsPerLeaf, mData.nodeCount[0]); cudaCheck(cudaMemcpyAsync(&mMaxPointsPerLeaf, d_maxPointsPerLeaf, sizeof(uint32_t), cudaMemcpyDeviceToHost, mStream)); @@ -748,25 +790,25 @@ void PointsToGrid::countNodes(const PtrT points, size_t point if (mMaxPointsPerLeaf > std::numeric_limits::max()) { throw std::runtime_error("Too many points per leaf: "+std::to_string(mMaxPointsPerLeaf)); } - mResource->deallocate_async(d_maxPointsPerLeaf, sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); } cudaCheck(cudaEventSynchronize(copyEvent)); - mData.pointsPerLeafPrefix = static_cast(mResource->allocate_async(mData.nodeCount[0]*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + mPointsPerLeafPrefixBuf = BufT(mStream, this->ref(), mData.nodeCount[0], nanovdb::cuda::noInit); + mData.pointsPerLeafPrefix = mPointsPerLeafPrefixBuf.data(); CALL_CUBS(DeviceScan::ExclusiveSum, mData.pointsPerLeaf, mData.pointsPerLeafPrefix, mData.nodeCount[0]); cudaCheck(cudaStreamSynchronize(mStream)); - mData.d_leaf_keys = static_cast(mResource->allocate_async(mData.nodeCount[0]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + mLeafKeysBuf = BufT(mStream, this->ref(), mData.nodeCount[0], nanovdb::cuda::noInit); + mData.d_leaf_keys = mLeafKeysBuf.data(); cudaCheck(cudaMemcpyAsync(mData.d_leaf_keys, d_keys, mData.nodeCount[0]*sizeof(uint64_t), cudaMemcpyDeviceToDevice, mStream)); CALL_CUBS(DeviceSelect::Unique, thrust::make_transform_iterator(mData.d_leaf_keys, ShiftRight<12>()), d_keys, d_node_count+1, mData.nodeCount[0]);// count lower nodes cudaCheck(cudaMemcpyAsync(mData.nodeCount+1, d_node_count+1, sizeof(uint32_t), cudaMemcpyDeviceToHost, mStream)); cudaCheck(cudaStreamSynchronize(mStream)); - mData.d_lower_keys = static_cast(mResource->allocate_async(mData.nodeCount[1]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + mLowerKeysBuf = BufT(mStream, this->ref(), mData.nodeCount[1], nanovdb::cuda::noInit); + mData.d_lower_keys = mLowerKeysBuf.data(); cudaCheck(cudaMemcpyAsync(mData.d_lower_keys, d_keys, mData.nodeCount[1]*sizeof(uint64_t), cudaMemcpyDeviceToDevice, mStream)); - mResource->deallocate_async(d_keys, pointCount*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - mResource->deallocate_async(d_node_count, 3*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); if (mVerbose==2) mTimer.stop(); cudaCheck(cudaEventDestroy(copyEvent)); @@ -1011,7 +1053,7 @@ inline void PointsToGrid::processUpperNodes() util::cuda::lambdaKernel<<>>(mData.nodeCount[2], BuildUpperNodesFunctor(), mDeviceData); cudaCheckError(); - mResource->deallocate_async(mData.d_tile_keys, mData.nodeCount[2]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + mTileKeysBuf.destroy(); mData.d_tile_keys = nullptr; const uint64_t valueCount = mData.nodeCount[2] << 15; util::cuda::lambdaKernel<<>>(valueCount, SetUpperBackgroundValuesFunctor(), mDeviceData); @@ -1146,11 +1188,11 @@ inline void PointsToGrid::processLeafNodes(size_t pointCount) util::cuda::lambdaKernel<<>>(mData.voxelCount, SetLeafActiveVoxelStateAndValuesFunctor(), mDeviceData); cudaCheckError(); - mResource->deallocate_async(mData.d_keys, pointCount*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - mResource->deallocate_async(mData.pointsPerVoxel, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - mResource->deallocate_async(mData.pointsPerVoxelPrefix, mData.voxelCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - mResource->deallocate_async(mData.pointsPerLeafPrefix, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - mResource->deallocate_async(mData.pointsPerLeaf,mData.nodeCount[0]*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + mKeysBuf.destroy(); mData.d_keys = nullptr; + mPointsPerVoxelBuf.destroy(); mData.pointsPerVoxel = nullptr; + mPointsPerVoxelPrefixBuf.destroy();mData.pointsPerVoxelPrefix = nullptr; + mPointsPerLeafPrefixBuf.destroy(); mData.pointsPerLeafPrefix = nullptr; + mPointsPerLeafBuf.destroy(); mData.pointsPerLeaf = nullptr; if (mVerbose==2) mTimer.restart("set inactive voxel values"); const uint64_t denseVoxelCount = mData.nodeCount[0] << 9; @@ -1159,15 +1201,16 @@ inline void PointsToGrid::processLeafNodes(size_t pointCount) if constexpr(BuildTraits::is_onindex) { if (mVerbose==2) mTimer.restart("prefix-sum for index grid"); - auto devValueIndex = static_cast(mResource->allocate_async(mData.nodeCount[0]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); - auto devValueIndexPrefix = static_cast(mResource->allocate_async(mData.nodeCount[0]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + BufT valueIndexScratch(mStream, this->ref(), mData.nodeCount[0], nanovdb::cuda::noInit); + BufT valueIndexPrefixScratch(mStream, this->ref(), mData.nodeCount[0], nanovdb::cuda::noInit); + auto devValueIndex = valueIndexScratch.data(); + auto devValueIndexPrefix = valueIndexPrefixScratch.data(); kernels::fillValueIndexKernel<<>>(mData.nodeCount[0], 0, devValueIndex, mDeviceData); cudaCheckError(); CALL_CUBS(DeviceScan::InclusiveSum, devValueIndex, devValueIndexPrefix, mData.nodeCount[0]); - mResource->deallocate_async(devValueIndex, mData.nodeCount[0]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + valueIndexScratch.destroy(); devValueIndex = nullptr; kernels::leafPrefixSumKernel<<>>(mData.nodeCount[0], 0, devValueIndexPrefix, mDeviceData); cudaCheckError(); - mResource->deallocate_async(devValueIndexPrefix, mData.nodeCount[0]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); } if (mVerbose==2) mTimer.stop(); @@ -1190,7 +1233,7 @@ template inline void PointsToGrid::processPoints(const PtrT points, size_t pointCount) { if constexpr(util::is_same::value) this->encodePoints(points, pointCount); - mResource->deallocate_async(mData.d_indx, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + mIndxBuf.destroy(); mData.d_indx = nullptr; }// PointsToGrid::processPoints template @@ -1326,7 +1369,7 @@ inline void PointsToGrid::processBBox() // update and propagate bbox from leaf -> lower/parent nodes util::cuda::lambdaKernel<<>>(mData.nodeCount[0], UpdateAndPropagateLeafBBoxFunctor(), mDeviceData); - mResource->deallocate_async(mData.d_leaf_keys, mData.nodeCount[0]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + mLeafKeysBuf.destroy(); mData.d_leaf_keys = nullptr; cudaCheckError(); // reset bbox in upper nodes @@ -1335,7 +1378,7 @@ inline void PointsToGrid::processBBox() // propagate bbox from lower -> upper/parent node util::cuda::lambdaKernel<<>>(mData.nodeCount[1], PropagateLowerBBoxFunctor(), mDeviceData); - mResource->deallocate_async(mData.d_lower_keys, mData.nodeCount[1]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + mLowerKeysBuf.destroy(); mData.d_lower_keys = nullptr; cudaCheckError() // propagate bbox from upper -> root/parent node diff --git a/pendingchanges/nanovdb.txt b/pendingchanges/nanovdb.txt index c676c3be27..8ddfebaf28 100644 --- a/pendingchanges/nanovdb.txt +++ b/pendingchanges/nanovdb.txt @@ -4,7 +4,7 @@ NanoVDB: - Added new _hostdev_ function named nanovdb::math::isoCrossing, which intersects a ray against a user-defined iso-surface. Improvements: - - The GPU builders now allocate all scratch through an injectable memory resource: tools::cuda::TopologyBuilder and tools::cuda::MeshToGrid gained a ResourceT template parameter (defaulted, so existing code is unaffected), joining PointsToGrid. Added nanovdb::cuda::SyncFromAsync, a CRTP base that derives the synchronous half of the resource concept from the stream-ordered half, and nanovdb::cuda::ResourceRef, a non-owning reference to a resource that is itself a resource, for containers that hold their resource by value. + - The GPU builders now allocate all scratch through an injectable memory resource: tools::cuda::TopologyBuilder and tools::cuda::MeshToGrid gained a ResourceT template parameter (defaulted, so existing code is unaffected), joining PointsToGrid. Added nanovdb::cuda::SyncFromAsync, a CRTP base that derives the synchronous half of the resource concept from the stream-ordered half, and nanovdb::cuda::ResourceRef, a non-owning reference to a resource that is itself a resource, for containers that hold their resource by value. PointsToGrid's device scratch and intermediate arrays are now owned by cuda::Buffer as well, replacing all of its hand-paired allocate/free calls. - The bug-fix to the nanovdb::ReadAccessor (see below) improves random-access performance in some use-cases (especially on the CPU). Fixes: From c7302bf4bfd2aca74f4878ab190316f794714ba6 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 5 Aug 2026 05:02:18 +0000 Subject: [PATCH 08/17] NanoVDB: free TopologyBuilder scratch on the caller's stream Each release site sits in a function that receives the stream, so pass it to destroy rather than relying on the retained stream matching -- they are the same on every current path, but the explicit form does not depend on that staying true. Also drop the cudaGetDevice calls whose result the Buffer conversion left unused. Co-Authored-By: Claude Opus 5 (1M context) Signed-off-by: Mark Harris --- .../nanovdb/tools/cuda/TopologyBuilder.cuh | 22 ++++++------------- 1 file changed, 7 insertions(+), 15 deletions(-) diff --git a/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh b/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh index 1aa5a68c28..13eb4714f4 100644 --- a/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh +++ b/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh @@ -147,8 +147,6 @@ void TopologyBuilder::allocateInternalMaskBuffers(cudaStream_ // 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 - int device = 0; - cudaGetDevice(&device); 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); @@ -176,8 +174,6 @@ void TopologyBuilder::countNodes(cudaStream_t stream) // as well as the tile table at the root. std::size_t size = processedTileCount*Mask<5>::SIZE; - int device = 0; - cudaGetDevice(&device); 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); @@ -452,8 +448,6 @@ inline void TopologyBuilder::processLowerNodes(cudaStream_t s using CountType = uint32_t (*)[Mask<5>::SIZE]; if (processedTileCount) { // Unless output grid is empty - int device = 0; - cudaGetDevice(&device); std::size_t lowerCount = data()->nodeCount[1]; mLowerParents = ScratchT(stream, *mResource, lowerCount*sizeof(uint32_t), nanovdb::cuda::noInit); std::size_t leafCount = data()->nodeCount[0]; @@ -475,10 +469,10 @@ inline void TopologyBuilder::processLowerNodes(cudaStream_t s } mProcessedRoot.clear(stream); - mUpperMasks.destroy(); - mLowerMasks.destroy(); - mLowerOffsets.destroy(); - mLeafOffsets.destroy(); + mUpperMasks.destroy(stream); + mLowerMasks.destroy(stream); + mLowerOffsets.destroy(stream); + mLeafOffsets.destroy(stream); }// TopologyBuilder::processLowerNodes //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- @@ -516,8 +510,6 @@ struct UpdateLeafVoxelOffsetsFunctor template inline void TopologyBuilder::processLeafOffsets(cudaStream_t stream) { - int device = 0; - cudaGetDevice(&device); 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); @@ -600,13 +592,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())); - mLeafParents.destroy(); + 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())); - mLowerParents.destroy(); + mLowerParents.destroy(stream); cudaCheckError(); // propagate bbox from upper -> root/parent node @@ -644,7 +636,7 @@ inline void TopologyBuilder::postProcessGridTree(cudaStream_t 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())); cudaCheckError(); - mVoxelOffsets.destroy(); + mVoxelOffsets.destroy(stream); tools::cuda::updateChecksum((GridData*)data()->d_bufferPtr, mChecksum, stream); }// TopologyBuilder::postProcessGridTree From f619056bbef1c5e7caeea1ba9638777d705dd058 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 5 Aug 2026 05:03:46 +0000 Subject: [PATCH 09/17] NanoVDB: borrow TopologyBuilder scratch through ResourceRef The scratch buffers held their resource by value, so each of the eight carried its own copy -- fine for the stateless default, wrong for a stateful resource, whose accounting would be split across copies while the caller's instance saw nothing. Borrow through ResourceRef instead, the same reconciliation TempPool uses. Assert the stream-ordered requirement directly in TopologyBuilder and MeshToGrid so a synchronous-only resource fails with a diagnostic that names the builder, not just the pool inside it. Note SyncFromAsync's synchronize cost on its allocate. Co-Authored-By: Claude Opus 5 (1M context) Signed-off-by: Mark Harris --- nanovdb/nanovdb/cuda/DeviceResource.h | 2 ++ nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh | 3 +++ .../nanovdb/tools/cuda/TopologyBuilder.cuh | 25 ++++++++++++++----- 3 files changed, 24 insertions(+), 6 deletions(-) diff --git a/nanovdb/nanovdb/cuda/DeviceResource.h b/nanovdb/nanovdb/cuda/DeviceResource.h index 0fcd3d0328..bef41f96d9 100644 --- a/nanovdb/nanovdb/cuda/DeviceResource.h +++ b/nanovdb/nanovdb/cuda/DeviceResource.h @@ -171,6 +171,8 @@ struct SyncFromAsync /// @brief Allocates @c bytes usable on any stream when this returns. /// @param bytes number of bytes to allocate /// @param alignment requested alignment + /// @note Every call synchronizes the null stream; on hot paths prefer the + /// stream-ordered pair. void* allocate(size_t bytes, size_t alignment) { void* p = static_cast(*this).allocate_async(bytes, alignment, cudaStream_t{0}); diff --git a/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh b/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh index b93db43d9b..7c6353dff2 100644 --- a/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh @@ -59,6 +59,9 @@ struct alignas(16) MeshToGridBoxTrianglePair { // sizeof = 16B template class MeshToGrid { + static_assert(nanovdb::cuda::is_async_resource::value, + "MeshToGrid allocates stream-ordered scratch and requires an AsyncResource"); + using PointT = nanovdb::Vec3f; using TriangleIndexT = nanovdb::Vec3i; using TriangleT = Triangle; diff --git a/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh b/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh index 13eb4714f4..45607f9a32 100644 --- a/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh +++ b/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh @@ -56,11 +56,16 @@ class TopologyBuilder using LowerT = NanoLower; using LeafT = NanoLeaf; - /// @brief Device-only scratch storage, allocated from the injected - /// resource. These buffers are never read on the host, so they use - /// the single-space Buffer rather than the dual DeviceBuffer, whose - /// host pointer and per-device array they would leave unused. - using ScratchT = nanovdb::cuda::Buffer; + static_assert(nanovdb::cuda::is_async_resource::value, + "TopologyBuilder allocates stream-ordered scratch and requires an AsyncResource"); + + /// @brief Device-only scratch storage, borrowing the injected resource + /// through a ResourceRef so all traffic reaches the caller's + /// instance (which may be stateful) rather than a copy. These + /// buffers are never read on the host, so they use the single-space + /// Buffer rather than the dual DeviceBuffer, whose host pointer and + /// per-device array they would leave unused. + using ScratchT = nanovdb::cuda::Buffer>; public: @@ -68,7 +73,15 @@ public: /// @param resource resource instance all device scratch is allocated from; /// must outlive this builder TopologyBuilder(cudaStream_t stream, ResourceT& resource = nanovdb::cuda::default_resource()) - : mResource(&resource) + : mUpperMasks(stream, resource, 0, nanovdb::cuda::noInit) + , mLowerMasks(stream, resource, 0, nanovdb::cuda::noInit) + , mUpperOffsets(stream, resource, 0, nanovdb::cuda::noInit) + , mLowerOffsets(stream, resource, 0, nanovdb::cuda::noInit) + , mLeafOffsets(stream, resource, 0, nanovdb::cuda::noInit) + , mVoxelOffsets(stream, resource, 0, nanovdb::cuda::noInit) + , mLowerParents(stream, resource, 0, nanovdb::cuda::noInit) + , mLeafParents(stream, resource, 0, nanovdb::cuda::noInit) + , mResource(&resource) , mTempDevicePool(resource) { mData = nanovdb::cuda::DeviceBuffer::create(sizeof(Data)); From 49575823678d3c9aec0455b371ec86baa0a5b596 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 5 Aug 2026 05:07:28 +0000 Subject: [PATCH 10/17] NanoVDB: guard PointsToGrid's copy event with a scope owner The too-many-points-per-leaf throw unwinds past the event's manual destroy, leaking the handle. Own it with a small guard so unwinding releases it, consistent with the buffer ownership in this function. Also assert the stream-ordered resource requirement on the class, so a synchronous-only resource fails naming PointsToGrid rather than the pool inside it. Co-Authored-By: Claude Opus 5 (1M context) Signed-off-by: Mark Harris --- nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh | 15 ++++++++++++--- 1 file changed, 12 insertions(+), 3 deletions(-) diff --git a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh index 5fd5640d2f..542efcc5a4 100644 --- a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh @@ -294,6 +294,8 @@ struct PointsToGridData { template class PointsToGrid { + static_assert(nanovdb::cuda::is_async_resource::value, + "PointsToGrid allocates stream-ordered scratch and requires an AsyncResource"); public: /// @brief Map constructor, which other constructors might call /// @param map Map to be used for the output device grid @@ -640,8 +642,16 @@ void PointsToGrid::countNodes(const PtrT points, size_t point BufT nodeCountScratch(mStream, this->ref(), 0, nanovdb::cuda::noInit); uint64_t *d_keys = nullptr; uint32_t *d_indx = nullptr, *d_points_per_tile = nullptr, *d_node_count = nullptr; - cudaEvent_t copyEvent; - cudaCheck(cudaEventCreate(©Event)); + // Owns the copy event so the too-many-points-per-leaf throw below cannot + // leak the handle. + struct EventGuard { + cudaEvent_t event; + EventGuard() { cudaCheck(cudaEventCreate(&event)); } + ~EventGuard() { cudaCheck(cudaEventDestroy(event)); } + EventGuard(const EventGuard&) = delete; + EventGuard& operator=(const EventGuard&) = delete; + } eventGuard; + cudaEvent_t copyEvent = eventGuard.event; // Bisection search for the voxel size dx that yields the target point // density: each iteration builds tile and voxel keys at the current dx, @@ -810,7 +820,6 @@ void PointsToGrid::countNodes(const PtrT points, size_t point cudaCheck(cudaMemcpyAsync(mData.d_lower_keys, d_keys, mData.nodeCount[1]*sizeof(uint64_t), cudaMemcpyDeviceToDevice, mStream)); if (mVerbose==2) mTimer.stop(); - cudaCheck(cudaEventDestroy(copyEvent)); //printf("Leaf count = %u, lower count = %u, upper count = %u\n", mData.nodeCount[0], mData.nodeCount[1], mData.nodeCount[2]); }// PointsToGrid::countNodes From 6dc37e67e385e1a1bebe6f958a423373eed10b45 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 5 Aug 2026 05:27:07 +0000 Subject: [PATCH 11/17] NanoVDB: run the builders on a synchronous memory resource The builders require a stream-ordered resource, which a device without memory-pool support cannot provide through cudaMallocAsync. Two additions close the gap. MallocResource is a synchronous cudaMalloc/cudaFree resource that works on any device. AsyncFromSync presents a synchronous resource as a stream-ordered one -- the mirror of SyncFromAsync, and the analog of cuda::mr's synchronous_resource_adapter: allocate_async forwards, since synchronously allocated memory is already valid on every stream, and deallocate_async synchronizes the stream first, which makes the synchronous contract's quiescence requirement hold. The serialization cost of that synchronize is documented at the type, so a synchronous backend is an explicit caller choice rather than a silent substitution. The new test drives PointsToGrid end-to-end with an injected AsyncFromSync over a stateful synchronous resource: every scratch allocation routes through cudaMalloc and is freed through the caller's instance, never touching stream-ordered allocation -- which is exactly what a pool-less device requires of the scratch path. Co-Authored-By: Claude Opus 5 (1M context) Signed-off-by: Mark Harris --- nanovdb/nanovdb/cuda/DeviceResource.h | 71 +++++++++++++++++++ .../nanovdb/unittest/TestMemoryResource.cu | 50 +++++++++++++ pendingchanges/nanovdb.txt | 2 +- 3 files changed, 122 insertions(+), 1 deletion(-) diff --git a/nanovdb/nanovdb/cuda/DeviceResource.h b/nanovdb/nanovdb/cuda/DeviceResource.h index bef41f96d9..f25f908fd2 100644 --- a/nanovdb/nanovdb/cuda/DeviceResource.h +++ b/nanovdb/nanovdb/cuda/DeviceResource.h @@ -192,6 +192,77 @@ struct SyncFromAsync } }; +/// @brief Synchronous device memory resource backed by cudaMalloc/cudaFree. +/// Models only the Resource concept: it never touches stream-ordered +/// allocation, so it works on devices without memory-pool support +/// (cudaDevAttrMemoryPoolsSupported == 0), where DeviceResource's +/// cudaMallocAsync path fails by design. Pair with AsyncFromSync to +/// drive the stream-ordered builders on such a device. +class MallocResource +{ +public: + // cudaMalloc aligns memory to 256 bytes by default + static constexpr size_t DEFAULT_ALIGNMENT = 256; + + /// @brief Allocates @c bytes with cudaMalloc; valid on every stream when + /// this returns. A zero request returns nullptr. + void* allocate(size_t bytes, size_t) + { + if (bytes == 0) return nullptr; + void* p = nullptr; + cudaCheck(cudaMalloc(&p, bytes)); + return p; + } + + /// @brief Frees @c p with cudaFree; the caller guarantees that device work + /// touching the memory has completed. + void deallocate(void* p, size_t, size_t) { cudaCheck(cudaFree(p)); } +};// MallocResource + +/// @brief Wrapper presenting a synchronous resource as a stream-ordered one, +/// so it can drive components that require the AsyncResource concept +/// (TempPool and the GPU builders). +/// @tparam R the wrapped synchronous resource, held by value; wrap a +/// ResourceRef to borrow a stateful instance instead. +/// @details The mirror of SyncFromAsync, and the analog of cuda::mr's +/// synchronous_resource_adapter. allocate_async forwards to +/// R::allocate, whose memory is immediately valid on every stream -- +/// a stronger guarantee than stream-ordering requires. +/// deallocate_async synchronizes @c stream before R::deallocate, +/// establishing the quiescence the synchronous contract demands. +/// @warning Every deallocation synchronizes its stream, so expect +/// serialization relative to a genuinely stream-ordered resource. +/// That is the unavoidable cost of a synchronous backend under a +/// stream-ordered algorithm; this wrapper exists so the cost is +/// explicit and chosen by the caller -- e.g. on a device without +/// memory-pool support -- rather than silently substituted. +template +struct AsyncFromSync +{ + static_assert(is_resource::value, + "AsyncFromSync requires R to model the synchronous Resource concept"); + + static constexpr size_t DEFAULT_ALIGNMENT = R::DEFAULT_ALIGNMENT; + + R resource; + + /// @brief Allocates through the synchronous resource; the result is valid + /// on every stream, hence trivially valid on @c stream. + void* allocate_async(size_t bytes, size_t alignment, cudaStream_t) { return resource.allocate(bytes, alignment); } + + /// @brief Synchronizes @c stream, then frees through the synchronous + /// resource -- the synchronize makes the quiescence contract hold. + void deallocate_async(void* p, size_t bytes, size_t alignment, cudaStream_t stream) + { + cudaCheck(cudaStreamSynchronize(stream)); + resource.deallocate(p, bytes, alignment); + } + + /// @brief Synchronous pair, forwarding to the wrapped resource. + void* allocate(size_t bytes, size_t alignment) { return resource.allocate(bytes, alignment); } + void deallocate(void* p, size_t bytes, size_t alignment) { resource.deallocate(p, bytes, alignment); } +};// AsyncFromSync + /// @brief Non-owning reference to a memory resource that is itself a resource: /// copying the ref shares the underlying resource rather than copying it. /// @tparam R the referenced resource type diff --git a/nanovdb/nanovdb/unittest/TestMemoryResource.cu b/nanovdb/nanovdb/unittest/TestMemoryResource.cu index ed5591055b..5171cc658e 100644 --- a/nanovdb/nanovdb/unittest/TestMemoryResource.cu +++ b/nanovdb/nanovdb/unittest/TestMemoryResource.cu @@ -136,6 +136,56 @@ TEST(TestMemoryResource, PinnedResource_DefaultResourceRoundTrip) // retained stream (the stream of the most recent reallocate), not the null stream. //====================================================================== +// Synchronous-only and stateful: the shape of a vGPU or arena backend. +struct SyncCountingResource +{ + static constexpr size_t DEFAULT_ALIGNMENT = nanovdb::cuda::MallocResource::DEFAULT_ALIGNMENT; + int allocs = 0, deallocs = 0; + void* allocate(size_t bytes, size_t alignment) { + void* p = nanovdb::cuda::MallocResource{}.allocate(bytes, alignment); + if (p) ++allocs; + return p; + } + void deallocate(void* p, size_t bytes, size_t alignment) { + if (p) ++deallocs; + nanovdb::cuda::MallocResource{}.deallocate(p, bytes, alignment); + } +}; + +static_assert(nanovdb::cuda::is_resource::value, + "MallocResource must model the synchronous Resource concept"); +static_assert(!nanovdb::cuda::is_async_resource::value, + "MallocResource must not claim the AsyncResource concept"); +static_assert(nanovdb::cuda::is_async_resource< + nanovdb::cuda::AsyncFromSync>::value, + "AsyncFromSync must lift a synchronous resource to AsyncResource"); + +TEST(TestMemoryResource, PointsToGrid_RunsOnSynchronousResource) +{ + // The pool-less-device path: every builder allocation routes through + // cudaMalloc/cudaFree via AsyncFromSync, never touching cudaMallocAsync. + // The grid handle's own buffer is separate from the injected resource. + using RefT = nanovdb::cuda::ResourceRef; + using VgpuT = nanovdb::cuda::AsyncFromSync; + SyncCountingResource base; + VgpuT res{RefT(base)}; + + const std::vector voxels = {{0,0,0},{1,2,3},{8,8,8},{100,100,100},{-50,20,7}}; + nanovdb::Coord* d_voxels = nullptr; + ASSERT_EQ(cudaMalloc(&d_voxels, voxels.size()*sizeof(nanovdb::Coord)), cudaSuccess); + ASSERT_EQ(cudaMemcpy(d_voxels, voxels.data(), voxels.size()*sizeof(nanovdb::Coord), cudaMemcpyHostToDevice), cudaSuccess); + { + nanovdb::tools::cuda::PointsToGrid converter(nanovdb::Map(1.0), cudaStream_t{0}, res); + auto handle = converter.getHandle(d_voxels, voxels.size()); + auto* grid = handle.deviceGrid(); + EXPECT_NE(grid, nullptr); + } + ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess); + ASSERT_EQ(cudaFree(d_voxels), cudaSuccess); + EXPECT_GT(base.allocs, 0); // scratch really routed through the sync resource + EXPECT_EQ(base.allocs, base.deallocs); // and every allocation was freed through it +} + TEST(TestMemoryResource, TempPool_FreesOnRetainedStream) { cudaStream_t s = nullptr; diff --git a/pendingchanges/nanovdb.txt b/pendingchanges/nanovdb.txt index 906e6fdaeb..2b1a300b28 100644 --- a/pendingchanges/nanovdb.txt +++ b/pendingchanges/nanovdb.txt @@ -5,7 +5,7 @@ NanoVDB: - Added nanovdb::cuda::Buffer and nanovdb::cuda::BufferView (CUDA): a typed, resource-aware, stream-ordered container that allocates from an injectable memory resource and frees on its retained stream, and a non-owning view over externally managed memory that a GridHandle can wrap without copying. Member names follow cuda::buffer (destroy, set_stream, swap). Also added the synchronous resource concept nanovdb::cuda::is_resource alongside is_async_resource. Improvements: - - The GPU builders now allocate all scratch through an injectable memory resource: tools::cuda::TopologyBuilder and tools::cuda::MeshToGrid gained a ResourceT template parameter (defaulted, so existing code is unaffected), joining PointsToGrid. Added nanovdb::cuda::SyncFromAsync, a CRTP base that derives the synchronous half of the resource concept from the stream-ordered half, and nanovdb::cuda::ResourceRef, a non-owning reference to a resource that is itself a resource, for containers that hold their resource by value. PointsToGrid's device scratch and intermediate arrays are now owned by cuda::Buffer as well, replacing all of its hand-paired allocate/free calls. + - The GPU builders now allocate all scratch through an injectable memory resource: tools::cuda::TopologyBuilder and tools::cuda::MeshToGrid gained a ResourceT template parameter (defaulted, so existing code is unaffected), joining PointsToGrid. Added nanovdb::cuda::SyncFromAsync, a CRTP base that derives the synchronous half of the resource concept from the stream-ordered half, and nanovdb::cuda::ResourceRef, a non-owning reference to a resource that is itself a resource, for containers that hold their resource by value. PointsToGrid's device scratch and intermediate arrays are now owned by cuda::Buffer as well, replacing all of its hand-paired allocate/free calls. Added nanovdb::cuda::MallocResource, a synchronous cudaMalloc-backed resource that works on devices without memory-pool support, and nanovdb::cuda::AsyncFromSync, which presents any synchronous resource as a stream-ordered one by synchronizing before each deallocation, so the builders can run with an injected synchronous resource. - The bug-fix to the nanovdb::ReadAccessor (see below) improves random-access performance in some use-cases (especially on the CPU). Fixes: From 2608ea8758e4c4e952ea131f0c2a0f03073b4d8f Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 5 Aug 2026 05:29:56 +0000 Subject: [PATCH 12/17] NanoVDB: assert the scratch alignment TopologyBuilder relies on The byte scratch is reinterpreted as word-sized types, which is valid for every resource whose DEFAULT_ALIGNMENT is at least word alignment -- all CUDA allocation paths give 256 -- but nothing said so. Assert it, so a custom resource with a weaker guarantee fails at compile time instead of misaligning on the device. Co-Authored-By: Claude Opus 5 (1M context) Signed-off-by: Mark Harris --- nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh | 2 ++ 1 file changed, 2 insertions(+) diff --git a/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh b/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh index 45607f9a32..9159002927 100644 --- a/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh +++ b/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh @@ -58,6 +58,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"); /// @brief Device-only scratch storage, borrowing the injected resource /// through a ResourceRef so all traffic reaches the caller's From 2e1e237518b047f0ad32b5ee0f61fa8490555a23 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 5 Aug 2026 05:29:58 +0000 Subject: [PATCH 13/17] NanoVDB: terminate a bare cudaCheckError with a semicolon Legal without one -- the macro expands to a braced block -- but every other use in the file spells it as a statement. Co-Authored-By: Claude Opus 5 (1M context) Signed-off-by: Mark Harris --- nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh index 542efcc5a4..4f81a51a54 100644 --- a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh @@ -1388,7 +1388,7 @@ inline void PointsToGrid::processBBox() // propagate bbox from lower -> upper/parent node util::cuda::lambdaKernel<<>>(mData.nodeCount[1], PropagateLowerBBoxFunctor(), mDeviceData); mLowerKeysBuf.destroy(); mData.d_lower_keys = nullptr; - cudaCheckError() + cudaCheckError(); // propagate bbox from upper -> root/parent node util::cuda::lambdaKernel<<>>(mData.nodeCount[2], PropagateUpperBBoxFunctor(), mDeviceData); From aac7f0050e7d01e5c11b58a4427ee52e7e2938dd Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 5 Aug 2026 05:42:36 +0000 Subject: [PATCH 14/17] NanoVDB: make AsyncFromSync's null deallocation a no-op Null-free is a no-op everywhere else, so there is no quiescence to establish and no reason to synchronize the stream for it. Also correct the sync-resource test's comment: the grid handle's output buffer is allocated through BufferT, not the injected resource, so "every builder allocation" overstated what the test demonstrates -- it is the scratch allocations that never touch stream-ordered allocation. Co-Authored-By: Claude Opus 5 (1M context) Signed-off-by: Mark Harris --- nanovdb/nanovdb/cuda/DeviceResource.h | 2 ++ nanovdb/nanovdb/unittest/TestMemoryResource.cu | 5 +++-- 2 files changed, 5 insertions(+), 2 deletions(-) diff --git a/nanovdb/nanovdb/cuda/DeviceResource.h b/nanovdb/nanovdb/cuda/DeviceResource.h index f25f908fd2..21da9b0ca2 100644 --- a/nanovdb/nanovdb/cuda/DeviceResource.h +++ b/nanovdb/nanovdb/cuda/DeviceResource.h @@ -252,8 +252,10 @@ struct AsyncFromSync /// @brief Synchronizes @c stream, then frees through the synchronous /// resource -- the synchronize makes the quiescence contract hold. + /// Null is a no-op and skips the synchronize. void deallocate_async(void* p, size_t bytes, size_t alignment, cudaStream_t stream) { + if (p == nullptr) return; cudaCheck(cudaStreamSynchronize(stream)); resource.deallocate(p, bytes, alignment); } diff --git a/nanovdb/nanovdb/unittest/TestMemoryResource.cu b/nanovdb/nanovdb/unittest/TestMemoryResource.cu index 5171cc658e..a0a6ae3128 100644 --- a/nanovdb/nanovdb/unittest/TestMemoryResource.cu +++ b/nanovdb/nanovdb/unittest/TestMemoryResource.cu @@ -162,9 +162,10 @@ static_assert(nanovdb::cuda::is_async_resource< TEST(TestMemoryResource, PointsToGrid_RunsOnSynchronousResource) { - // The pool-less-device path: every builder allocation routes through + // The pool-less-device path: every scratch allocation routes through // cudaMalloc/cudaFree via AsyncFromSync, never touching cudaMallocAsync. - // The grid handle's own buffer is separate from the injected resource. + // The grid handle's output buffer is the exception -- getHandle allocates + // it through BufferT::create, not through the injected resource. using RefT = nanovdb::cuda::ResourceRef; using VgpuT = nanovdb::cuda::AsyncFromSync; SyncCountingResource base; From 5fda0ac64e1c1340272074accc849af8a86260ed Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 5 Aug 2026 07:14:51 +0000 Subject: [PATCH 15/17] NanoVDB: set_stream matches cuda::buffer's contract, not a divergence cuda::buffer's set_stream is deliberately non-synchronizing -- its documented synchronization is a stale note left behind when the synchronization was removed -- so our non-synchronizing set_stream matches it in both name and contract, and the comment claiming a divergence was wrong. Co-Authored-By: Claude Fable 5 Signed-off-by: Mark Harris --- nanovdb/nanovdb/cuda/Buffer.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/nanovdb/nanovdb/cuda/Buffer.h b/nanovdb/nanovdb/cuda/Buffer.h index d9144591ca..44c2f44f0a 100644 --- a/nanovdb/nanovdb/cuda/Buffer.h +++ b/nanovdb/nanovdb/cuda/Buffer.h @@ -180,8 +180,8 @@ class Buffer : private detail::StreamHolder::value> /// @param stream cuda stream subsequent deallocation is ordered on /// @warning The caller is responsible for ordering @c stream after any /// in-flight work that uses the buffer's memory. This deliberately - /// does not synchronize, matching cuda::buffer's - /// set_stream_unsynchronized rather than its set_stream. + /// does not synchronize, matching cuda::buffer's set_stream, which + /// avoids implicit synchronization in fundamental primitives. template::value, int> = 0> void set_stream(cudaStream_t stream) { this->mStream = stream; } From b7dbe1fe4fedbc5f38d2404761de95adab0ddf15 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Sat, 8 Aug 2026 00:35:45 +0000 Subject: [PATCH 16/17] NanoVDB: mark cuda::Buffer::clear deprecated in the documentation Documentation-level only: the [[deprecated]] attribute would warn from GridHandle::reset and NodeManager::reset, template members in our own headers that must keep calling clear() until every buffer type provides destroy() -- an unactionable diagnostic for callers of reset(), and a build break under -Werror. The attribute lands when those callers migrate. Co-Authored-By: Claude Fable 5 Signed-off-by: Mark Harris --- nanovdb/nanovdb/cuda/Buffer.h | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/nanovdb/nanovdb/cuda/Buffer.h b/nanovdb/nanovdb/cuda/Buffer.h index 44c2f44f0a..c4d90bda13 100644 --- a/nanovdb/nanovdb/cuda/Buffer.h +++ b/nanovdb/nanovdb/cuda/Buffer.h @@ -275,6 +275,10 @@ class Buffer : private detail::StreamHolder::value> } /// @brief Frees the buffer memory (if any) and resets to the empty state. + /// @deprecated Use destroy(). Documentation-level only for now: the + /// [[deprecated]] attribute would fire from GridHandle::reset + /// and NodeManager::reset, which must keep calling clear() + /// until every buffer type provides destroy(). /// @note Transitional, and not the name to use: it exists only because /// GridHandle::reset still calls clear() on its buffer. It goes away /// when the legacy dual buffers do and GridHandle moves to destroy(). From 7d58a33521a947b8df58f7056b264924dcaa8fae Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Mon, 10 Aug 2026 22:08:50 +0000 Subject: [PATCH 17/17] NanoVDB: exercise the concept-required members of the test resources The trait checks detect allocate/deallocate through unevaluated contexts, which never odr-use them, so nvcc warned that every synchronous pair and stub member was declared but never referenced (#177-D). The attribute route is closed -- nvcc's front end ignores [[maybe_unused]] for this diagnostic -- so reference them the honest way: a test that verifies the synchronous halves of the counting and stream-recording doubles behave like their stream-ordered halves, and one that pins the trait probes' stub behavior. The TU now builds with no warnings at all. Co-Authored-By: Claude Fable 5 Signed-off-by: Mark Harris --- nanovdb/nanovdb/unittest/TestBuffer.cu | 44 ++++++++++++++++++++++++++ 1 file changed, 44 insertions(+) diff --git a/nanovdb/nanovdb/unittest/TestBuffer.cu b/nanovdb/nanovdb/unittest/TestBuffer.cu index 40be6878a6..1ba042193f 100644 --- a/nanovdb/nanovdb/unittest/TestBuffer.cu +++ b/nanovdb/nanovdb/unittest/TestBuffer.cu @@ -151,6 +151,50 @@ static_assert(nanovdb::cuda::is_resource::value, static_assert(nanovdb::cuda::is_async_resource::value, "DualResource must satisfy the AsyncResource concept"); +//====================================================================== +// The members required by the concepts are exercised directly: the trait +// checks above detect them without odr-using them, and an unreferenced +// member of a file-local struct is dead code to the compiler. + +TEST(TestBuffer, TestResourcesSynchronousPair) +{ + // The synchronous halves of the counting and stream-recording doubles are + // required by the AsyncResource refinement; verify they behave like their + // stream-ordered halves. + Counters c; + CountingResource counting{&c}; + void* p = counting.allocate(256, CountingResource::DEFAULT_ALIGNMENT); + EXPECT_NE(p, nullptr); + EXPECT_EQ(c.allocs, 1); + counting.deallocate(p, 256, CountingResource::DEFAULT_ALIGNMENT); + EXPECT_EQ(c.deallocs, 1); + + StreamLog log; + StreamRecordingResource recording{&log}; + p = recording.allocate(256, StreamRecordingResource::DEFAULT_ALIGNMENT); + EXPECT_NE(p, nullptr); + ASSERT_EQ(log.allocStreams.size(), 1u); + EXPECT_EQ(log.allocStreams[0], cudaStream_t(0));// sync pair delegates through the null stream + recording.deallocate(p, 256, StreamRecordingResource::DEFAULT_ALIGNMENT); + ASSERT_EQ(log.deallocStreams.size(), 1u); + EXPECT_EQ(log.deallocStreams[0], cudaStream_t(0)); +} + +TEST(TestBuffer, TraitDoubleStubs) +{ + // DualResource and AsyncOnlyResource exist as concept probes; their stub + // members allocate nothing and are safe to call with null arguments. + DualResource dual; + EXPECT_EQ(dual.allocate(0, DualResource::DEFAULT_ALIGNMENT), nullptr); + dual.deallocate(nullptr, 0, DualResource::DEFAULT_ALIGNMENT); + EXPECT_EQ(dual.allocate_async(0, DualResource::DEFAULT_ALIGNMENT, cudaStream_t(0)), nullptr); + dual.deallocate_async(nullptr, 0, DualResource::DEFAULT_ALIGNMENT, cudaStream_t(0)); + + AsyncOnlyResource asyncOnly; + EXPECT_EQ(asyncOnly.allocate_async(0, AsyncOnlyResource::DEFAULT_ALIGNMENT, cudaStream_t(0)), nullptr); + asyncOnly.deallocate_async(nullptr, 0, AsyncOnlyResource::DEFAULT_ALIGNMENT, cudaStream_t(0)); +} + TEST(TestBuffer, ResourceTraits) { // The static_asserts above are the real test; this anchors them in a