Skip to content

NanoVDB: expose ResourceT on the five TopologyBuilder consumers (CUDA) #2273

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
swahtz merged 13 commits into from
Aug 18, 2026
Merged

NanoVDB: expose ResourceT on the five TopologyBuilder consumers (CUDA) #2273

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
13 commits
Select commit Hold shift + click to select a range
04fd4bc
NanoVDB: align cuda::Buffer member names with cuda::buffer
harrism Aug 5, 2026
a5ad1fc
NanoVDB: allocate TopologyBuilder scratch from an injected resource
harrism Aug 5, 2026
1e69333
NanoVDB: add SyncFromAsync, and give MeshToGrid a resource seam
harrism Aug 5, 2026
35d61dd
NanoVDB: add ResourceRef and route TempPool's scratch through cuda::B…
harrism Aug 5, 2026
fb05c71
NanoVDB: note cuda::Buffer and cuda::BufferView in pendingchanges
harrism Aug 5, 2026
c7302bf
NanoVDB: free TopologyBuilder scratch on the caller's stream
harrism Aug 5, 2026
f90cb26
Merge branch 'nanovdb-buffer-retrofit' into nanovdb-temppool-buffer
harrism Aug 5, 2026
f619056
NanoVDB: borrow TopologyBuilder scratch through ResourceRef
harrism Aug 5, 2026
2608ea8
NanoVDB: assert the scratch alignment TopologyBuilder relies on
harrism Aug 5, 2026
5fda0ac
NanoVDB: set_stream matches cuda::buffer's contract, not a divergence
harrism Aug 5, 2026
2db0993
Merge branch 'nanovdb-buffer-retrofit' into nanovdb-temppool-buffer
harrism Aug 5, 2026
0c174b9
NanoVDB: expose ResourceT on the five TopologyBuilder consumers
swahtz Aug 7, 2026
b292512
Merge branch 'master' into nanovdb-consumer-resource-seams
swahtz Aug 18, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
46 changes: 26 additions & 20 deletions nanovdb/nanovdb/tools/cuda/CoarsenGrid.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,9 +30,12 @@ namespace nanovdb {

namespace tools::cuda {

template <typename BuildT>
template <typename BuildT, typename ResourceT = nanovdb::cuda::DeviceResource>
class CoarsenGrid
{
static_assert(nanovdb::cuda::is_async_resource<ResourceT>::value,
"CoarsenGrid allocates stream-ordered scratch and requires an AsyncResource");

using GridT = NanoGrid<BuildT>;
using TreeT = NanoTree<BuildT>;
using RootT = NanoRoot<BuildT>;
Expand All @@ -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<ResourceT>())
: 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
Expand Down Expand Up @@ -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<BuildT> mBuilder;
TopologyBuilder<BuildT, ResourceT> mBuilder;
cudaStream_t mStream{0};
util::cuda::Timer mTimer;
int mVerbose{0};
const GridT *mDeviceSrcGrid;
TreeData mSrcTreeData;
};// tools::cuda::CoarsenGrid<BuildT>
};// tools::cuda::CoarsenGrid<BuildT, ResourceT>

//-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------

template<typename BuildT>
template<typename BuildT, typename ResourceT>
template<typename BufferT>
GridHandle<BufferT>
CoarsenGrid<BuildT>::getHandle(const BufferT &pool)
CoarsenGrid<BuildT, ResourceT>::getHandle(const BufferT &pool)
{
// Copy TreeData from GPU -> CPU
cudaStreamSynchronize(mStream);
Expand Down Expand Up @@ -147,12 +153,12 @@ CoarsenGrid<BuildT>::getHandle(const BufferT &pool)
cudaStreamSynchronize(mStream);

return GridHandle<BufferT>(std::move(buffer));
}// CoarsenGrid<BuildT>::getHandle
}// CoarsenGrid<BuildT, ResourceT>::getHandle

//-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------

