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 01f9da9b67..5c1bb36e5e 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.