From 06afbf0d0132fed3c400395209fc15ea7799fc84 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Tue, 18 Aug 2026 23:28:40 +0000 Subject: [PATCH 1/2] NanoVDB: own the small builders' device scratch with cuda::Buffer (CUDA) Route the last raw mallocAsync/freeAsync pairs in the GPU tools through RAII cuda::Buffer scratch over an injectable resource, completing the audit that PointsToGrid, TopologyBuilder, MeshToGrid and the topology-op consumers already went through: IndexToGrid (node accessor + grid name), addBlindData (byte-size scratch), GridStats (per-node statistics scratch) and SignedFloodFill (node counts). Each gains a defaulted ResourceT template parameter following the established conventions -- classes take a must-outlive resource instance in the constructor, free functions take the template parameter and route through the per-type default instance. Every replaced free keeps its position in stream order via destroy(stream) at the old freeAsync site, and GridStats' mode-dependent allocation becomes a zero-element buffer when the mode has no average, so the resource sees identical traffic. IndexToGrid's nested NodeAccessor is hoisted to the namespace-scope IndexToGridNodeAccessor so the kernels that name it do not acquire a spurious ResourceT dependency; a public alias preserves the nested spelling. Seam tests with exact allocation counts cover all four (including the grid-name path, the Extrema zero-allocation path, and free-function forwarding to the default instance). Part of #2232 (step 2). Co-Authored-By: Claude Fable 5 Signed-off-by: Mark Harris --- nanovdb/nanovdb/tools/cuda/AddBlindData.cuh | 15 ++- nanovdb/nanovdb/tools/cuda/GridStats.cuh | 48 ++++--- nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh | 79 +++++++----- .../nanovdb/tools/cuda/SignedFloodFill.cuh | 38 ++++-- .../nanovdb/unittest/TestMemoryResource.cu | 117 +++++++++++++++++- pendingchanges/nanovdb.txt | 1 + 6 files changed, 232 insertions(+), 66 deletions(-) diff --git a/nanovdb/nanovdb/tools/cuda/AddBlindData.cuh b/nanovdb/nanovdb/tools/cuda/AddBlindData.cuh index 299f060ceb..9becafc803 100644 --- a/nanovdb/nanovdb/tools/cuda/AddBlindData.cuh +++ b/nanovdb/nanovdb/tools/cuda/AddBlindData.cuh @@ -18,6 +18,7 @@ #define NANOVDB_TOOLS_CUDA_ADDBLINDDATA_CUH_HAS_BEEN_INCLUDED #include +#include #include #include #include @@ -34,6 +35,7 @@ namespace tools::cuda {// ============================================ /// @tparam BuildT Build type of the grid /// @tparam BlindDataT Type of the blind data /// @tparam BufferT Type of the buffer used for allocation +/// @tparam ResourceT Template type of optional resource used for internal temporary memory /// @param d_grid Pointer to device grid /// @param d_blindData Pointer to device blind data /// @param valueCount number of values in the blind data @@ -43,7 +45,7 @@ namespace tools::cuda {// ============================================ /// @param pool optional pool used for allocation /// @param stream optional CUDA stream (defaults to CUDA stream 0) /// @return GridHandle with blind data appended -template +template GridHandle addBlindData(const NanoGrid *d_grid, const BlindDataT *d_blindData, @@ -60,11 +62,16 @@ addBlindData(const NanoGrid *d_grid, // old grid old meta new meta old data new data static_assert(BufferTraits::hasDeviceDual, "Expected BufferT to support device allocation"); + static_assert(nanovdb::cuda::is_async_resource::value, + "addBlindData allocates stream-ordered scratch and requires an AsyncResource"); // extract byte sizes of the grid, blind meta data and blind data enum {GRID=0, META=1, DATA=2, CHECKSUM=3}; - uint64_t tmp[4], *d_tmp; - cudaCheck(util::cuda::mallocAsync((void**)&d_tmp, 4*sizeof(uint64_t), stream)); + using ScratchT = nanovdb::cuda::Buffer>; + auto &resource = nanovdb::cuda::default_resource(); + uint64_t tmp[4]; + ScratchT tmpBuf(stream, nanovdb::cuda::ResourceRef(resource), 4, nanovdb::cuda::noInit); + uint64_t *d_tmp = tmpBuf.data(); util::cuda::lambdaKernel<<<1, 1, 0, stream>>>(1, [=] __device__(size_t) { if (auto count = d_grid->blindDataCount()) { d_tmp[GRID] = util::PtrDiff(&d_grid->blindMetaData(0), d_grid); @@ -116,7 +123,7 @@ addBlindData(const NanoGrid *d_grid, for (uint32_t i=0, n=grid.mBlindMetadataCount-1; imDataOffset += sizeof(GridBlindMetaData); grid.mGridSize += sizeof(GridBlindMetaData) + meta->blindDataSize();// expansion with 32 byte alignment }); cudaCheckError(); - cudaCheck(util::cuda::freeAsync(d_tmp, stream)); + tmpBuf.destroy(stream); Checksum cs(tmp[CHECKSUM]); cuda::updateChecksum(reinterpret_cast(d_data), cs.mode(), stream); diff --git a/nanovdb/nanovdb/tools/cuda/GridStats.cuh b/nanovdb/nanovdb/tools/cuda/GridStats.cuh index 1705c629f2..140e02445d 100644 --- a/nanovdb/nanovdb/tools/cuda/GridStats.cuh +++ b/nanovdb/nanovdb/tools/cuda/GridStats.cuh @@ -16,6 +16,7 @@ #define NANOVDB_TOOLS_CUDA_GRIDSTATS_CUH_HAS_BEEN_INCLUDED #include +#include #include #include // for cub::Uninitialized @@ -26,18 +27,22 @@ namespace tools::cuda { /// @brief Update, i.e. re-compute, grid statistics like min/max, stats and bbox /// information for an existing NanoVDB Grid. -/// @param grid Grid whose stats to update +/// @tparam ResourceT Template type of optional resource used for internal temporary memory +/// @param d_grid Grid whose stats to update /// @param mode Mode of computation for the statistics. /// @param stream Optional cuda stream (defaults to zero) -template +template void updateGridStats(NanoGrid *d_grid, StatsMode mode = StatsMode::Default, cudaStream_t stream = 0); //================================================================================================ -/// @brief Allows for the construction of NanoVDB grids without any dependency -template::ValueType>> +/// @brief Re-computes statistics (of type @c StatsT) for each node of an existing NanoVDB grid on the device +template::ValueType>, typename ResourceT = nanovdb::cuda::DeviceResource> class GridStats { + static_assert(nanovdb::cuda::is_async_resource::value, + "GridStats allocates stream-ordered scratch and requires an AsyncResource"); + using GridT = NanoGrid; using TreeT = typename GridT::TreeType; using ValueT = typename TreeT::ValueType; @@ -48,9 +53,19 @@ class GridStats static_assert(util::is_same::value, "Mismatching type"); ValueT mDelta; // skip rendering of node if: node.max < -mDelta || node.min > mDelta + ResourceT* mResource;// non-owning; all device scratch routes through this resource instance + + template + using BufT = nanovdb::cuda::Buffer>; + nanovdb::cuda::ResourceRef ref() { return nanovdb::cuda::ResourceRef(*mResource); } public: - GridStats(ValueT delta = ValueT(0)) : mDelta(delta) {} + /// @brief Constructor + /// @param delta skip rendering of nodes when node.max < -delta or node.min > delta + /// @param resource resource instance all device scratch is allocated from; + /// must outlive this instance (defaults to the per-type default resource) + GridStats(ValueT delta = ValueT(0), ResourceT& resource = nanovdb::cuda::default_resource()) + : mDelta(delta), mResource(&resource) {} void update(GridT *d_grid, cudaStream_t stream = 0); @@ -285,8 +300,8 @@ __global__ void processRootAndGrid(NodeManager *d_nodeMgr, StatsT *d_sta //================================================================================================ -template -void GridStats::update(NanoGrid *d_grid, cudaStream_t stream) +template +void GridStats::update(NanoGrid *d_grid, cudaStream_t stream) { static const uint32_t threadsPerBlock = 128; auto blocksPerGrid = [&](uint32_t count)->uint32_t{return (count + (threadsPerBlock - 1)) / threadsPerBlock;}; @@ -298,11 +313,12 @@ void GridStats::update(NanoGrid *d_grid, cudaStream_t st cudaCheck(cudaMemcpyAsync(nodeCount, (char*)d_grid + sizeof(GridData) + 4*sizeof(uint64_t), 3*sizeof(uint32_t), cudaMemcpyDeviceToHost, stream)); //cudaStreamSynchronize(stream);// finish all device tasks in stream - StatsT *d_stats = nullptr; - // One d_stats slot per node (leaves, then lower, then upper) so every node - // has its own slot - see processInternal. - if constexpr(StatsT::hasAverage()) cudaCheck(util::cuda::mallocAsync((void**)&d_stats, (nodeCount[0]+nodeCount[1]+nodeCount[2])*sizeof(StatsT), stream)); + // has its own slot - see processInternal. Zero elements when the mode has + // no average, so nothing is allocated. + BufT statsBuf(stream, this->ref(), + StatsT::hasAverage() ? nodeCount[0]+nodeCount[1]+nodeCount[2] : 0, nanovdb::cuda::noInit); + StatsT *d_stats = statsBuf.data(); // warp per leaf (4 warps per 128-thread block); block per internal node if (nodeCount[0]) processLeaf<<>>(d_nodeMgr, d_stats); @@ -313,25 +329,25 @@ void GridStats::update(NanoGrid *d_grid, cudaStream_t st processRootAndGrid<<<1, 1, 0, stream>>>(d_nodeMgr, d_stats); - if constexpr(StatsT::hasAverage()) cudaCheck(util::cuda::freeAsync(d_stats, stream)); + statsBuf.destroy(stream); } // cuda::GridStats::update( Grid ) //================================================================================================ -template +template void updateGridStats(NanoGrid *d_grid, StatsMode mode, cudaStream_t stream) { if (d_grid == nullptr && mode == StatsMode::Disable) { return; } else if (mode == StatsMode::BBox || util::is_same::value) { - GridStats > stats; + GridStats, ResourceT> stats; stats.update(d_grid, stream); } else if (mode == StatsMode::MinMax) { - GridStats > stats; + GridStats, ResourceT> stats; stats.update(d_grid, stream); } else if (mode == StatsMode::All) { - GridStats > stats; + GridStats, ResourceT> stats; stats.update(d_grid, stream); } else { throw std::runtime_error("GridStats: Unsupported statistics mode."); diff --git a/nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh b/nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh index 31cc0aa507..e9ee47c426 100644 --- a/nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh @@ -18,6 +18,7 @@ #define NANOVDB_TOOLS_CUDA_INDEXTOGRID_CUH_HAS_BEEN_INCLUDED #include +#include #include #include #include @@ -31,42 +32,49 @@ namespace tools::cuda {// ====================================================== /// @tparam DstBuildT Build time of the destination/output Grid /// @tparam SrcBuildT Build type of the source/input IndexGrid /// @tparam BufferT Type of the buffer used for allocation of the destination Grid +/// @tparam ResourceT Template type of optional resource used for internal temporary memory /// @param d_srcGrid Device pointer to source/input IndexGrid, i.e. SrcBuildT={ValueIndex,ValueOnIndex} /// @param d_srcValues Device pointer to an array of values /// @param pool Memory pool used to create a buffer for the destination/output Grid -/// @param stream optional CUDA stream (defaults to CUDA stream 0 +/// @param stream optional CUDA stream (defaults to CUDA stream 0) /// @note If d_srcGrid has stats (min,max,avg,std-div), the d_srcValues is also assumed /// to have the same information, all of which are then copied to the destination/output grid. /// An exception to this rule is if the type of d_srcValues is different from the stats type /// NanoRoot::FloatType, e.g. if DstBuildT=Vec3f then NanoRoot::FloatType=float, /// in which case average and standard-deviation is undefined in the output grid. /// @return returns handle to grid that combined IndexGrid and values -template +template typename util::enable_if::is_index, GridHandle>::type indexToGrid(const NanoGrid *d_srcGrid, const typename BuildToValueMap::type *d_srcValues, const BufferT &pool = BufferT(), cudaStream_t stream = 0); -template +template typename util::enable_if::is_index, GridHandle>::type createNanoGrid(const NanoGrid *d_srcGrid, const typename BuildToValueMap::type *d_srcValues, const BufferT &pool = BufferT(), cudaStream_t stream = 0) { - return indexToGrid(d_srcGrid, d_srcValues, pool, stream); + return indexToGrid(d_srcGrid, d_srcValues, pool, stream); } namespace {// anonymous namespace -template +template struct IndexToGridNodeAccessor; + +template class IndexToGrid { + static_assert(nanovdb::cuda::is_async_resource::value, + "IndexToGrid allocates stream-ordered scratch and requires an AsyncResource"); using SrcGridT = NanoGrid; public: - struct NodeAccessor; + using NodeAccessor = IndexToGridNodeAccessor; /// @brief Constructor from a source IndeGrid - /// @param srcGrid Device pointer to IndexGrid used as the source - IndexToGrid(const SrcGridT *d_srcGrid, cudaStream_t stream = 0); - - ~IndexToGrid() {cudaCheck(util::cuda::freeAsync(mDevNodeAcc, mStream));} + /// @param d_srcGrid Device pointer to IndexGrid used as the source + /// @param stream optional CUDA stream (defaults to CUDA stream 0) + /// @param resource resource instance all device scratch is allocated from; + /// must outlive this instance (defaults to the per-type default resource) + IndexToGrid(const SrcGridT *d_srcGrid, cudaStream_t stream = 0, + ResourceT& resource = nanovdb::cuda::default_resource()); /// @brief Toggle on and off verbose mode /// @param on if true verbose is turned on @@ -90,7 +98,14 @@ private: util::cuda::Timer mTimer; std::string mGridName; bool mVerbose{false}; - NodeAccessor mNodeAcc, *mDevNodeAcc; + ResourceT* mResource;// non-owning; all device scratch routes through this resource instance + template + using BufT = nanovdb::cuda::Buffer>; + nanovdb::cuda::ResourceRef ref() { return nanovdb::cuda::ResourceRef(*mResource); } + NodeAccessor mNodeAcc; + BufT mDevNodeAccBuf; + NodeAccessor *mDevNodeAcc{nullptr};// view of mDevNodeAccBuf + BufT mGridNameBuf;// owns the device copy of mGridName template BufferT getBuffer(const BufferT &pool); @@ -99,8 +114,9 @@ private: //================================================================================================ template -struct IndexToGrid::NodeAccessor +struct IndexToGridNodeAccessor { + using SrcGridT = NanoGrid; uint64_t grid, tree, root, node[3], meta, blind, size;// byte offsets, node: 0=leaf,1=lower, 2=upper const SrcGridT *d_srcGrid;// device point to source IndexGrid void *d_dstPtr;// device pointer to buffer with destination Grid @@ -125,12 +141,12 @@ struct IndexToGrid::NodeAccessor __device__ typename NanoNode::type& dstNode(int i) const { return *(util::PtrAdd::type>(d_dstPtr, node[LEVEL])+i); } -};// IndexToGrid::NodeAccessor +};// IndexToGridNodeAccessor //================================================================================================ template -__global__ void processGridTreeRootKernel(typename IndexToGrid::NodeAccessor *nodeAcc, +__global__ void processGridTreeRootKernel(IndexToGridNodeAccessor *nodeAcc, const typename BuildToValueMap::type *srcValues) { using SrcValueT = typename BuildToValueMap::type; @@ -173,7 +189,7 @@ __global__ void processGridTreeRootKernel(typename IndexToGrid::NodeA //================================================================================================ template -__global__ void processRootTilesKernel(typename IndexToGrid::NodeAccessor *nodeAcc, +__global__ void processRootTilesKernel(IndexToGridNodeAccessor *nodeAcc, const typename BuildToValueMap::type *srcValues) { const auto tileID = blockIdx.x, tileCount = nodeAcc->nodeCount[3];// note: tileID != childID! @@ -205,7 +221,7 @@ __global__ void processRootTilesKernel(typename IndexToGrid::NodeAcce //================================================================================================ template -__global__ void processNodesKernel(typename IndexToGrid::NodeAccessor *nodeAcc, +__global__ void processNodesKernel(IndexToGridNodeAccessor *nodeAcc, const typename BuildToValueMap::type *srcValues) { using SrcNodeT = typename NanoNode::type; @@ -257,7 +273,7 @@ __global__ void processNodesKernel(typename IndexToGrid::NodeAccessor //================================================================================================ template -__global__ void processLeafsKernel(typename IndexToGrid::NodeAccessor *nodeAcc, +__global__ void processLeafsKernel(IndexToGridNodeAccessor *nodeAcc, const typename BuildToValueMap::type *srcValues) { using SrcValueT = typename BuildToValueMap::type; @@ -299,7 +315,7 @@ __global__ void processLeafsKernel(typename IndexToGrid::NodeAccessor template __global__ void cpyNodeCountKernel(const NanoGrid *srcGrid, - typename IndexToGrid::NodeAccessor *nodeAcc) + IndexToGridNodeAccessor *nodeAcc) { assert(srcGrid->isSequential()); nodeAcc->d_srcGrid = srcGrid; @@ -311,12 +327,14 @@ __global__ void cpyNodeCountKernel(const NanoGrid *srcGrid, //================================================================================================ -template -IndexToGrid::IndexToGrid(const SrcGridT *d_srcGrid, cudaStream_t stream) - : mStream(stream), mTimer(stream) +template +IndexToGrid::IndexToGrid(const SrcGridT *d_srcGrid, cudaStream_t stream, ResourceT& resource) + : mStream(stream), mTimer(stream), mResource(&resource), + mDevNodeAccBuf(stream, nanovdb::cuda::ResourceRef(resource), 1, nanovdb::cuda::noInit), + mGridNameBuf(stream, nanovdb::cuda::ResourceRef(resource), 0, nanovdb::cuda::noInit) { NANOVDB_ASSERT(d_srcGrid); - cudaCheck(util::cuda::mallocAsync((void**)&mDevNodeAcc, sizeof(NodeAccessor), mStream)); + mDevNodeAcc = mDevNodeAccBuf.data(); cpyNodeCountKernel<<<1, 1, 0, mStream>>>(d_srcGrid, mDevNodeAcc); cudaCheckError(); cudaCheck(cudaMemcpyAsync(&mNodeAcc, mDevNodeAcc, sizeof(NodeAccessor), cudaMemcpyDeviceToHost, mStream));// mNodeAcc = *mDevNodeAcc @@ -324,9 +342,9 @@ IndexToGrid::IndexToGrid(const SrcGridT *d_srcGrid, cudaStream_t stre //================================================================================================ -template +template template -GridHandle IndexToGrid::getHandle(const typename BuildToValueMap::type *srcValues, +GridHandle IndexToGrid::getHandle(const typename BuildToValueMap::type *srcValues, const BufferT &pool) { if (mVerbose) mTimer.start("Initiate buffer"); @@ -340,7 +358,7 @@ GridHandle IndexToGrid::getHandle(const typename BuildToValu processRootTilesKernel<<>>(mDevNodeAcc, srcValues); cudaCheckError(); - cudaCheck(util::cuda::freeAsync(mNodeAcc.d_gridName, mStream)); + mGridNameBuf.destroy(mStream); if (mVerbose) mTimer.restart("Process upper internal nodes"); processNodesKernel<<>>(mDevNodeAcc, srcValues); @@ -365,9 +383,9 @@ GridHandle IndexToGrid::getHandle(const typename BuildToValu //================================================================================================ -template +template template -inline BufferT IndexToGrid::getBuffer(const BufferT &pool) +inline BufferT IndexToGrid::getBuffer(const BufferT &pool) { mNodeAcc.grid = 0;// grid is always stored at the start of the buffer! mNodeAcc.tree = NanoGrid::memUsage(); // grid ends and tree begins @@ -394,7 +412,8 @@ inline BufferT IndexToGrid::getBuffer(const BufferT &pool) cudaCheck(cudaMemsetAsync(mNodeAcc.d_dstPtr, 0, mNodeAcc.node[0], mStream)); if (size_t size = mGridName.size()) { - cudaCheck(util::cuda::mallocAsync((void**)&mNodeAcc.d_gridName, size, mStream)); + mGridNameBuf = BufT(mStream, this->ref(), size, nanovdb::cuda::noInit); + mNodeAcc.d_gridName = mGridNameBuf.data(); cudaCheck(cudaMemcpyAsync(mNodeAcc.d_gridName, mGridName.data(), size, cudaMemcpyHostToDevice, mStream)); } else { mNodeAcc.d_gridName = nullptr; @@ -405,11 +424,11 @@ inline BufferT IndexToGrid::getBuffer(const BufferT &pool) //================================================================================================ -template +template typename util::enable_if::is_index, GridHandle>::type indexToGrid(const NanoGrid *d_srcGrid, const typename BuildToValueMap::type *d_srcValues, const BufferT &pool, cudaStream_t stream) { - IndexToGrid converter(d_srcGrid, stream); + IndexToGrid converter(d_srcGrid, stream); return converter.template getHandle(d_srcValues, pool); } diff --git a/nanovdb/nanovdb/tools/cuda/SignedFloodFill.cuh b/nanovdb/nanovdb/tools/cuda/SignedFloodFill.cuh index 69b7aa4e00..1ebda2d542 100644 --- a/nanovdb/nanovdb/tools/cuda/SignedFloodFill.cuh +++ b/nanovdb/nanovdb/tools/cuda/SignedFloodFill.cuh @@ -23,6 +23,7 @@ #define NANOVDB_TOOLS_CUDA_SIGNEDFLOODFILL_CUH_HAS_BEEN_INCLUDED #include +#include #include #include #include @@ -36,19 +37,28 @@ namespace tools::cuda { /// @brief Performs signed flood-fill operation on the hierarchical tree structure on the device /// @tparam BuildT Build type of the grid to be flood-filled +/// @tparam ResourceT Template type of optional resource used for internal temporary memory /// @param d_grid Non-const device pointer to the grid that will be flood-filled /// @param verbose If true timing information will be printed to the terminal /// @param stream optional cuda stream -template +template typename util::enable_if::is_float, void>::type signedFloodFill(NanoGrid *d_grid, bool verbose = false, cudaStream_t stream = 0); -template +template class SignedFloodFill { + static_assert(nanovdb::cuda::is_async_resource::value, + "SignedFloodFill allocates stream-ordered scratch and requires an AsyncResource"); public: - SignedFloodFill(bool verbose = false, cudaStream_t stream = 0) - : mStream(stream), mVerbose(verbose) {} + /// @brief Constructor + /// @param verbose if true timing information is printed to the terminal + /// @param stream optional CUDA stream (defaults to CUDA stream 0) + /// @param resource resource instance all device scratch is allocated from; + /// must outlive this instance (defaults to the per-type default resource) + SignedFloodFill(bool verbose = false, cudaStream_t stream = 0, + ResourceT& resource = nanovdb::cuda::default_resource()) + : mStream(stream), mVerbose(verbose), mResource(&resource) {} /// @brief Toggle on and off verbose mode /// @param on if true verbose is turned on @@ -60,6 +70,11 @@ private: cudaStream_t mStream{0}; util::cuda::Timer mTimer; bool mVerbose{false}; + ResourceT* mResource;// non-owning; all device scratch routes through this resource instance + + template + using BufT = nanovdb::cuda::Buffer>; + nanovdb::cuda::ResourceRef ref() { return nanovdb::cuda::ResourceRef(*mResource); } };// SignedFloodFill @@ -185,17 +200,18 @@ __global__ void cpyNodeCountKernel(NanoGrid *d_grid, uint64_t *d_count) //================================================================================================ -template -void SignedFloodFill::operator()(NanoGrid *d_grid) +template +void SignedFloodFill::operator()(NanoGrid *d_grid) { static_assert(BuildTraits::is_float, "cuda::SignedFloodFill only works on float grids"); NANOVDB_ASSERT(d_grid); - uint64_t count[4], *d_count = nullptr; - cudaCheck(util::cuda::mallocAsync((void**)&d_count, 4*sizeof(uint64_t), mStream)); + uint64_t count[4]; + BufT countBuf(mStream, this->ref(), 4, nanovdb::cuda::noInit); + uint64_t *d_count = countBuf.data(); kernels::cpyNodeCountKernel<<<1, 1, 0, mStream>>>(d_grid, d_count); cudaCheckError(); cudaCheck(cudaMemcpyAsync(&count, d_count, 4*sizeof(uint64_t), cudaMemcpyDeviceToHost, mStream)); - cudaCheck(util::cuda::freeAsync(d_count, mStream)); + countBuf.destroy(mStream); static const int threadsPerBlock = 128; auto blocksPerGrid = [&](size_t count)->uint32_t{return (count + (threadsPerBlock - 1)) / threadsPerBlock;}; @@ -221,11 +237,11 @@ void SignedFloodFill::operator()(NanoGrid *d_grid) //================================================================================================ -template +template typename util::enable_if::is_float, void>::type signedFloodFill(NanoGrid *d_grid, bool verbose, cudaStream_t stream) { - SignedFloodFill sff(verbose, stream); + SignedFloodFill sff(verbose, stream); sff(d_grid); auto *d_gridData = d_grid->data(); Checksum cs = getChecksum(d_gridData, stream); diff --git a/nanovdb/nanovdb/unittest/TestMemoryResource.cu b/nanovdb/nanovdb/unittest/TestMemoryResource.cu index d50d87bf22..394a38f5a1 100644 --- a/nanovdb/nanovdb/unittest/TestMemoryResource.cu +++ b/nanovdb/nanovdb/unittest/TestMemoryResource.cu @@ -4,8 +4,9 @@ /// @file TestMemoryResource.cu /// /// @brief Unit tests for the CUDA memory-resource concept (cuda::DeviceResource, -/// cuda::PinnedResource) and the cuda::TempPool / tools::cuda::PointsToGrid -/// resource plumbing. +/// cuda::PinnedResource) and the resource plumbing of cuda::TempPool and +/// the tools::cuda builders (PointsToGrid, the TopologyBuilder consumers, +/// IndexToGrid, GridStats, SignedFloodFill, addBlindData). #include #include @@ -16,6 +17,10 @@ #include #include #include +#include +#include +#include +#include #include #include @@ -330,21 +335,25 @@ TEST(TestMemoryResource, PointsToGrid_PointEncodedWithCustomResource) // injected resource instance (B5, openvdb #2232). //====================================================================== -/// @brief Build a small ValueOnIndex device grid from a handful of voxels. +/// @brief Build a small device grid of type @c BuildT 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. +template static nanovdb::GridHandle -buildIndexGrid(const std::vector& voxels) +buildGrid(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)); + nanovdb::tools::cuda::PointsToGrid converter(nanovdb::Map(1.0)); auto handle = converter.getHandle(d_voxels, voxels.size()); cudaCheck(cudaFree(d_voxels)); return handle; } +static nanovdb::GridHandle +buildIndexGrid(const std::vector& voxels) { return buildGrid(voxels); } + TEST(TestMemoryResource, DilateGrid_InjectedResourceSeam) { // The dilated grid handle's output buffer goes through BufferT::create, and the @@ -458,4 +467,102 @@ TEST(TestMemoryResource, CoarsenGrid_InjectedResourceSeam) EXPECT_EQ(res.allocs, res.deallocs); } +//====================================================================== +// Small builders (IndexToGrid, SignedFloodFill, GridStats, AddBlindData) +// route their device scratch through cuda::Buffer over an injectable +// resource instead of raw mallocAsync/freeAsync pairs (openvdb #2232). +// Their scratch is small and fixed, so the counts are exact. +//====================================================================== + +TEST(TestMemoryResource, IndexToGrid_InjectedResourceSeam) +{ + auto src = buildIndexGrid({{0,0,0},{1,2,3},{4,4,4}}); + auto* d_srcGrid = src.deviceGrid(); + ASSERT_NE(d_srcGrid, nullptr); + + float* d_values = nullptr;// one value per index; over-provisioned + cudaCheck(cudaMalloc(&d_values, 64 * sizeof(float))); + cudaCheck(cudaMemset(d_values, 0, 64 * sizeof(float))); + + CountingResource res; + { + nanovdb::tools::cuda::IndexToGrid op(d_srcGrid, 0, res); + op.setGridName("seam");// exercises the device grid-name buffer as well + auto handle = op.getHandle(d_values); + ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess); + EXPECT_TRUE(handle.deviceData()); + EXPECT_EQ(res.allocs, 2); // node accessor + grid name + } // op destroyed -> node accessor freed via res + ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess); + EXPECT_EQ(res.allocs, res.deallocs); + ASSERT_EQ(cudaFree(d_values), cudaSuccess); +} + +TEST(TestMemoryResource, SignedFloodFill_InjectedResourceSeam) +{ + auto src = buildGrid({{0,0,0},{1,2,3},{4,4,4}}); + auto* d_grid = src.deviceGrid(); + ASSERT_NE(d_grid, nullptr); + + CountingResource res; + nanovdb::tools::cuda::SignedFloodFill op(false, 0, res); + op(d_grid); + ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess); + EXPECT_EQ(res.allocs, 1); // the node-count scratch + EXPECT_EQ(res.deallocs, 1); +} + +TEST(TestMemoryResource, GridStats_InjectedResourceSeam) +{ + auto src = buildGrid({{0,0,0},{1,2,3},{4,4,4}}); + auto* d_grid = src.deviceGrid(); + ASSERT_NE(d_grid, nullptr); + + CountingResource res; + { // Stats has an average, so the per-node scratch is allocated + nanovdb::tools::cuda::GridStats, CountingResource> stats(0.0f, res); + stats.update(d_grid); + ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess); + EXPECT_EQ(res.allocs, 1); + EXPECT_EQ(res.deallocs, 1); + } + { // Extrema has no average: the zero-element buffer must not allocate + nanovdb::tools::cuda::GridStats, CountingResource> stats(0.0f, res); + stats.update(d_grid); + ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess); + EXPECT_EQ(res.allocs, 1); // unchanged + EXPECT_EQ(res.deallocs, 1); + } + { // The free function forwards ResourceT to the per-type default instance + auto& def = nanovdb::cuda::default_resource(); + const int a0 = def.allocs, d0 = def.deallocs; + nanovdb::tools::cuda::updateGridStats(d_grid, nanovdb::tools::StatsMode::All); + ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess); + EXPECT_EQ(def.allocs - a0, 1); + EXPECT_EQ(def.deallocs - d0, 1); + } +} + +TEST(TestMemoryResource, AddBlindData_InjectedResourceSeam) +{ + auto src = buildGrid({{0,0,0},{1,2,3},{4,4,4}}); + auto* d_grid = src.deviceGrid(); + ASSERT_NE(d_grid, nullptr); + + float* d_blind = nullptr; + cudaCheck(cudaMalloc(&d_blind, 8 * sizeof(float))); + cudaCheck(cudaMemset(d_blind, 0, 8 * sizeof(float))); + + // The free function routes through the per-type default resource instance. + auto& res = nanovdb::cuda::default_resource(); + const int a0 = res.allocs, d0 = res.deallocs; + auto handle = nanovdb::tools::cuda::addBlindData( + d_grid, d_blind, 8); + ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess); + EXPECT_TRUE(handle.deviceData()); + EXPECT_EQ(res.allocs - a0, 1); // the byte-size scratch + EXPECT_EQ(res.deallocs - d0, 1); + ASSERT_EQ(cudaFree(d_blind), cudaSuccess); +} + } // unnamed namespace diff --git a/pendingchanges/nanovdb.txt b/pendingchanges/nanovdb.txt index f36be8529e..2a2d6c820b 100644 --- a/pendingchanges/nanovdb.txt +++ b/pendingchanges/nanovdb.txt @@ -7,6 +7,7 @@ NanoVDB: Improvements: - The GPU builders now allocate all scratch through an injectable memory resource: tools::cuda::TopologyBuilder and tools::cuda::MeshToGrid gained a ResourceT template parameter (defaulted, so existing code is unaffected), joining PointsToGrid. Added nanovdb::cuda::SyncFromAsync, a CRTP base that derives the synchronous half of the resource concept from the stream-ordered half, and nanovdb::cuda::ResourceRef, a non-owning reference to a resource that is itself a resource, for containers that hold their resource by value. PointsToGrid's device scratch and intermediate arrays are now owned by cuda::Buffer as well, replacing all of its hand-paired allocate/free calls. Added nanovdb::cuda::MallocResource, a synchronous cudaMalloc-backed resource that works on devices without memory-pool support, and nanovdb::cuda::AsyncFromSync, which presents any synchronous resource as a stream-ordered one by synchronizing before each deallocation, so the builders can run with an injected synchronous resource. - Added the CMake option NANOVDB_CUDA_WERROR (default ON, also implied by OPENVDB_CXX_STRICT) which passes --Werror=all-warnings to NVCC so that every device-side diagnostic becomes an error when building the NanoVDB tests, tools and examples. Disable it if a newer CUDA toolkit introduces diagnostics that block your build. + - The remaining GPU tools now allocate their device scratch through an injectable memory resource as well: tools::cuda::indexToGrid, tools::cuda::updateGridStats, tools::cuda::signedFloodFill and tools::cuda::addBlindData (and their implementation classes) gained a defaulted ResourceT template parameter, and their internal mallocAsync/freeAsync pairs are replaced by RAII nanovdb::cuda::Buffer scratch. Existing code is unaffected. - The bug-fix to the nanovdb::ReadAccessor (see below) improves random-access performance in some use-cases (especially on the CPU). Fixes: From e8126f298a33a78ecee732a569888ca1bd96256a Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 19 Aug 2026 01:05:53 +0000 Subject: [PATCH 2/2] NanoVDB: document GridStats' NodeManager allocation boundary The temporary NodeManager in GridStats::update allocates through the dual-space DeviceBuffer, which cannot yet take a resource; make that visible on update() so users injecting a resource for accounting or limits know this allocation bypasses it until the single-space handle work lands. Co-Authored-By: Claude Fable 5 Signed-off-by: Mark Harris --- nanovdb/nanovdb/tools/cuda/GridStats.cuh | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/nanovdb/nanovdb/tools/cuda/GridStats.cuh b/nanovdb/nanovdb/tools/cuda/GridStats.cuh index 140e02445d..8f78646ea4 100644 --- a/nanovdb/nanovdb/tools/cuda/GridStats.cuh +++ b/nanovdb/nanovdb/tools/cuda/GridStats.cuh @@ -67,6 +67,11 @@ public: GridStats(ValueT delta = ValueT(0), ResourceT& resource = nanovdb::cuda::default_resource()) : mDelta(delta), mResource(&resource) {} + /// @note The per-node statistics scratch is allocated through the injected + /// resource. The temporary NodeManager this method builds still + /// allocates through the dual-space DeviceBuffer, which does not yet + /// accept a resource; that allocation bypasses @c ResourceT until the + /// single-space handle work on the roadmap (openvdb #2232) lands. void update(GridT *d_grid, cudaStream_t stream = 0); }; // cuda::GridStats