diff --git a/nanovdb/nanovdb/cuda/Buffer.h b/nanovdb/nanovdb/cuda/Buffer.h index 41d15d511a..d9144591ca 100644 --- a/nanovdb/nanovdb/cuda/Buffer.h +++ b/nanovdb/nanovdb/cuda/Buffer.h @@ -50,7 +50,7 @@ struct StreamHolder { cudaStream_t mStream = 0; }; /// is_resource). When @c R provides both interfaces the stream-ordered /// one is used. /// @details With a stream-ordered resource the Buffer retains the stream of -/// the most recent allocation (or the one supplied via setStream) +/// the most recent allocation (or the one supplied via set_stream) /// and orders its deallocation on that stream. Buffer is move-only. template class Buffer : private detail::StreamHolder::value> @@ -133,7 +133,7 @@ class Buffer : private detail::StreamHolder::value> Buffer& operator=(Buffer&& other) noexcept { if (this != &other) { - this->clear(); + this->destroy(); static_cast&>(*this) = other; mResource = std::move(other.mResource); mData = other.mData; @@ -168,7 +168,7 @@ class Buffer : private detail::StreamHolder::value> /// @brief D-tor. A stream-ordered resource frees on the retained stream; /// a synchronous resource frees immediately. - ~Buffer() { this->clear(); } + ~Buffer() { this->destroy(); } /// @brief Returns the retained stream, i.e. the stream the buffer's memory /// will be freed on. @@ -179,9 +179,11 @@ class Buffer : private detail::StreamHolder::value> /// deallocation (and destruction) is ordered on @c stream instead. /// @param stream cuda stream subsequent deallocation is ordered on /// @warning The caller is responsible for ordering @c stream after any - /// in-flight work that uses the buffer's memory. + /// in-flight work that uses the buffer's memory. This deliberately + /// does not synchronize, matching cuda::buffer's + /// set_stream_unsynchronized rather than its set_stream. template::value, int> = 0> - void setStream(cudaStream_t stream) { this->mStream = stream; } + void set_stream(cudaStream_t stream) { this->mStream = stream; } /// @brief Resizes the buffer to @c count elements, preserving the leading /// min(old, new) elements. Every operation — the new allocation, the @@ -219,7 +221,7 @@ class Buffer : private detail::StreamHolder::value> mSize = count; } else { - this->mStream = stream; // no reallocation: setStream semantics + this->mStream = stream; // no reallocation: set_stream semantics } } @@ -263,13 +265,46 @@ class Buffer : private detail::StreamHolder::value> bool empty() const { return mSize == 0; } /// @brief Frees the buffer memory (if any) and resets to the empty state. - void clear() + /// A stream-ordered resource frees on the retained stream. + /// @note Spelled destroy to match cuda::buffer. This is the name to use. + void destroy() { this->deallocate(mData, mSize); mData = nullptr; mSize = 0; } + /// @brief Frees the buffer memory (if any) and resets to the empty state. + /// @note Transitional, and not the name to use: it exists only because + /// GridHandle::reset still calls clear() on its buffer. It goes away + /// when the legacy dual buffers do and GridHandle moves to destroy(). + void clear() { this->destroy(); } + + /// @brief Frees the buffer memory (if any) on @c stream and resets to the + /// empty state. @c stream becomes the retained stream. + /// @param stream cuda stream the deallocation is ordered on + /// @warning The caller is responsible for ordering @c stream after any + /// in-flight work that uses the buffer's memory. + template::value, int> = 0> + void destroy(cudaStream_t stream) + { + this->mStream = stream; + this->destroy(); + } + + /// @brief Exchanges the contents of this buffer with @c other. Neither + /// buffer allocates, frees, or copies element data. + /// @param other buffer to exchange contents with + void swap(Buffer& other) noexcept + { + auto& lhs = static_cast&>(*this); + auto& rhs = static_cast&>(other); + std::swap(lhs, rhs); + std::swap(mResource, other.mResource); + std::swap(mData, other.mData); + std::swap(mSize, other.mSize); + } + private: /// @brief Returns @c count * sizeof(T), throwing std::runtime_error if the /// byte size would overflow size_t instead of silently wrapping into diff --git a/nanovdb/nanovdb/tools/cuda/CoarsenGrid.cuh b/nanovdb/nanovdb/tools/cuda/CoarsenGrid.cuh index d22e5462b4..f0686eac84 100644 --- a/nanovdb/nanovdb/tools/cuda/CoarsenGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/CoarsenGrid.cuh @@ -210,7 +210,7 @@ void CoarsenGrid::coarsenInternalNodes() if (auto srcLeafCount = mSrcTreeData.mNodeCount[0]) { // Unless it's an empty grid util::cuda::lambdaKernel<<>>( srcLeafCount, util::morphology::cuda::CoarsenInternalNodesFunctor(), - mDeviceSrcGrid, mBuilder.deviceProcessedRoot(), mBuilder.mUpperMasks.deviceData(), mBuilder.mLowerMasks.deviceData() ); + mDeviceSrcGrid, mBuilder.deviceProcessedRoot(), mBuilder.mUpperMasks.data(), mBuilder.mLowerMasks.data() ); } }// CoarsenGrid::coarsenInternalNodes diff --git a/nanovdb/nanovdb/tools/cuda/PruneGrid.cuh b/nanovdb/nanovdb/tools/cuda/PruneGrid.cuh index 113acb6837..cb445c23ec 100644 --- a/nanovdb/nanovdb/tools/cuda/PruneGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/PruneGrid.cuh @@ -213,7 +213,7 @@ void PruneGrid::pruneInternalNodes() if (auto srcLeafCount = mSrcTreeData.mNodeCount[0]) { // Unless it's an empty grid util::cuda::lambdaKernel<<>>( srcLeafCount, util::morphology::cuda::PruneInternalNodesFunctor(), - mDeviceSrcGrid, mBuilder.deviceProcessedRoot(), mDeviceSrcLeafMask, mBuilder.mUpperMasks.deviceData(), mBuilder.mLowerMasks.deviceData() ); + mDeviceSrcGrid, mBuilder.deviceProcessedRoot(), mDeviceSrcLeafMask, mBuilder.mUpperMasks.data(), mBuilder.mLowerMasks.data() ); } }// PruneGrid::pruneInternalNodes diff --git a/nanovdb/nanovdb/tools/cuda/RefineGrid.cuh b/nanovdb/nanovdb/tools/cuda/RefineGrid.cuh index d611629abd..07704aa16d 100644 --- a/nanovdb/nanovdb/tools/cuda/RefineGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/RefineGrid.cuh @@ -225,7 +225,7 @@ void RefineGrid::refineInternalNodes() if (auto srcLeafCount = mSrcTreeData.mNodeCount[0]) { // Unless it's an empty grid util::cuda::lambdaKernel<<>>( srcLeafCount, util::morphology::cuda::RefineInternalNodesFunctor(), - mDeviceSrcGrid, mBuilder.deviceProcessedRoot(), mBuilder.mUpperMasks.deviceData(), mBuilder.mLowerMasks.deviceData() ); + mDeviceSrcGrid, mBuilder.deviceProcessedRoot(), mBuilder.mUpperMasks.data(), mBuilder.mLowerMasks.data() ); } }// RefineGrid::refineInternalNodes diff --git a/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh b/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh index 6ccbb25e6e..13eb4714f4 100644 --- a/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh +++ b/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh @@ -17,6 +17,8 @@ #include #include +#include +#include #include #include @@ -24,7 +26,25 @@ namespace nanovdb { namespace tools::cuda { +/// @brief Shared grid/tree offsets and node counts handed to the device +/// functors. Independent of the resource the builder allocates from, +/// so it lives outside TopologyBuilder and stays one type across every +/// ResourceT instantiation. template +struct TopologyBuilderData { + void *d_bufferPtr; + uint64_t grid, tree, root, upper, lower, leaf, size;// byte offsets to nodes in buffer + uint32_t nodeCount[3];// 0=leaf,1=lower, 2=upper + uint32_t *d_upperOffsets; + __hostdev__ NanoGrid& getGrid() const {return *util::PtrAdd>(d_bufferPtr, grid);} + __hostdev__ NanoTree& getTree() const {return *util::PtrAdd>(d_bufferPtr, tree);} + __hostdev__ NanoRoot& getRoot() const {return *util::PtrAdd>(d_bufferPtr, root);} + __hostdev__ NanoUpper& getUpper(int i) const {return *(util::PtrAdd>(d_bufferPtr, upper)+i);} + __hostdev__ NanoLower& getLower(int i) const {return *(util::PtrAdd>(d_bufferPtr, lower)+i);} + __hostdev__ NanoLeaf& getLeaf(int i) const {return *(util::PtrAdd>(d_bufferPtr, leaf)+i);} +};// TopologyBuilderData + +template class TopologyBuilder { static_assert(nanovdb::BuildTraits::is_onindex);// For now, only OnIndexGrids supported @@ -36,25 +56,25 @@ class TopologyBuilder using LowerT = NanoLower; using LeafT = NanoLeaf; + /// @brief Device-only scratch storage, allocated from the injected + /// resource. These buffers are never read on the host, so they use + /// the single-space Buffer rather than the dual DeviceBuffer, whose + /// host pointer and per-device array they would leave unused. + using ScratchT = nanovdb::cuda::Buffer; + public: - TopologyBuilder(cudaStream_t stream) + /// @param stream cuda stream the scratch allocations are ordered on + /// @param resource resource instance all device scratch is allocated from; + /// must outlive this builder + TopologyBuilder(cudaStream_t stream, ResourceT& resource = nanovdb::cuda::default_resource()) + : mResource(&resource) + , mTempDevicePool(resource) { mData = nanovdb::cuda::DeviceBuffer::create(sizeof(Data)); } - struct Data { - void *d_bufferPtr; - uint64_t grid, tree, root, upper, lower, leaf, size;// byte offsets to nodes in buffer - uint32_t nodeCount[3];// 0=leaf,1=lower, 2=upper - uint32_t *d_upperOffsets; - __hostdev__ GridT& getGrid() const {return *util::PtrAdd(d_bufferPtr, grid);} - __hostdev__ TreeT& getTree() const {return *util::PtrAdd(d_bufferPtr, tree);} - __hostdev__ RootT& getRoot() const {return *util::PtrAdd(d_bufferPtr, root);} - __hostdev__ UpperT& getUpper(int i) const {return *(util::PtrAdd(d_bufferPtr, upper)+i);} - __hostdev__ LowerT& getLower(int i) const {return *(util::PtrAdd(d_bufferPtr, lower)+i);} - __hostdev__ LeafT& getLeaf(int i) const {return *(util::PtrAdd(d_bufferPtr, leaf)+i);} - };// Data + using Data = TopologyBuilderData; void allocateInternalMaskBuffers(cudaStream_t stream); @@ -74,21 +94,21 @@ public: void postProcessGridTree(cudaStream_t stream); nanovdb::cuda::DeviceBuffer mProcessedRoot; - nanovdb::cuda::DeviceBuffer mUpperMasks; - nanovdb::cuda::DeviceBuffer mLowerMasks; - nanovdb::cuda::DeviceBuffer mUpperOffsets; - nanovdb::cuda::DeviceBuffer mLowerOffsets; - nanovdb::cuda::DeviceBuffer mLeafOffsets; - nanovdb::cuda::DeviceBuffer mVoxelOffsets; - nanovdb::cuda::DeviceBuffer mLowerParents; - nanovdb::cuda::DeviceBuffer mLeafParents; + ScratchT mUpperMasks; + ScratchT mLowerMasks; + ScratchT mUpperOffsets; + ScratchT mLowerOffsets; + ScratchT mLeafOffsets; + ScratchT mVoxelOffsets; + ScratchT mLowerParents; + ScratchT mLeafParents; nanovdb::cuda::DeviceBuffer mData; CheckMode mChecksum{CheckMode::Disable}; auto deviceProcessedRoot() { return static_cast(mProcessedRoot.deviceData()); } auto hostProcessedRoot() { return static_cast(mProcessedRoot.data()); } - void* deviceUpperMasks() { return mUpperMasks.deviceData(); } - void* deviceLowerMasks() { return mLowerMasks.deviceData(); } + void* deviceUpperMasks() { return mUpperMasks.data(); } + void* deviceLowerMasks() { return mLowerMasks.data(); } Data* data() { return static_cast(mData.data()); } Data* deviceData() { return static_cast(mData.deviceData()); } @@ -96,8 +116,9 @@ private: static constexpr unsigned int mNumThreads = 128;// for kernels spawned via lambdaKernel (others may specialize) static unsigned int numBlocks(unsigned int n) {return (n + mNumThreads - 1) / mNumThreads;} - nanovdb::cuda::TempDevicePool mTempDevicePool; -};// tools::cuda::TopologyBuilder + ResourceT* mResource;// non-owning; all device scratch routes through this instance + nanovdb::cuda::TempPool mTempDevicePool; +};// tools::cuda::TopologyBuilder //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- @@ -118,30 +139,28 @@ private: //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template -void TopologyBuilder::allocateInternalMaskBuffers(cudaStream_t stream) +template +void TopologyBuilder::allocateInternalMaskBuffers(cudaStream_t stream) { if (hostProcessedRoot()->tileCount() == 0) return; // Processing empty grid(s); nothing to allocate // Allocate (and zero-fill) buffers large enough to hold: // (a) The serialized masks of all upper nodes, for all tiles in the updated root node, and // (b) The serialized masks of all densified lower nodes, as if every upper node had a full set of 32^3 lower children - int device = 0; - cudaGetDevice(&device); uint64_t upperSize = hostProcessedRoot()->tileCount() * sizeof(Mask<5>); uint64_t lowerSize = hostProcessedRoot()->tileCount() * Mask<5>::SIZE * sizeof(Mask<4>); - mUpperMasks = nanovdb::cuda::DeviceBuffer::create(upperSize, nullptr, device, stream); - if (mUpperMasks.deviceData() == nullptr) throw std::runtime_error("Failed to allocate upper mask buffer on device"); - cudaCheck(cudaMemsetAsync(mUpperMasks.deviceData(), 0, upperSize, stream)); - mLowerMasks = nanovdb::cuda::DeviceBuffer::create( lowerSize, nullptr, device, stream ); - if (mLowerMasks.deviceData() == nullptr) throw std::runtime_error("Failed to allocate lower mask buffer on device"); - cudaCheck(cudaMemsetAsync(mLowerMasks.deviceData(), 0, lowerSize, stream)); -}// TopologyBuilder::allocateInternalMaskBuffers + mUpperMasks = ScratchT(stream, *mResource, upperSize, nanovdb::cuda::noInit); + if (mUpperMasks.data() == nullptr) throw std::runtime_error("Failed to allocate upper mask buffer on device"); + cudaCheck(cudaMemsetAsync(mUpperMasks.data(), 0, upperSize, stream)); + mLowerMasks = ScratchT(stream, *mResource, lowerSize, nanovdb::cuda::noInit); + if (mLowerMasks.data() == nullptr) throw std::runtime_error("Failed to allocate lower mask buffer on device"); + cudaCheck(cudaMemsetAsync(mLowerMasks.data(), 0, lowerSize, stream)); +}// TopologyBuilder::allocateInternalMaskBuffers //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template -void TopologyBuilder::countNodes(cudaStream_t stream) +template +void TopologyBuilder::countNodes(cudaStream_t stream) { auto processedTileCount = hostProcessedRoot()->tileCount(); if (processedTileCount == 0) { // Processing empty grid(s); zero nodes at all levels @@ -155,59 +174,57 @@ void TopologyBuilder::countNodes(cudaStream_t stream) // as well as the tile table at the root. std::size_t size = processedTileCount*Mask<5>::SIZE; - int device = 0; - cudaGetDevice(&device); - nanovdb::cuda::DeviceBuffer upperCountsBuffer = nanovdb::cuda::DeviceBuffer::create(processedTileCount*sizeof(uint32_t), nullptr, device, stream); - nanovdb::cuda::DeviceBuffer lowerCountsBuffer = nanovdb::cuda::DeviceBuffer::create(size*sizeof(uint32_t), nullptr, device, stream); - nanovdb::cuda::DeviceBuffer leafCountsBuffer = nanovdb::cuda::DeviceBuffer::create(size*sizeof(uint32_t), nullptr, device, stream); + ScratchT upperCountsBuffer = ScratchT(stream, *mResource, processedTileCount*sizeof(uint32_t), nanovdb::cuda::noInit); + ScratchT lowerCountsBuffer = ScratchT(stream, *mResource, size*sizeof(uint32_t), nanovdb::cuda::noInit); + ScratchT leafCountsBuffer = ScratchT(stream, *mResource, size*sizeof(uint32_t), nanovdb::cuda::noInit); using CountType = uint32_t (*)[Mask<5>::SIZE]; - auto lowerCounts = reinterpret_cast( lowerCountsBuffer.deviceData() ); - auto leafCounts = reinterpret_cast( leafCountsBuffer.deviceData() ); + auto lowerCounts = reinterpret_cast(lowerCountsBuffer.data()); + auto leafCounts = reinterpret_cast(leafCountsBuffer.data()); using Op = util::morphology::cuda::EnumerateNodesFunctor; util::cuda::operatorKernel <<>> (deviceUpperMasks(), deviceLowerMasks(), lowerCounts, leafCounts); - mUpperOffsets = nanovdb::cuda::DeviceBuffer::create((processedTileCount+1)*sizeof(uint32_t), nullptr, device, stream); - mLowerOffsets = nanovdb::cuda::DeviceBuffer::create((size+1)*sizeof(uint32_t), nullptr, device, stream); - mLeafOffsets = nanovdb::cuda::DeviceBuffer::create((size+1)*sizeof(uint32_t), nullptr, device, stream); + mUpperOffsets = ScratchT(stream, *mResource, (processedTileCount+1)*sizeof(uint32_t), nanovdb::cuda::noInit); + mLowerOffsets = ScratchT(stream, *mResource, (size+1)*sizeof(uint32_t), nanovdb::cuda::noInit); + mLeafOffsets = ScratchT(stream, *mResource, (size+1)*sizeof(uint32_t), nanovdb::cuda::noInit); - cudaCheck(cudaMemsetAsync(mLowerOffsets.deviceData(), 0, sizeof(uint32_t), stream)); + cudaCheck(cudaMemsetAsync(mLowerOffsets.data(), 0, sizeof(uint32_t), stream)); CALL_CUBS(DeviceScan::InclusiveSum, - static_cast(lowerCountsBuffer.deviceData()), - static_cast(mLowerOffsets.deviceData())+1, + reinterpret_cast(lowerCountsBuffer.data()), + reinterpret_cast(mLowerOffsets.data())+1, size); - cudaCheck(cudaMemcpyAsync(&data()->nodeCount[1], static_cast(mLowerOffsets.deviceData())+size, sizeof(uint32_t), cudaMemcpyDeviceToHost, stream)); + cudaCheck(cudaMemcpyAsync(&data()->nodeCount[1], reinterpret_cast(mLowerOffsets.data())+size, sizeof(uint32_t), cudaMemcpyDeviceToHost, stream)); - cudaCheck(cudaMemsetAsync(mLeafOffsets.deviceData(), 0, sizeof(uint32_t), stream)); + cudaCheck(cudaMemsetAsync(mLeafOffsets.data(), 0, sizeof(uint32_t), stream)); CALL_CUBS(DeviceScan::InclusiveSum, - static_cast(leafCountsBuffer.deviceData()), - static_cast(mLeafOffsets.deviceData())+1, + reinterpret_cast(leafCountsBuffer.data()), + reinterpret_cast(mLeafOffsets.data())+1, size); - cudaCheck(cudaMemcpyAsync(&data()->nodeCount[0], static_cast(mLeafOffsets.deviceData())+size, sizeof(uint32_t), cudaMemcpyDeviceToHost, stream)); + cudaCheck(cudaMemcpyAsync(&data()->nodeCount[0], reinterpret_cast(mLeafOffsets.data())+size, sizeof(uint32_t), cudaMemcpyDeviceToHost, stream)); util::cuda::lambdaKernel<<>>( processedTileCount, [] __device__(size_t tileID, CountType lowerOffsets, uint32_t* upperCounts) { upperCounts[tileID] = (lowerOffsets[tileID+1][0] > lowerOffsets[tileID][0]) ? 1 : 0; }, - static_cast(mLowerOffsets.deviceData()), - static_cast(upperCountsBuffer.deviceData())); + reinterpret_cast(mLowerOffsets.data()), + reinterpret_cast(upperCountsBuffer.data())); - cudaCheck(cudaMemsetAsync( mUpperOffsets.deviceData(), 0, sizeof(uint32_t), stream)); + cudaCheck(cudaMemsetAsync( mUpperOffsets.data(), 0, sizeof(uint32_t), stream)); CALL_CUBS(DeviceScan::InclusiveSum, - static_cast(upperCountsBuffer.deviceData()), - static_cast(mUpperOffsets.deviceData())+1, + reinterpret_cast(upperCountsBuffer.data()), + reinterpret_cast(mUpperOffsets.data())+1, processedTileCount); - cudaCheck(cudaMemcpyAsync(&data()->nodeCount[2], static_cast(mUpperOffsets.deviceData())+processedTileCount, sizeof(uint32_t), cudaMemcpyDeviceToHost, stream)); -}// TopologyBuilder::countNodes + cudaCheck(cudaMemcpyAsync(&data()->nodeCount[2], reinterpret_cast(mUpperOffsets.data())+processedTileCount, sizeof(uint32_t), cudaMemcpyDeviceToHost, stream)); +}// TopologyBuilder::countNodes //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template +template template -BufferT TopologyBuilder::getBuffer(const BufferT &pool, cudaStream_t stream) +BufferT TopologyBuilder::getBuffer(const BufferT &pool, cudaStream_t stream) { // Allocates a device buffer for the destination grid, once the topology/size of the tree is known data()->grid = 0;// grid is always stored at the start of the buffer! @@ -226,11 +243,11 @@ BufferT TopologyBuilder::getBuffer(const BufferT &pool, cudaStream_t str data()->d_bufferPtr = buffer.deviceData(); if (data()->d_bufferPtr == nullptr) throw std::runtime_error("Failed to allocate grid buffer on the device"); if (data()->nodeCount[2] != 0) // Unless the result is an empty grid - data()->d_upperOffsets = static_cast(mUpperOffsets.deviceData()); + data()->d_upperOffsets = reinterpret_cast(mUpperOffsets.data()); mData.deviceUpload(device, stream, false); return buffer; -}// TopologyBuilder::getBuffer +}// TopologyBuilder::getBuffer //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- @@ -240,7 +257,7 @@ template struct BuildGridTreeRootFunctor { __device__ - void operator()(size_t, typename TopologyBuilder::Data *d_data) { + void operator()(size_t, TopologyBuilderData *d_data) { // process Root auto &root = d_data->getRoot(); @@ -320,7 +337,7 @@ struct InitGridTreeRootFunctor Map map; // transform to embed in the output grid __device__ - void operator()(size_t, typename TopologyBuilder::Data *d_data) { + void operator()(size_t, TopologyBuilderData *d_data) { // process Root (identical to BuildGridTreeRootFunctor) auto &root = d_data->getRoot(); @@ -389,7 +406,7 @@ template struct BuildUpperNodesFunctor { __device__ - void operator()(size_t processedTileID, typename TopologyBuilder::Data *d_data, NanoRoot *d_processedRoot) { + void operator()(size_t processedTileID, TopologyBuilderData *d_data, NanoRoot *d_processedRoot) { uint32_t tileID = d_data->d_upperOffsets[processedTileID]; if (tileID != d_data->d_upperOffsets[processedTileID+1]) // if the offsets are the same, this was a speculatively introduced tile which was not necessary { @@ -406,8 +423,8 @@ struct BuildUpperNodesFunctor }// namespace topology::detail -template -inline void TopologyBuilder::processUpperNodes(cudaStream_t stream) +template +inline void TopologyBuilder::processUpperNodes(cudaStream_t stream) { // Connect all newly allocated upper nodes to their respective tiles // Also fill in any necessary part of the preamble (in InternalData) of upper nodes @@ -418,12 +435,12 @@ inline void TopologyBuilder::processUpperNodes(cudaStream_t stream) processedTileCount, topology::detail::BuildUpperNodesFunctor(), deviceData(), deviceProcessedRoot()); cudaCheckError(); } -}// TopologyBuilder::processUpperNodes +}// TopologyBuilder::processUpperNodes //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template -inline void TopologyBuilder::processLowerNodes(cudaStream_t stream) +template +inline void TopologyBuilder::processLowerNodes(cudaStream_t stream) { // Fill out the contents of all newly allocated lower nodes (using the densified upper/lower mask arrays) // Also fill in the preamble (most of LeafData) for their leaf children @@ -431,34 +448,32 @@ inline void TopologyBuilder::processLowerNodes(cudaStream_t stream) using CountType = uint32_t (*)[Mask<5>::SIZE]; if (processedTileCount) { // Unless output grid is empty - int device = 0; - cudaGetDevice(&device); std::size_t lowerCount = data()->nodeCount[1]; - mLowerParents = nanovdb::cuda::DeviceBuffer::create(lowerCount*sizeof(uint32_t), nullptr, device, stream); + mLowerParents = ScratchT(stream, *mResource, lowerCount*sizeof(uint32_t), nanovdb::cuda::noInit); std::size_t leafCount = data()->nodeCount[0]; - mLeafParents = nanovdb::cuda::DeviceBuffer::create(leafCount*sizeof(uint32_t), nullptr, device, stream); + mLeafParents = ScratchT(stream, *mResource, leafCount*sizeof(uint32_t), nanovdb::cuda::noInit); using Op = util::morphology::cuda::ProcessLowerNodesFunctor; util::cuda::operatorKernel <<>>( deviceUpperMasks(), deviceLowerMasks(), - static_cast(mUpperOffsets.deviceData()), - static_cast(mLowerOffsets.deviceData()), - static_cast(mLeafOffsets.deviceData()), + reinterpret_cast(mUpperOffsets.data()), + reinterpret_cast(mLowerOffsets.data()), + reinterpret_cast(mLeafOffsets.data()), static_cast(data()->d_bufferPtr), - static_cast(mLowerParents.deviceData()), - static_cast(mLeafParents.deviceData()) + reinterpret_cast(mLowerParents.data()), + reinterpret_cast(mLeafParents.data()) ); cudaCheckError(); } mProcessedRoot.clear(stream); - mUpperMasks.clear(stream); - mLowerMasks.clear(stream); - mLowerOffsets.clear(stream); - mLeafOffsets.clear(stream); -}// TopologyBuilder::processLowerNodes + mUpperMasks.destroy(stream); + mLowerMasks.destroy(stream); + mLowerOffsets.destroy(stream); + mLeafOffsets.destroy(stream); +}// TopologyBuilder::processLowerNodes //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- @@ -468,7 +483,7 @@ template struct UpdateLeafVoxelCountsAndPrefixSumFunctor { __device__ - void operator()(size_t leafID, typename TopologyBuilder::Data *d_data, uint64_t *d_voxelCounts) { + void operator()(size_t leafID, TopologyBuilderData *d_data, uint64_t *d_voxelCounts) { auto &leaf = d_data->getGrid().tree().template getFirstNode<0>()[leafID]; const uint64_t *w = leaf.mValueMask.words(); uint64_t prefixSum = 0, sum = util::countOn(*w++); @@ -485,32 +500,30 @@ template struct UpdateLeafVoxelOffsetsFunctor { __device__ - void operator()(size_t leafID, typename TopologyBuilder::Data *d_data, uint64_t *d_voxelOffsets) { + void operator()(size_t leafID, TopologyBuilderData *d_data, uint64_t *d_voxelOffsets) { auto &leaf = d_data->getGrid().tree().template getFirstNode<0>()[leafID]; leaf.mOffset = d_voxelOffsets[leafID]+1; } }; }// namespace topology::detail -template -inline void TopologyBuilder::processLeafOffsets(cudaStream_t stream) +template +inline void TopologyBuilder::processLeafOffsets(cudaStream_t stream) { - int device = 0; - cudaGetDevice(&device); std::size_t leafCount = data()->nodeCount[0]; if (leafCount) { // Unless output grid is empty - mVoxelOffsets = nanovdb::cuda::DeviceBuffer::create((leafCount+1)*sizeof(uint64_t), nullptr, device, stream); - cudaCheck(cudaMemsetAsync(mVoxelOffsets.deviceData(), 0, sizeof(uint64_t), stream)); + mVoxelOffsets = ScratchT(stream, *mResource, (leafCount+1)*sizeof(uint64_t), nanovdb::cuda::noInit); + cudaCheck(cudaMemsetAsync(mVoxelOffsets.data(), 0, sizeof(uint64_t), stream)); util::cuda::lambdaKernel<<>>( - leafCount, topology::detail::UpdateLeafVoxelCountsAndPrefixSumFunctor(), deviceData(), static_cast(mVoxelOffsets.deviceData())+1); + leafCount, topology::detail::UpdateLeafVoxelCountsAndPrefixSumFunctor(), deviceData(), reinterpret_cast(mVoxelOffsets.data())+1); CALL_CUBS(DeviceScan::InclusiveSum, - static_cast(mVoxelOffsets.deviceData())+1, - static_cast(mVoxelOffsets.deviceData())+1, + reinterpret_cast(mVoxelOffsets.data())+1, + reinterpret_cast(mVoxelOffsets.data())+1, leafCount); util::cuda::lambdaKernel<<>>( - leafCount, topology::detail::UpdateLeafVoxelOffsetsFunctor(), deviceData(), static_cast(mVoxelOffsets.deviceData())); + leafCount, topology::detail::UpdateLeafVoxelOffsetsFunctor(), deviceData(), reinterpret_cast(mVoxelOffsets.data())); } -}// TopologyBuilder::processLeafOffsets +}// TopologyBuilder::processLeafOffsets //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- @@ -527,7 +540,7 @@ template struct UpdateAndPropagateLeafBBoxFunctor { __device__ - void operator()(size_t tid, typename TopologyBuilder::Data *d_data, const uint32_t* leafParents) { + void operator()(size_t tid, TopologyBuilderData *d_data, const uint32_t* leafParents) { auto &lower = d_data->getLower(leafParents[tid]); auto &leaf = d_data->getLeaf(tid); leaf.updateBBox(); @@ -539,7 +552,7 @@ template struct PropagateLowerBBoxFunctor { __device__ - void operator()(size_t tid, typename TopologyBuilder::Data *d_data, const uint32_t* lowerParents) { + void operator()(size_t tid, TopologyBuilderData *d_data, const uint32_t* lowerParents) { auto &upper = d_data->getUpper(lowerParents[tid]); auto &lower = d_data->getLower(tid); upper.mBBox.expandAtomic(lower.bbox()); } @@ -549,7 +562,7 @@ template struct PropagateUpperBBoxFunctor { __device__ - void operator()(size_t tid, typename TopologyBuilder::Data *d_data) { + void operator()(size_t tid, TopologyBuilderData *d_data) { d_data->getRoot().mBBox.expandAtomic(d_data->getUpper(tid).bbox()); } }; @@ -558,7 +571,7 @@ template struct UpdateRootWorldBBoxFunctor { __device__ - void operator()(size_t tid, typename TopologyBuilder::Data *d_data) { + void operator()(size_t tid, TopologyBuilderData *d_data) { // TODO: check that the correct semantics are followed in this transformation auto BBox = d_data->getRoot().mBBox; BBox.max() += 1; @@ -569,8 +582,8 @@ struct UpdateRootWorldBBoxFunctor }// namespace topology::detail -template -inline void TopologyBuilder::processBBox(cudaStream_t stream) +template +inline void TopologyBuilder::processBBox(cudaStream_t stream) { if (data()->nodeCount[0] == 0) return; // Output grid is empty; retain empty bounding box @@ -578,14 +591,14 @@ inline void TopologyBuilder::processBBox(cudaStream_t stream) // update and propagate bbox from leaf -> lower/parent nodes util::cuda::lambdaKernel<<nodeCount[0]), mNumThreads, 0, stream>>>( - data()->nodeCount[0], topology::detail::UpdateAndPropagateLeafBBoxFunctor(), deviceData(), static_cast(mLeafParents.deviceData())); - mLeafParents.clear(stream); + data()->nodeCount[0], topology::detail::UpdateAndPropagateLeafBBoxFunctor(), deviceData(), reinterpret_cast(mLeafParents.data())); + mLeafParents.destroy(stream); cudaCheckError(); // propagate bbox from lower -> upper/parent node util::cuda::lambdaKernel<<nodeCount[1]), mNumThreads, 0, stream>>>( - data()->nodeCount[1], topology::detail::PropagateLowerBBoxFunctor(), deviceData(), static_cast(mLowerParents.deviceData())); - mLowerParents.clear(stream); + data()->nodeCount[1], topology::detail::PropagateLowerBBoxFunctor(), deviceData(), reinterpret_cast(mLowerParents.data())); + mLowerParents.destroy(stream); cudaCheckError(); // propagate bbox from upper -> root/parent node @@ -595,7 +608,7 @@ inline void TopologyBuilder::processBBox(cudaStream_t stream) // update the world-bbox in the root node util::cuda::lambdaKernel<<<1, 1, 0, stream>>>(1, topology::detail::UpdateRootWorldBBoxFunctor(), deviceData()); cudaCheckError(); -}// TopologyBuilder::processBBox +}// TopologyBuilder::processBBox //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- @@ -605,7 +618,7 @@ template struct PostProcessGridTreeFunctor { __device__ - void operator()(size_t tid, typename TopologyBuilder::Data *d_data, uint64_t* d_voxelOffsets) { + void operator()(size_t tid, TopologyBuilderData *d_data, uint64_t* d_voxelOffsets) { auto& grid = d_data->getGrid(); auto& tree = grid.tree(); auto leafCount = tree.mNodeCount[0]; @@ -616,17 +629,17 @@ struct PostProcessGridTreeFunctor }// namespace topology::detail -template -inline void TopologyBuilder::postProcessGridTree(cudaStream_t stream) +template +inline void TopologyBuilder::postProcessGridTree(cudaStream_t stream) { // Finish updates to GridData/TreeData and (optionally) update checksum if (data()->nodeCount[0]) // if grid is empty, the default values are correct - util::cuda::lambdaKernel<<<1, 1, 0, stream>>>(1, topology::detail::PostProcessGridTreeFunctor(), deviceData(), static_cast(mVoxelOffsets.deviceData())); + util::cuda::lambdaKernel<<<1, 1, 0, stream>>>(1, topology::detail::PostProcessGridTreeFunctor(), deviceData(), reinterpret_cast(mVoxelOffsets.data())); cudaCheckError(); - mVoxelOffsets.clear(stream); + mVoxelOffsets.destroy(stream); tools::cuda::updateChecksum((GridData*)data()->d_bufferPtr, mChecksum, stream); -}// TopologyBuilder::postProcessGridTree +}// TopologyBuilder::postProcessGridTree //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- diff --git a/nanovdb/nanovdb/unittest/TestBuffer.cu b/nanovdb/nanovdb/unittest/TestBuffer.cu index 6a871802c6..40be6878a6 100644 --- a/nanovdb/nanovdb/unittest/TestBuffer.cu +++ b/nanovdb/nanovdb/unittest/TestBuffer.cu @@ -258,9 +258,65 @@ TEST(TestBuffer, ClearFreesAndEmpties) ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess); } +TEST(TestBuffer, DestroyIsTheSpellingClearDelegatesTo) +{ + Counters c; + nanovdb::cuda::Buffer buf(0, CountingResource{&c}, 64, nanovdb::cuda::noInit); + ASSERT_EQ(c.allocs, 1); + buf.destroy(); // cuda::buffer's spelling + EXPECT_EQ(buf.data(), nullptr); + EXPECT_EQ(buf.size(), 0u); + EXPECT_EQ(c.deallocs, 1); + buf.destroy(); // idempotent + EXPECT_EQ(c.deallocs, 1); + ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess); +} + +TEST(TestBuffer, DestroyOnStreamRetargetsTheFree) +{ + cudaStream_t a, b; + ASSERT_EQ(cudaStreamCreate(&a), cudaSuccess); + ASSERT_EQ(cudaStreamCreate(&b), cudaSuccess); + { + StreamLog log; + nanovdb::cuda::Buffer buf(a, StreamRecordingResource{&log}, 32, nanovdb::cuda::noInit); + ASSERT_EQ(log.allocStreams.size(), 1u); + EXPECT_EQ(log.allocStreams[0], a); // allocated on a + buf.destroy(b); // explicit stream overload + ASSERT_EQ(log.deallocStreams.size(), 1u); + EXPECT_EQ(log.deallocStreams[0], b); // freed on b, not the retained stream a + EXPECT_EQ(buf.stream(), b); // b is retained afterwards + } + ASSERT_EQ(cudaStreamSynchronize(a), cudaSuccess); + ASSERT_EQ(cudaStreamSynchronize(b), cudaSuccess); + cudaStreamDestroy(a); + cudaStreamDestroy(b); +} + +TEST(TestBuffer, SwapExchangesWithoutAllocatingOrFreeing) +{ + Counters c; + nanovdb::cuda::Buffer x(0, CountingResource{&c}, 128, nanovdb::cuda::noInit); + nanovdb::cuda::Buffer y(0, CountingResource{&c}, 64, nanovdb::cuda::noInit); + ASSERT_EQ(c.allocs, 2); + auto* px = x.data(); + auto* py = y.data(); + const int allocs = c.allocs, deallocs = c.deallocs; + + x.swap(y); + + EXPECT_EQ(c.allocs, allocs); // no allocation + EXPECT_EQ(c.deallocs, deallocs);// no free + EXPECT_EQ(x.data(), py); + EXPECT_EQ(y.data(), px); + EXPECT_EQ(x.size(), 64u); + EXPECT_EQ(y.size(), 128u); + ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess); +} + //====================================================================== // Stream retention: the destructor and resize free on the retained stream -// (stream of the most recent allocation, or the one supplied via setStream) +// (stream of the most recent allocation, or the one supplied via set_stream) //====================================================================== TEST(TestBuffer, DestructorFreesOnAllocationStream) @@ -288,7 +344,7 @@ TEST(TestBuffer, SetStreamRedirectsTheFree) StreamLog log; { nanovdb::cuda::Buffer buf(a, StreamRecordingResource{&log}, 32, nanovdb::cuda::noInit); - buf.setStream(b); // member update only, no synchronization + buf.set_stream(b); // member update only, no synchronization EXPECT_EQ(buf.stream(), b); } ASSERT_EQ(log.allocStreams.size(), 1u); @@ -473,7 +529,7 @@ TEST(TestBuffer, AsyncPathIsGraphCapturable) template struct HasSetStream : std::false_type {}; template -struct HasSetStream().setStream(cudaStream_t{0}))>> : std::true_type {}; +struct HasSetStream().set_stream(cudaStream_t{0}))>> : std::true_type {}; template struct HasStreamGetter : std::false_type {}; @@ -484,9 +540,9 @@ using PinnedBufferF = nanovdb::cuda::Buffer; // A Buffer over a synchronous resource exposes no stream API at all. -static_assert(!HasSetStream::value, "sync-resource Buffer must not expose setStream"); +static_assert(!HasSetStream::value, "sync-resource Buffer must not expose set_stream"); static_assert(!HasStreamGetter::value, "sync-resource Buffer must not expose stream()"); -static_assert(HasSetStream::value, "async-resource Buffer exposes setStream"); +static_assert(HasSetStream::value, "async-resource Buffer exposes set_stream"); static_assert(HasStreamGetter::value, "async-resource Buffer exposes stream()"); TEST(TestBuffer, PinnedBufferIsPageLocked) diff --git a/pendingchanges/nanovdb.txt b/pendingchanges/nanovdb.txt index 8e3bdbc0ca..f342f5c47a 100644 --- a/pendingchanges/nanovdb.txt +++ b/pendingchanges/nanovdb.txt @@ -2,6 +2,7 @@ NanoVDB: New Features: - Added new _hostdev_ function named nanovdb::math::isoCrossing, which intersects a ray against a user-defined iso-surface. + - Added nanovdb::cuda::Buffer and nanovdb::cuda::BufferView (CUDA): a typed, resource-aware, stream-ordered container that allocates from an injectable memory resource and frees on its retained stream, and a non-owning view over externally managed memory that a GridHandle can wrap without copying. Member names follow cuda::buffer (destroy, set_stream, swap). Also added the synchronous resource concept nanovdb::cuda::is_resource alongside is_async_resource. Improvements: - The bug-fix to the nanovdb::ReadAccessor (see below) improves random-access performance in some use-cases (especially on the CPU).