From 04fd4bc3174f189b4a624cef8abed30a56884f7a Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 5 Aug 2026 01:18:58 +0000 Subject: [PATCH 01/10] 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/10] 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/10] 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/10] 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/10] 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 c7302bf4bfd2aca74f4878ab190316f794714ba6 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 5 Aug 2026 05:02:18 +0000 Subject: [PATCH 06/10] 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 07/10] 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 2608ea8758e4c4e952ea131f0c2a0f03073b4d8f Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 5 Aug 2026 05:29:56 +0000 Subject: [PATCH 08/10] 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 5fda0ac64e1c1340272074accc849af8a86260ed Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 5 Aug 2026 07:14:51 +0000 Subject: [PATCH 09/10] 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 0c174b9d5f874f76008601c8fb4081efc1f9ebdb Mon Sep 17 00:00:00 2001 From: Jonathan Swartz Date: Fri, 7 Aug 2026 16:16:56 +1200 Subject: [PATCH 10/10] NanoVDB: expose ResourceT on the five TopologyBuilder consumers DilateGrid, MergeGrids, PruneGrid, RefineGrid and CoarsenGrid each hold a TopologyBuilder that was constructed as mBuilder(stream), binding the per-type default DeviceResource with no injection path. This extends the seam pattern established for PointsToGrid (#2268) and MeshToGrid (#2269) to these five operators so their stream-ordered device scratch can route through an injected memory resource -- the last piece downstream (fvdb-core) needs before it can delete its forked headers (openvdb #2232, B5). Each class gains a `ResourceT` template parameter (defaulted to nanovdb::cuda::DeviceResource) with an is_async_resource guard assert, its TopologyBuilder member becomes TopologyBuilder, and its constructor(s) gain a trailing defaulted `ResourceT& resource` argument forwarded to the builder. All out-of-class member definitions are updated to the two-parameter form. The change is additive-in-behavior: ResourceT defaults everywhere, so every existing caller compiles unchanged with identical allocation behavior. The deliberately-untouched dual-space buffers (each operator's mProcessedRoot and the builder's mData DeviceBuffer) remain on DeviceBuffer and are Step 3's dual-space problem. Adds one injected-resource test per operator to TestMemoryResource.cu, mirroring PointsToGrid_RunsOnSynchronousResource: build a small ValueOnIndex source grid, run the op with a CountingResource, and assert allocs > 0 && allocs == deallocs. Local GPU verification (RTX PRO 6000 Blackwell, CUDA 13.0): nanovdb_test_cuda_memory_resource: 15/15 (10 existing + 5 new) nanovdb_test_cuda_buffer: 27/27 nanovdb_test_cuda: 52/53 (UnifiedBuffer_IO fails on baseline too -- missing data/3_spheres.nvdb) Co-Authored-By: Claude Opus 4.8 (1M context) Signed-off-by: Jonathan Swartz --- nanovdb/nanovdb/tools/cuda/CoarsenGrid.cuh | 46 +++--- nanovdb/nanovdb/tools/cuda/DilateGrid.cuh | 46 +++--- nanovdb/nanovdb/tools/cuda/MergeGrids.cuh | 53 ++++--- nanovdb/nanovdb/tools/cuda/PruneGrid.cuh | 46 +++--- nanovdb/nanovdb/tools/cuda/RefineGrid.cuh | 46 +++--- .../nanovdb/unittest/TestMemoryResource.cu | 139 ++++++++++++++++++ pendingchanges/nanovdbresourceseams.txt | 4 + 7 files changed, 278 insertions(+), 102 deletions(-) create mode 100644 pendingchanges/nanovdbresourceseams.txt diff --git a/nanovdb/nanovdb/tools/cuda/CoarsenGrid.cuh b/nanovdb/nanovdb/tools/cuda/CoarsenGrid.cuh index f0686eac84..0f4cebd15c 100644 --- a/nanovdb/nanovdb/tools/cuda/CoarsenGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/CoarsenGrid.cuh @@ -30,9 +30,12 @@ namespace nanovdb { namespace tools::cuda { -template +template class CoarsenGrid { + static_assert(nanovdb::cuda::is_async_resource::value, + "CoarsenGrid allocates stream-ordered scratch and requires an AsyncResource"); + using GridT = NanoGrid; using TreeT = NanoTree; using RootT = NanoRoot; @@ -43,8 +46,11 @@ public: /// @brief Constructor /// @param deviceGrid source device grid to be coarsened /// @param stream optional CUDA stream (defaults to CUDA stream 0) - CoarsenGrid(const GridT* d_srcGrid, cudaStream_t stream = 0) - : mBuilder(stream), mStream(stream), mTimer(stream), mDeviceSrcGrid(d_srcGrid) {} + /// @param resource resource instance all device scratch is allocated from; + /// must outlive this operator (defaults to the per-type default resource) + CoarsenGrid(const GridT* d_srcGrid, cudaStream_t stream = 0, + ResourceT& resource = nanovdb::cuda::default_resource()) + : mBuilder(stream, resource), mStream(stream), mTimer(stream), mDeviceSrcGrid(d_srcGrid) {} /// @brief Toggle on and off verbose mode /// @param level Verbose level: 0=quiet, 1=timing, 2=benchmarking @@ -74,20 +80,20 @@ 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}; util::cuda::Timer mTimer; int mVerbose{0}; const GridT *mDeviceSrcGrid; TreeData mSrcTreeData; -};// tools::cuda::CoarsenGrid +};// tools::cuda::CoarsenGrid //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template +template template GridHandle -CoarsenGrid::getHandle(const BufferT &pool) +CoarsenGrid::getHandle(const BufferT &pool) { // Copy TreeData from GPU -> CPU cudaStreamSynchronize(mStream); @@ -147,12 +153,12 @@ CoarsenGrid::getHandle(const BufferT &pool) cudaStreamSynchronize(mStream); return GridHandle(std::move(buffer)); -}// CoarsenGrid::getHandle +}// CoarsenGrid::getHandle //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template -void CoarsenGrid::coarsenRoot() +template +void CoarsenGrid::coarsenRoot() { // This method coarsens the root tiles, to accommodate for the overall downsamping operation. @@ -197,12 +203,12 @@ void CoarsenGrid::coarsenRoot() for (const auto& [key, tile] : coarsenedTiles) *coarsenedRootPtr->tile(t++) = tile; mBuilder.mProcessedRoot.deviceUpload(device, mStream, false); -}// CoarsenGrid::coarsenRoot +}// CoarsenGrid::coarsenRoot //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template -void CoarsenGrid::coarsenInternalNodes() +template +void CoarsenGrid::coarsenInternalNodes() { // Computes the masks of upper and (densified) lower internal nodes, as a result of the coarsening operation // Masks of lower internal nodes are densified in the sense that a serialized array of them is allocated, @@ -212,24 +218,24 @@ void CoarsenGrid::coarsenInternalNodes() srcLeafCount, util::morphology::cuda::CoarsenInternalNodesFunctor(), mDeviceSrcGrid, mBuilder.deviceProcessedRoot(), mBuilder.mUpperMasks.data(), mBuilder.mLowerMasks.data() ); } -}// CoarsenGrid::coarsenInternalNodes +}// CoarsenGrid::coarsenInternalNodes //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template -void CoarsenGrid::processGridTreeRoot() +template +void CoarsenGrid::processGridTreeRoot() { // Copy GridData from source grid // By convention: this will duplicate grid name and map. Others will be reset later cudaCheck(cudaMemcpyAsync(&mBuilder.data()->getGrid(), mDeviceSrcGrid->data(), GridT::memUsage(), cudaMemcpyDeviceToDevice, mStream)); util::cuda::lambdaKernel<<<1, 1, 0, mStream>>>(1, topology::detail::BuildGridTreeRootFunctor(), mBuilder.deviceData()); cudaCheckError(); -}// CoarsenGrid::processGridTreeRoot +}// CoarsenGrid::processGridTreeRoot //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template -void CoarsenGrid::coarsenLeafNodes() +template +void CoarsenGrid::coarsenLeafNodes() { // Coarsens the active masks of the source grid (as indicated at the leaf level), into a new grid that // has been already topologically coarsened to include all necessary leaf nodes. @@ -241,7 +247,7 @@ void CoarsenGrid::coarsenLeafNodes() // Update leaf offsets and prefix sums mBuilder.processLeafOffsets(mStream); -}// CoarsenGrid::coarsenLeafNodes +}// CoarsenGrid::coarsenLeafNodes //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- diff --git a/nanovdb/nanovdb/tools/cuda/DilateGrid.cuh b/nanovdb/nanovdb/tools/cuda/DilateGrid.cuh index f13ea5df32..37155b0cc2 100644 --- a/nanovdb/nanovdb/tools/cuda/DilateGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/DilateGrid.cuh @@ -30,9 +30,12 @@ namespace nanovdb { namespace tools::cuda { -template +template class DilateGrid { + static_assert(nanovdb::cuda::is_async_resource::value, + "DilateGrid allocates stream-ordered scratch and requires an AsyncResource"); + using GridT = NanoGrid; using TreeT = NanoTree; using RootT = NanoRoot; @@ -43,8 +46,11 @@ public: /// @brief Constructor /// @param deviceGrid source device grid to be dilated /// @param stream optional CUDA stream (defaults to CUDA stream 0) - DilateGrid(const GridT* d_srcGrid, cudaStream_t stream = 0) - : mBuilder(stream), mStream(stream), mTimer(stream), mDeviceSrcGrid(d_srcGrid) {} + /// @param resource resource instance all device scratch is allocated from; + /// must outlive this operator (defaults to the per-type default resource) + DilateGrid(const GridT* d_srcGrid, cudaStream_t stream = 0, + ResourceT& resource = nanovdb::cuda::default_resource()) + : mBuilder(stream, resource), mStream(stream), mTimer(stream), mDeviceSrcGrid(d_srcGrid) {} /// @brief Toggle on and off verbose mode /// @param level Verbose level: 0=quiet, 1=timing, 2=benchmarking @@ -78,21 +84,21 @@ 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}; util::cuda::Timer mTimer; int mVerbose{0}; const GridT *mDeviceSrcGrid; morphology::NearestNeighbors mOp{morphology::NN_FACE_EDGE_VERTEX}; TreeData mSrcTreeData; -};// tools::cuda::DilateGrid +};// tools::cuda::DilateGrid //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template +template template GridHandle -DilateGrid::getHandle(const BufferT &pool) +DilateGrid::getHandle(const BufferT &pool) { // Copy TreeData from GPU -> CPU cudaStreamSynchronize(mStream); @@ -152,12 +158,12 @@ DilateGrid::getHandle(const BufferT &pool) cudaStreamSynchronize(mStream); return GridHandle(std::move(buffer)); -}// DilateGrid::getHandle +}// DilateGrid::getHandle //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template -void DilateGrid::dilateRoot() +template +void DilateGrid::dilateRoot() { // This method conservatively and speculatively dilates the root tiles, to accommodate // any new root nodes that might be introduced by the dilation operation. @@ -220,12 +226,12 @@ void DilateGrid::dilateRoot() for (const auto& [key, tile] : dilatedTiles) *dilatedRootPtr->tile(t++) = tile; mBuilder.mProcessedRoot.deviceUpload(device, mStream, false); -}// DilateGrid::dilateRoot +}// DilateGrid::dilateRoot //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template -void DilateGrid::dilateInternalNodes() +template +void DilateGrid::dilateInternalNodes() { // Computes the masks of upper and (densified) lower internal nodes, as a result of the dilation operation // Masks of lower internal nodes are densified in the sense that a serialized array of them is allocated, @@ -247,24 +253,24 @@ void DilateGrid::dilateInternalNodes() <<>> (mDeviceSrcGrid, mBuilder.deviceProcessedRoot(), mBuilder.deviceUpperMasks(), mBuilder.deviceLowerMasks()); } } -}// DilateGrid::dilateInternalNodes +}// DilateGrid::dilateInternalNodes //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template -void DilateGrid::processGridTreeRoot() +template +void DilateGrid::processGridTreeRoot() { // Copy GridData from source grid // By convention: this will duplicate grid name and map. Others will be reset later cudaCheck(cudaMemcpyAsync(&mBuilder.data()->getGrid(), mDeviceSrcGrid->data(), GridT::memUsage(), cudaMemcpyDeviceToDevice, mStream)); util::cuda::lambdaKernel<<<1, 1, 0, mStream>>>(1, topology::detail::BuildGridTreeRootFunctor(), mBuilder.deviceData()); cudaCheckError(); -}// DilateGrid::processGridTreeRoot +}// DilateGrid::processGridTreeRoot //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template -void DilateGrid::dilateLeafNodes() +template +void DilateGrid::dilateLeafNodes() { // Dilates the active masks of the source grid (as indicated at the leaf level), into a new grid that // has been already topologically dilated to include all necessary leaf nodes. @@ -285,7 +291,7 @@ void DilateGrid::dilateLeafNodes() // Update leaf offsets and prefix sums mBuilder.processLeafOffsets(mStream); -}// DilateGrid::dilateLeafNodes +}// DilateGrid::dilateLeafNodes //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- diff --git a/nanovdb/nanovdb/tools/cuda/MergeGrids.cuh b/nanovdb/nanovdb/tools/cuda/MergeGrids.cuh index c26f8a3240..70149db567 100644 --- a/nanovdb/nanovdb/tools/cuda/MergeGrids.cuh +++ b/nanovdb/nanovdb/tools/cuda/MergeGrids.cuh @@ -33,9 +33,12 @@ namespace nanovdb { namespace tools::cuda { -template +template class MergeGrids { + static_assert(nanovdb::cuda::is_async_resource::value, + "MergeGrids allocates stream-ordered scratch and requires an AsyncResource"); + using GridT = NanoGrid; using TreeT = NanoTree; using RootT = NanoRoot; @@ -46,15 +49,21 @@ public: /// @brief Constructor for an N-ary merge /// @param d_srcGrids list of source device grids to be merged (active-mask union) /// @param stream optional CUDA stream (defaults to CUDA stream 0) - MergeGrids(const std::vector& d_srcGrids, cudaStream_t stream = 0) - : mBuilder(stream), mStream(stream), mTimer(stream), mDeviceSrcGrids(d_srcGrids) {} + /// @param resource resource instance all device scratch is allocated from; + /// must outlive this operator (defaults to the per-type default resource) + MergeGrids(const std::vector& d_srcGrids, cudaStream_t stream = 0, + ResourceT& resource = nanovdb::cuda::default_resource()) + : mBuilder(stream, resource), mStream(stream), mTimer(stream), mDeviceSrcGrids(d_srcGrids) {} /// @brief Convenience constructor for the common binary merge /// @param d_srcGrid1 first source device grid to be merged /// @param d_srcGrid2 second source device grid to be merged /// @param stream optional CUDA stream (defaults to CUDA stream 0) - MergeGrids(const GridT* d_srcGrid1, const GridT* d_srcGrid2, cudaStream_t stream = 0) - : MergeGrids(std::vector{d_srcGrid1, d_srcGrid2}, stream) {} + /// @param resource resource instance all device scratch is allocated from; + /// must outlive this operator (defaults to the per-type default resource) + MergeGrids(const GridT* d_srcGrid1, const GridT* d_srcGrid2, cudaStream_t stream = 0, + ResourceT& resource = nanovdb::cuda::default_resource()) + : MergeGrids(std::vector{d_srcGrid1, d_srcGrid2}, stream, resource) {} /// @brief Toggle on and off verbose mode /// @param level Verbose level: 0=quiet, 1=timing, 2=benchmarking @@ -84,20 +93,20 @@ 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}; util::cuda::Timer mTimer; int mVerbose{0}; std::vector mDeviceSrcGrids; std::vector mSrcTreeData; -};// tools::cuda::MergeGrids +};// tools::cuda::MergeGrids //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template +template template GridHandle -MergeGrids::getHandle(const BufferT &pool) +MergeGrids::getHandle(const BufferT &pool) { if (mDeviceSrcGrids.empty()) throw std::runtime_error("MergeGrids: no input grids"); @@ -163,12 +172,12 @@ MergeGrids::getHandle(const BufferT &pool) cudaStreamSynchronize(mStream); return GridHandle(std::move(buffer)); -}// MergeGrids::getHandle +}// MergeGrids::getHandle //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template -void MergeGrids::mergeRoot() +template +void MergeGrids::mergeRoot() { // Creates a new merged tree root with the merged tiles of the two input root topologies @@ -218,12 +227,12 @@ void MergeGrids::mergeRoot() for (const auto& [key, tile] : mergedTiles) *mergedRootPtr->tile(t++) = tile; mBuilder.mProcessedRoot.deviceUpload(device, mStream, false); -}// MergeGrids::mergeRoot +}// MergeGrids::mergeRoot //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template -void MergeGrids::mergeInternalNodes() +template +void MergeGrids::mergeInternalNodes() { // Merges the masks of upper and lower nodes from both input topologies into the // densified, pre-allocated mask arrays of the merged result @@ -236,12 +245,12 @@ void MergeGrids::mergeInternalNodes() <<>> (mDeviceSrcGrids[i], mBuilder.deviceProcessedRoot(), mBuilder.deviceUpperMasks(), mBuilder.deviceLowerMasks()); } -}// MergeGrids::mergeInternalNodes +}// MergeGrids::mergeInternalNodes //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template -void MergeGrids::processGridTreeRoot() +template +void MergeGrids::processGridTreeRoot() { // Copy GridData from the first source grid // TODO: Check for instances where extra processing is needed @@ -249,12 +258,12 @@ void MergeGrids::processGridTreeRoot() cudaCheck(cudaMemcpyAsync(&mBuilder.data()->getGrid(), mDeviceSrcGrids.front()->data(), GridT::memUsage(), cudaMemcpyDeviceToDevice, mStream)); util::cuda::lambdaKernel<<<1, 1, 0, mStream>>>(1, topology::detail::BuildGridTreeRootFunctor(), mBuilder.deviceData()); cudaCheckError(); -}// MergeGrids::processGridTreeRoot +}// MergeGrids::processGridTreeRoot //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template -void MergeGrids::mergeLeafNodes() +template +void MergeGrids::mergeLeafNodes() { using Op = util::morphology::cuda::MergeLeafNodesFunctor; // Each input ORs its leaf active masks into the merged leaf topology. @@ -267,7 +276,7 @@ void MergeGrids::mergeLeafNodes() // Update leaf offsets and prefix sums mBuilder.processLeafOffsets(mStream); -}// MergeGrids::mergeLeafNodes +}// MergeGrids::mergeLeafNodes //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- diff --git a/nanovdb/nanovdb/tools/cuda/PruneGrid.cuh b/nanovdb/nanovdb/tools/cuda/PruneGrid.cuh index cb445c23ec..ee4f3f6e1f 100644 --- a/nanovdb/nanovdb/tools/cuda/PruneGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/PruneGrid.cuh @@ -30,9 +30,12 @@ namespace nanovdb { namespace tools::cuda { -template +template class PruneGrid { + static_assert(nanovdb::cuda::is_async_resource::value, + "PruneGrid allocates stream-ordered scratch and requires an AsyncResource"); + using GridT = NanoGrid; using TreeT = NanoTree; using RootT = NanoRoot; @@ -44,8 +47,11 @@ public: /// @param d_srcGrid source device grid to be pruned /// @param d_srcLeafMask sidecar array of leaf masks for voxels to retain /// @param stream optional CUDA stream (defaults to CUDA stream 0) - PruneGrid(const GridT* d_srcGrid, const Mask<3>* d_srcLeafMask, cudaStream_t stream = 0) - : mBuilder(stream), mStream(stream), mTimer(stream), mDeviceSrcGrid(d_srcGrid), mDeviceSrcLeafMask(d_srcLeafMask) {} + /// @param resource resource instance all device scratch is allocated from; + /// must outlive this operator (defaults to the per-type default resource) + PruneGrid(const GridT* d_srcGrid, const Mask<3>* d_srcLeafMask, cudaStream_t stream = 0, + ResourceT& resource = nanovdb::cuda::default_resource()) + : mBuilder(stream, resource), mStream(stream), mTimer(stream), mDeviceSrcGrid(d_srcGrid), mDeviceSrcLeafMask(d_srcLeafMask) {} /// @brief Toggle on and off verbose mode /// @param level Verbose level: 0=quiet, 1=timing, 2=benchmarking @@ -75,21 +81,21 @@ 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}; util::cuda::Timer mTimer; int mVerbose{0}; const GridT *mDeviceSrcGrid; const Mask<3> *mDeviceSrcLeafMask; TreeData mSrcTreeData; -};// tools::cuda::PruneGrid +};// tools::cuda::PruneGrid //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template +template template GridHandle -PruneGrid::getHandle(const BufferT &pool) +PruneGrid::getHandle(const BufferT &pool) { // Copy TreeData from GPU -> CPU cudaStreamSynchronize(mStream); @@ -150,12 +156,12 @@ PruneGrid::getHandle(const BufferT &pool) cudaStreamSynchronize(mStream); return GridHandle(std::move(buffer)); -}// PruneGrid::getHandle +}// PruneGrid::getHandle //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template -void PruneGrid::pruneRoot() +template +void PruneGrid::pruneRoot() { // This method conservatively (and trivially) prunes the root tile table. // For this simple approximation, it is assumed that all root tiles currently present will presist, @@ -200,12 +206,12 @@ void PruneGrid::pruneRoot() for (const auto& [key, tile] : prunedTiles) *prunedRootPtr->tile(t++) = tile; mBuilder.mProcessedRoot.deviceUpload(device, mStream, false); -}// PruneGrid::pruneRoot +}// PruneGrid::pruneRoot //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template -void PruneGrid::pruneInternalNodes() +template +void PruneGrid::pruneInternalNodes() { // Computes the masks of upper and (densified) lower internal nodes, as a result of the pruning operation // Masks of lower internal nodes are densified in the sense that a serialized array of them is allocated, @@ -215,24 +221,24 @@ void PruneGrid::pruneInternalNodes() srcLeafCount, util::morphology::cuda::PruneInternalNodesFunctor(), mDeviceSrcGrid, mBuilder.deviceProcessedRoot(), mDeviceSrcLeafMask, mBuilder.mUpperMasks.data(), mBuilder.mLowerMasks.data() ); } -}// PruneGrid::pruneInternalNodes +}// PruneGrid::pruneInternalNodes //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template -void PruneGrid::processGridTreeRoot() +template +void PruneGrid::processGridTreeRoot() { // Copy GridData from source grid // By convention: this will duplicate grid name and map. Others will be reset later cudaCheck(cudaMemcpyAsync(&mBuilder.data()->getGrid(), mDeviceSrcGrid->data(), GridT::memUsage(), cudaMemcpyDeviceToDevice, mStream)); util::cuda::lambdaKernel<<<1, 1, 0, mStream>>>(1, topology::detail::BuildGridTreeRootFunctor(), mBuilder.deviceData()); cudaCheckError(); -}// PruneGrid::processGridTreeRoot +}// PruneGrid::processGridTreeRoot //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template -void PruneGrid::pruneLeafNodes() +template +void PruneGrid::pruneLeafNodes() { // Prunes the active masks of the source grid to the intersection with the leaf-mask sidecar // followed by rebuilding the leaf offsets @@ -244,7 +250,7 @@ void PruneGrid::pruneLeafNodes() // Update leaf offsets and prefix sums mBuilder.processLeafOffsets(mStream); -}// PruneGrid::pruneLeafNodes +}// PruneGrid::pruneLeafNodes //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- diff --git a/nanovdb/nanovdb/tools/cuda/RefineGrid.cuh b/nanovdb/nanovdb/tools/cuda/RefineGrid.cuh index 07704aa16d..410173c49d 100644 --- a/nanovdb/nanovdb/tools/cuda/RefineGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/RefineGrid.cuh @@ -30,9 +30,12 @@ namespace nanovdb { namespace tools::cuda { -template +template class RefineGrid { + static_assert(nanovdb::cuda::is_async_resource::value, + "RefineGrid allocates stream-ordered scratch and requires an AsyncResource"); + using GridT = NanoGrid; using TreeT = NanoTree; using RootT = NanoRoot; @@ -43,8 +46,11 @@ public: /// @brief Constructor /// @param deviceGrid source device grid to be refined /// @param stream optional CUDA stream (defaults to CUDA stream 0) - RefineGrid(const GridT* d_srcGrid, cudaStream_t stream = 0) - : mBuilder(stream), mStream(stream), mTimer(stream), mDeviceSrcGrid(d_srcGrid) {} + /// @param resource resource instance all device scratch is allocated from; + /// must outlive this operator (defaults to the per-type default resource) + RefineGrid(const GridT* d_srcGrid, cudaStream_t stream = 0, + ResourceT& resource = nanovdb::cuda::default_resource()) + : mBuilder(stream, resource), mStream(stream), mTimer(stream), mDeviceSrcGrid(d_srcGrid) {} /// @brief Toggle on and off verbose mode /// @param level Verbose level: 0=quiet, 1=timing, 2=benchmarking @@ -74,20 +80,20 @@ 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}; util::cuda::Timer mTimer; int mVerbose{0}; const GridT *mDeviceSrcGrid; TreeData mSrcTreeData; -};// tools::cuda::RefineGrid +};// tools::cuda::RefineGrid //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template +template template GridHandle -RefineGrid::getHandle(const BufferT &pool) +RefineGrid::getHandle(const BufferT &pool) { // Copy TreeData from GPU -> CPU cudaStreamSynchronize(mStream); @@ -147,12 +153,12 @@ RefineGrid::getHandle(const BufferT &pool) cudaStreamSynchronize(mStream); return GridHandle(std::move(buffer)); -}// RefineGrid::getHandle +}// RefineGrid::getHandle //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template -void RefineGrid::refineRoot() +template +void RefineGrid::refineRoot() { // This method conservatively and speculatively refines the root tiles, to accommodate // any new root nodes that might be introduced by the upsampling operation. @@ -212,12 +218,12 @@ void RefineGrid::refineRoot() for (const auto& [key, tile] : refinedTiles) *refinedRootPtr->tile(t++) = tile; mBuilder.mProcessedRoot.deviceUpload(device, mStream, false); -}// RefineGrid::refineRoot +}// RefineGrid::refineRoot //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template -void RefineGrid::refineInternalNodes() +template +void RefineGrid::refineInternalNodes() { // Computes the masks of upper and (densified) lower internal nodes, as a result of the refinement operation // Masks of lower internal nodes are densified in the sense that a serialized array of them is allocated, @@ -227,24 +233,24 @@ void RefineGrid::refineInternalNodes() srcLeafCount, util::morphology::cuda::RefineInternalNodesFunctor(), mDeviceSrcGrid, mBuilder.deviceProcessedRoot(), mBuilder.mUpperMasks.data(), mBuilder.mLowerMasks.data() ); } -}// RefineGrid::refineInternalNodes +}// RefineGrid::refineInternalNodes //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template -void RefineGrid::processGridTreeRoot() +template +void RefineGrid::processGridTreeRoot() { // Copy GridData from source grid // By convention: this will duplicate grid name and map. Others will be reset later cudaCheck(cudaMemcpyAsync(&mBuilder.data()->getGrid(), mDeviceSrcGrid->data(), GridT::memUsage(), cudaMemcpyDeviceToDevice, mStream)); util::cuda::lambdaKernel<<<1, 1, 0, mStream>>>(1, topology::detail::BuildGridTreeRootFunctor(), mBuilder.deviceData()); cudaCheckError(); -}// RefineGrid::processGridTreeRoot +}// RefineGrid::processGridTreeRoot //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template -void RefineGrid::refineLeafNodes() +template +void RefineGrid::refineLeafNodes() { // Refines the active masks of the source grid (as indicated at the leaf level), into a new grid that // has been already topologically refined to include all necessary leaf nodes. @@ -256,7 +262,7 @@ void RefineGrid::refineLeafNodes() // Update leaf offsets and prefix sums mBuilder.processLeafOffsets(mStream); -}// RefineGrid::refineLeafNodes +}// RefineGrid::refineLeafNodes //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- diff --git a/nanovdb/nanovdb/unittest/TestMemoryResource.cu b/nanovdb/nanovdb/unittest/TestMemoryResource.cu index ed5591055b..c9fac7f460 100644 --- a/nanovdb/nanovdb/unittest/TestMemoryResource.cu +++ b/nanovdb/nanovdb/unittest/TestMemoryResource.cu @@ -11,6 +11,11 @@ #include #include #include +#include +#include +#include +#include +#include #include #include @@ -268,4 +273,138 @@ TEST(TestMemoryResource, PointsToGrid_PointEncodedWithCustomResource) EXPECT_GT(res.allocs, 0); } +//====================================================================== +// TopologyBuilder consumers (DilateGrid, MergeGrids, PruneGrid, RefineGrid, +// CoarsenGrid) route their builder's stream-ordered scratch through an +// injected resource instance (B5, openvdb #2232). +//====================================================================== + +/// @brief Build a small ValueOnIndex device grid from a handful of voxels. +/// The default resource is fine for this setup step; the CountingResource +/// of the op under test only observes that op's scratch. +static nanovdb::GridHandle +buildIndexGrid(const std::vector& voxels) +{ + nanovdb::Coord* d_voxels = nullptr; + cudaCheck(cudaMalloc(&d_voxels, voxels.size() * sizeof(nanovdb::Coord))); + cudaCheck(cudaMemcpy(d_voxels, voxels.data(), voxels.size() * sizeof(nanovdb::Coord), cudaMemcpyHostToDevice)); + nanovdb::tools::cuda::PointsToGrid converter(nanovdb::Map(1.0)); + auto handle = converter.getHandle(d_voxels, voxels.size()); + cudaCheck(cudaFree(d_voxels)); + return handle; +} + +TEST(TestMemoryResource, DilateGrid_InjectedResourceSeam) +{ + // The dilated grid handle's output buffer goes through BufferT::create, and the + // dual-space mProcessedRoot / builder mData go through DeviceBuffer -- none of these + // is observed by the CountingResource; only the builder's stream-ordered scratch is. + auto src = buildIndexGrid({{0,0,0},{1,2,3},{4,4,4}}); + auto* d_srcGrid = src.deviceGrid(); + ASSERT_NE(d_srcGrid, nullptr); + + CountingResource res; + { + nanovdb::tools::cuda::DilateGrid op(d_srcGrid, 0, res); + auto handle = op.getHandle(); + ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess); + EXPECT_TRUE(handle.deviceData()); + } + ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess); + EXPECT_GT(res.allocs, 0); // builder scratch routed through the injected instance + EXPECT_EQ(res.allocs, res.deallocs); // and every allocation was freed through it +} + +TEST(TestMemoryResource, MergeGrids_InjectedResourceSeam) +{ + // As above, the merged handle's output buffer and the dual-space mProcessedRoot / mData + // go through DeviceBuffer and are not counted -- only the builder's scratch is. + auto srcA = buildIndexGrid({{0,0,0},{1,1,1}}); + auto srcB = buildIndexGrid({{5,5,5},{6,6,6}}); + auto* gridA = srcA.deviceGrid(); + auto* gridB = srcB.deviceGrid(); + ASSERT_NE(gridA, nullptr); + ASSERT_NE(gridB, nullptr); + + CountingResource res; + { + nanovdb::tools::cuda::MergeGrids op(gridA, gridB, 0, res); + auto handle = op.getHandle(); + ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess); + EXPECT_TRUE(handle.deviceData()); + } + ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess); + EXPECT_GT(res.allocs, 0); + EXPECT_EQ(res.allocs, res.deallocs); +} + +TEST(TestMemoryResource, PruneGrid_InjectedResourceSeam) +{ + // Voxels are kept within a single 8^3 leaf, so the source grid has exactly one leaf + // and the retain-mask sidecar is a single all-on Mask<3>. As above, only the builder's + // scratch is counted; the output buffer and dual-space mProcessedRoot / mData are not. + auto src = buildIndexGrid({{0,0,0},{1,1,1},{2,2,2}}); + auto* d_srcGrid = src.deviceGrid(); + ASSERT_NE(d_srcGrid, nullptr); + + nanovdb::Mask<3> hostMask; + hostMask.setOn(); // retain every voxel + nanovdb::Mask<3>* d_mask = nullptr; + ASSERT_EQ(cudaMalloc(&d_mask, sizeof(nanovdb::Mask<3>)), cudaSuccess); + ASSERT_EQ(cudaMemcpy(d_mask, &hostMask, sizeof(nanovdb::Mask<3>), cudaMemcpyHostToDevice), cudaSuccess); + + CountingResource res; + { + nanovdb::tools::cuda::PruneGrid op(d_srcGrid, d_mask, 0, res); + auto handle = op.getHandle(); + ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess); + EXPECT_TRUE(handle.deviceData()); + } + ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess); + ASSERT_EQ(cudaFree(d_mask), cudaSuccess); + EXPECT_GT(res.allocs, 0); + EXPECT_EQ(res.allocs, res.deallocs); +} + +TEST(TestMemoryResource, RefineGrid_InjectedResourceSeam) +{ + // As above, only the builder's scratch is counted; the output buffer and dual-space + // mProcessedRoot / mData go through DeviceBuffer and are not. + auto src = buildIndexGrid({{0,0,0},{1,2,3},{4,4,4}}); + auto* d_srcGrid = src.deviceGrid(); + ASSERT_NE(d_srcGrid, nullptr); + + CountingResource res; + { + nanovdb::tools::cuda::RefineGrid op(d_srcGrid, 0, res); + auto handle = op.getHandle(); + ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess); + EXPECT_TRUE(handle.deviceData()); + } + ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess); + EXPECT_GT(res.allocs, 0); + EXPECT_EQ(res.allocs, res.deallocs); +} + +TEST(TestMemoryResource, CoarsenGrid_InjectedResourceSeam) +{ + // Voxels span a 2x2x2 block of leaves (leaf DIM = 8) so coarsening is non-degenerate. + // As above, only the builder's scratch is counted; the output buffer and dual-space + // mProcessedRoot / mData go through DeviceBuffer and are not. + auto src = buildIndexGrid({{0,0,0},{8,0,0},{0,8,0},{0,0,8},{8,8,0},{8,0,8},{0,8,8},{8,8,8}}); + auto* d_srcGrid = src.deviceGrid(); + ASSERT_NE(d_srcGrid, nullptr); + + CountingResource res; + { + nanovdb::tools::cuda::CoarsenGrid op(d_srcGrid, 0, res); + auto handle = op.getHandle(); + ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess); + EXPECT_TRUE(handle.deviceData()); + } + ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess); + EXPECT_GT(res.allocs, 0); + EXPECT_EQ(res.allocs, res.deallocs); +} + } // unnamed namespace diff --git a/pendingchanges/nanovdbresourceseams.txt b/pendingchanges/nanovdbresourceseams.txt new file mode 100644 index 0000000000..6cb174ab37 --- /dev/null +++ b/pendingchanges/nanovdbresourceseams.txt @@ -0,0 +1,4 @@ +NanoVDB: + + Improvements: + - The five tools::cuda::TopologyBuilder consumers -- DilateGrid, MergeGrids, PruneGrid, RefineGrid and CoarsenGrid -- gained a ResourceT template parameter (defaulted to nanovdb::cuda::DeviceResource, so existing code is unaffected) and a trailing defaulted resource constructor argument, forwarded to their TopologyBuilder. This routes each operator's stream-ordered device scratch through an injected memory resource, mirroring the seam added to PointsToGrid and MeshToGrid (openvdb #2232, B5). The grid handle's output buffer and the dual-space mProcessedRoot / builder mData buffers still allocate through DeviceBuffer and are left for a later step.