template<typename BuildT>
void CoarsenGrid<BuildT>::coarsenRoot()
template<typename BuildT, typename ResourceT>
void CoarsenGrid<BuildT, ResourceT>::coarsenRoot()
{
// This method coarsens the root tiles, to accommodate for the overall downsamping operation.

Expand Down Expand Up @@ -197,12 +203,12 @@ void CoarsenGrid<BuildT>::coarsenRoot()
for (const auto& [key, tile] : coarsenedTiles)
*coarsenedRootPtr->tile(t++) = tile;
mBuilder.mProcessedRoot.deviceUpload(device, mStream, false);
}// CoarsenGrid<BuildT>::coarsenRoot
}// CoarsenGrid<BuildT, ResourceT>::coarsenRoot

//-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------

template<typename BuildT>
void CoarsenGrid<BuildT>::coarsenInternalNodes()
template<typename BuildT, typename ResourceT>
void CoarsenGrid<BuildT, ResourceT>::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,
Expand All @@ -212,24 +218,24 @@ void CoarsenGrid<BuildT>::coarsenInternalNodes()
srcLeafCount, util::morphology::cuda::CoarsenInternalNodesFunctor<BuildT>(),
mDeviceSrcGrid, mBuilder.deviceProcessedRoot(), mBuilder.mUpperMasks.data(), mBuilder.mLowerMasks.data() );
}
}// CoarsenGrid<BuildT>::coarsenInternalNodes
}// CoarsenGrid<BuildT, ResourceT>::coarsenInternalNodes

//-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------

template <typename BuildT>
void CoarsenGrid<BuildT>::processGridTreeRoot()
template <typename BuildT, typename ResourceT>
void CoarsenGrid<BuildT, ResourceT>::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<BuildT>(), mBuilder.deviceData());
cudaCheckError();
}// CoarsenGrid<BuildT>::processGridTreeRoot
}// CoarsenGrid<BuildT, ResourceT>::processGridTreeRoot

//-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------

template<typename BuildT>
void CoarsenGrid<BuildT>::coarsenLeafNodes()
template<typename BuildT, typename ResourceT>
void CoarsenGrid<BuildT, ResourceT>::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.
Expand All @@ -241,7 +247,7 @@ void CoarsenGrid<BuildT>::coarsenLeafNodes()

// Update leaf offsets and prefix sums
mBuilder.processLeafOffsets(mStream);
}// CoarsenGrid<BuildT>::coarsenLeafNodes
}// CoarsenGrid<BuildT, ResourceT>::coarsenLeafNodes

//-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------

Expand Down
46 changes: 26 additions & 20 deletions nanovdb/nanovdb/tools/cuda/DilateGrid.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,9 +30,12 @@ namespace nanovdb {

namespace tools::cuda {

template <typename BuildT>
template <typename BuildT, typename ResourceT = nanovdb::cuda::DeviceResource>
class DilateGrid
{
static_assert(nanovdb::cuda::is_async_resource<ResourceT>::value,
"DilateGrid allocates stream-ordered scratch and requires an AsyncResource");

using GridT = NanoGrid<BuildT>;
using TreeT = NanoTree<BuildT>;
using RootT = NanoRoot<BuildT>;
Expand All @@ -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<ResourceT>())
: 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
Expand Down Expand Up @@ -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<BuildT> mBuilder;
TopologyBuilder<BuildT, ResourceT> 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<BuildT>
};// tools::cuda::DilateGrid<BuildT, ResourceT>

//-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------

template<typename BuildT>
template<typename BuildT, typename ResourceT>
template<typename BufferT>
GridHandle<BufferT>
DilateGrid<BuildT>::getHandle(const BufferT &pool)
DilateGrid<BuildT, ResourceT>::getHandle(const BufferT &pool)
{
// Copy TreeData from GPU -> CPU
cudaStreamSynchronize(mStream);
Expand Down Expand Up @@ -152,12 +158,12 @@ DilateGrid<BuildT>::getHandle(const BufferT &pool)
cudaStreamSynchronize(mStream);

return GridHandle<BufferT>(std::move(buffer));
}// DilateGrid<BuildT>::getHandle
}// DilateGrid<BuildT, ResourceT>::getHandle

//-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------

