From 1e693339cbc99d9dd7cc6f16e8cf8f1f8a5ce01a Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 5 Aug 2026 01:52:07 +0000 Subject: [PATCH 1/4] 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 2/4] 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 f619056bbef1c5e7caeea1ba9638777d705dd058 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 5 Aug 2026 05:03:46 +0000 Subject: [PATCH 3/4] 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 4/4] 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