template<typename BuildT>
void DilateGrid<BuildT>::dilateRoot()
template<typename BuildT, typename ResourceT>
void DilateGrid<BuildT, ResourceT>::dilateRoot()
{
// This method conservatively and speculatively dilates the root tiles, to accommodate
// any new root nodes that might be introduced by the dilation operation.
Expand Down Expand Up @@ -220,12 +226,12 @@ void DilateGrid<BuildT>::dilateRoot()
for (const auto& [key, tile] : dilatedTiles)
*dilatedRootPtr->tile(t++) = tile;
mBuilder.mProcessedRoot.deviceUpload(device, mStream, false);
}// DilateGrid<BuildT>::dilateRoot
}// DilateGrid<BuildT, ResourceT>::dilateRoot

//-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------

template<typename BuildT>
void DilateGrid<BuildT>::dilateInternalNodes()
template<typename BuildT, typename ResourceT>
void DilateGrid<BuildT, ResourceT>::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,
Expand All @@ -247,24 +253,24 @@ void DilateGrid<BuildT>::dilateInternalNodes()
<<<dim3(mSrcTreeData.mNodeCount[1],Op::SlicesPerLowerNode,1), Op::MaxThreadsPerBlock, 0, mStream>>>
(mDeviceSrcGrid, mBuilder.deviceProcessedRoot(), mBuilder.deviceUpperMasks(), mBuilder.deviceLowerMasks()); }
}
}// DilateGrid<BuildT>::dilateInternalNodes
}// DilateGrid<BuildT, ResourceT>::dilateInternalNodes

//-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------

template <typename BuildT>
void DilateGrid<BuildT>::processGridTreeRoot()
template <typename BuildT, typename ResourceT>
void DilateGrid<BuildT, ResourceT>::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<BuildT>(), mBuilder.deviceData());
cudaCheckError();
}// DilateGrid<BuildT>::processGridTreeRoot
}// DilateGrid<BuildT, ResourceT>::processGridTreeRoot

//-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------

template<typename BuildT>
void DilateGrid<BuildT>::dilateLeafNodes()
template<typename BuildT, typename ResourceT>
void DilateGrid<BuildT, ResourceT>::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.
Expand All @@ -285,7 +291,7 @@ void DilateGrid<BuildT>::dilateLeafNodes()

// Update leaf offsets and prefix sums
mBuilder.processLeafOffsets(mStream);
}// DilateGrid<BuildT>::dilateLeafNodes
}// DilateGrid<BuildT, ResourceT>::dilateLeafNodes

//-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------

Expand Down
53 changes: 31 additions & 22 deletions nanovdb/nanovdb/tools/cuda/MergeGrids.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -33,9 +33,12 @@ namespace nanovdb {

namespace tools::cuda {

template <typename BuildT>
template <typename BuildT, typename ResourceT = nanovdb::cuda::DeviceResource>
class MergeGrids
{
static_assert(nanovdb::cuda::is_async_resource<ResourceT>::value,
"MergeGrids allocates stream-ordered scratch and requires an AsyncResource");

using GridT = NanoGrid<BuildT>;
using TreeT = NanoTree<BuildT>;
using RootT = NanoRoot<BuildT>;
Expand All @@ -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<const GridT*>& 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<const GridT*>& d_srcGrids, cudaStream_t stream = 0,
ResourceT& resource = nanovdb::cuda::default_resource<ResourceT>())
: 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<const GridT*>{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<ResourceT>())
: MergeGrids(std::vector<const GridT*>{d_srcGrid1, d_srcGrid2}, stream, resource) {}

/// @brief Toggle on and off verbose mode
/// @param level Verbose level: 0=quiet, 1=timing, 2=benchmarking
Expand Down Expand Up @@ -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<BuildT> mBuilder;
TopologyBuilder<BuildT, ResourceT> mBuilder;
cudaStream_t mStream{0};
util::cuda::Timer mTimer;
int mVerbose{0};
std::vector<const GridT*> mDeviceSrcGrids;
std::vector<TreeData> mSrcTreeData;
};// tools::cuda::MergeGrids<BuildT>
};// tools::cuda::MergeGrids<BuildT, ResourceT>

//-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------

template<typename BuildT>
template<typename BuildT, typename ResourceT>
template<typename BufferT>
GridHandle<BufferT>
MergeGrids<BuildT>::getHandle(const BufferT &pool)
MergeGrids<BuildT, ResourceT>::getHandle(const BufferT &pool)
{
if (mDeviceSrcGrids.empty())
throw std::runtime_error("MergeGrids: no input grids");
Expand Down Expand Up @@ -163,12 +172,12 @@ MergeGrids<BuildT>::getHandle(const BufferT &pool)
cudaStreamSynchronize(mStream);

return GridHandle<BufferT>(std::move(buffer));
}// MergeGrids<BuildT>::getHandle
}// MergeGrids<BuildT, ResourceT>::getHandle

//-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------

template<typename BuildT>
void MergeGrids<BuildT>::mergeRoot()
template<typename BuildT, typename ResourceT>
void MergeGrids<BuildT, ResourceT>::mergeRoot()
{
// Creates a new merged tree root with the merged tiles of the two input root topologies

Expand Down Expand Up @@ -218,12 +227,12 @@ void MergeGrids<BuildT>::mergeRoot()
for (const auto& [key, tile] : mergedTiles)
*mergedRootPtr->tile(t++) = tile;
mBuilder.mProcessedRoot.deviceUpload(device, mStream, false);
}// MergeGrids<BuildT>::mergeRoot
}// MergeGrids<BuildT, ResourceT>::mergeRoot

//-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------

template<typename BuildT>
void MergeGrids<BuildT>::mergeInternalNodes()
template<typename BuildT, typename ResourceT>
void MergeGrids<BuildT, ResourceT>::mergeInternalNodes()
{
// Merges the masks of upper and lower nodes from both input topologies into the
// densified, pre-allocated mask arrays of the merged result
Expand All @@ -236,25 +245,25 @@ void MergeGrids<BuildT>::mergeInternalNodes()
<<<mSrcTreeData[i].mNodeCount[1], Op::MaxThreadsPerBlock, 0, mStream>>>
(mDeviceSrcGrids[i], mBuilder.deviceProcessedRoot(), mBuilder.deviceUpperMasks(), mBuilder.deviceLowerMasks());
}
}// MergeGrids<BuildT>::mergeInternalNodes
}// MergeGrids<BuildT, ResourceT>::mergeInternalNodes

//-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------

template <typename BuildT>
void MergeGrids<BuildT>::processGridTreeRoot()
template <typename BuildT, typename ResourceT>
void MergeGrids<BuildT, ResourceT>::processGridTreeRoot()
{
// Copy GridData from the first source grid
// TODO: Check for instances where extra processing is needed
// TODO: check that the other grid inputs have consistent GridData, too
cudaCheck(cudaMemcpyAsync(&mBuilder.data()->getGrid(), mDeviceSrcGrids.front()->data(), GridT::memUsage(), cudaMemcpyDeviceToDevice, mStream));
util::cuda::lambdaKernel<<<1, 1, 0, mStream>>>(1, topology::detail::BuildGridTreeRootFunctor<BuildT>(), mBuilder.deviceData());
cudaCheckError();
}// MergeGrids<BuildT>::processGridTreeRoot
}// MergeGrids<BuildT, ResourceT>::processGridTreeRoot

//-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------

template<typename BuildT>
void MergeGrids<BuildT>::mergeLeafNodes()
template<typename BuildT, typename ResourceT>
void MergeGrids<BuildT, ResourceT>::mergeLeafNodes()
{
using Op = util::morphology::cuda::MergeLeafNodesFunctor<BuildT>;
// Each input ORs its leaf active masks into the merged leaf topology.
Expand All @@ -267,7 +276,7 @@ void MergeGrids<BuildT>::mergeLeafNodes()

// Update leaf offsets and prefix sums
mBuilder.processLeafOffsets(mStream);
}// MergeGrids<BuildT>::mergeLeafNodes
}// MergeGrids<BuildT, ResourceT>::mergeLeafNodes

//-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------

Expand Down
Loading
Loading