Skip to content

NanoVDB: own the small builders' device scratch with cuda::Buffer (CUDA) #2286

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 2 commits into from
Aug 19, 2026
Merged

NanoVDB: own the small builders' device scratch with cuda::Buffer (CUDA) #2286

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
06afbf0
NanoVDB: own the small builders' device scratch with cuda::Buffer (CUDA)
harrism Aug 18, 2026
e8126f2
NanoVDB: document GridStats' NodeManager allocation boundary
harrism Aug 19, 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
15 changes: 11 additions & 4 deletions nanovdb/nanovdb/tools/cuda/AddBlindData.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,6 +18,7 @@
#define NANOVDB_TOOLS_CUDA_ADDBLINDDATA_CUH_HAS_BEEN_INCLUDED

#include <nanovdb/NanoVDB.h>
#include <nanovdb/cuda/Buffer.h>
#include <nanovdb/cuda/DeviceBuffer.h>
#include <nanovdb/GridHandle.h>
#include <nanovdb/util/cuda/Util.h>
Expand All @@ -34,6 +35,7 @@ namespace tools::cuda {// ============================================
/// @tparam BuildT Build type of the grid
/// @tparam BlindDataT Type of the blind data
/// @tparam BufferT Type of the buffer used for allocation
/// @tparam ResourceT Template type of optional resource used for internal temporary memory
/// @param d_grid Pointer to device grid
/// @param d_blindData Pointer to device blind data
/// @param valueCount number of values in the blind data
Expand All @@ -43,7 +45,7 @@ namespace tools::cuda {// ============================================
/// @param pool optional pool used for allocation
/// @param stream optional CUDA stream (defaults to CUDA stream 0)
/// @return GridHandle with blind data appended
template<typename BuildT, typename BlindDataT, typename BufferT = nanovdb::cuda::DeviceBuffer>
template<typename BuildT, typename BlindDataT, typename BufferT = nanovdb::cuda::DeviceBuffer, typename ResourceT = nanovdb::cuda::DeviceResource>
GridHandle<BufferT>
addBlindData(const NanoGrid<BuildT> *d_grid,
const BlindDataT *d_blindData,
Expand All @@ -60,11 +62,16 @@ addBlindData(const NanoGrid<BuildT> *d_grid,
// old grid old meta new meta old data new data

static_assert(BufferTraits<BufferT>::hasDeviceDual, "Expected BufferT to support device allocation");
static_assert(nanovdb::cuda::is_async_resource<ResourceT>::value,
"addBlindData allocates stream-ordered scratch and requires an AsyncResource");

// extract byte sizes of the grid, blind meta data and blind data
enum {GRID=0, META=1, DATA=2, CHECKSUM=3};
uint64_t tmp[4], *d_tmp;
cudaCheck(util::cuda::mallocAsync((void**)&d_tmp, 4*sizeof(uint64_t), stream));
using ScratchT = nanovdb::cuda::Buffer<uint64_t, nanovdb::cuda::ResourceRef<ResourceT>>;
auto &resource = nanovdb::cuda::default_resource<ResourceT>();
uint64_t tmp[4];
ScratchT tmpBuf(stream, nanovdb::cuda::ResourceRef<ResourceT>(resource), 4, nanovdb::cuda::noInit);
uint64_t *d_tmp = tmpBuf.data();
util::cuda::lambdaKernel<<<1, 1, 0, stream>>>(1, [=] __device__(size_t) {
if (auto count = d_grid->blindDataCount()) {
d_tmp[GRID] = util::PtrDiff(&d_grid->blindMetaData(0), d_grid);
Expand Down Expand Up @@ -116,7 +123,7 @@ addBlindData(const NanoGrid<BuildT> *d_grid,
for (uint32_t i=0, n=grid.mBlindMetadataCount-1; i<n; ++i, ++meta) meta->mDataOffset += sizeof(GridBlindMetaData);
grid.mGridSize += sizeof(GridBlindMetaData) + meta->blindDataSize();// expansion with 32 byte alignment
}); cudaCheckError();
cudaCheck(util::cuda::freeAsync(d_tmp, stream));
tmpBuf.destroy(stream);

Checksum cs(tmp[CHECKSUM]);
cuda::updateChecksum(reinterpret_cast<GridData*>(d_data), cs.mode(), stream);
Expand Down
55 changes: 38 additions & 17 deletions nanovdb/nanovdb/tools/cuda/GridStats.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,6 +16,7 @@
#define NANOVDB_TOOLS_CUDA_GRIDSTATS_CUH_HAS_BEEN_INCLUDED

#include <nanovdb/NanoVDB.h>
#include <nanovdb/cuda/Buffer.h>
#include <nanovdb/tools/GridStats.h>

#include <cub/util_type.cuh>// for cub::Uninitialized
Expand All @@ -26,18 +27,22 @@ namespace tools::cuda {

/// @brief Update, i.e. re-compute, grid statistics like min/max, stats and bbox
/// information for an existing NanoVDB Grid.
/// @param grid Grid whose stats to update
/// @tparam ResourceT Template type of optional resource used for internal temporary memory
/// @param d_grid Grid whose stats to update
/// @param mode Mode of computation for the statistics.
/// @param stream Optional cuda stream (defaults to zero)
template<typename BuildT>
template<typename BuildT, typename ResourceT = nanovdb::cuda::DeviceResource>
void updateGridStats(NanoGrid<BuildT> *d_grid, StatsMode mode = StatsMode::Default, cudaStream_t stream = 0);

//================================================================================================

/// @brief Allows for the construction of NanoVDB grids without any dependency
template<typename BuildT, typename StatsT = Stats<typename NanoGrid<BuildT>::ValueType>>
/// @brief Re-computes statistics (of type @c StatsT) for each node of an existing NanoVDB grid on the device
template<typename BuildT, typename StatsT = Stats<typename NanoGrid<BuildT>::ValueType>, typename ResourceT = nanovdb::cuda::DeviceResource>
class GridStats
{
static_assert(nanovdb::cuda::is_async_resource<ResourceT>::value,
"GridStats allocates stream-ordered scratch and requires an AsyncResource");

using GridT = NanoGrid<BuildT>;
using TreeT = typename GridT::TreeType;
using ValueT = typename TreeT::ValueType;
Expand All @@ -48,10 +53,25 @@ class GridStats
static_assert(util::is_same<ValueT, typename StatsT::ValueType>::value, "Mismatching type");

ValueT mDelta; // skip rendering of node if: node.max < -mDelta || node.min > mDelta
ResourceT* mResource;// non-owning; all device scratch routes through this resource instance

public:
GridStats(ValueT delta = ValueT(0)) : mDelta(delta) {}
template<typename T>
using BufT = nanovdb::cuda::Buffer<T, nanovdb::cuda::ResourceRef<ResourceT>>;
nanovdb::cuda::ResourceRef<ResourceT> ref() { return nanovdb::cuda::ResourceRef<ResourceT>(*mResource); }

public:
/// @brief Constructor
/// @param delta skip rendering of nodes when node.max < -delta or node.min > delta
/// @param resource resource instance all device scratch is allocated from;
/// must outlive this instance (defaults to the per-type default resource)
GridStats(ValueT delta = ValueT(0), ResourceT& resource = nanovdb::cuda::default_resource<ResourceT>())
: mDelta(delta), mResource(&resource) {}

/// @note The per-node statistics scratch is allocated through the injected
/// resource. The temporary NodeManager this method builds still
/// allocates through the dual-space DeviceBuffer, which does not yet
/// accept a resource; that allocation bypasses @c ResourceT until the
/// single-space handle work on the roadmap (openvdb #2232) lands.
void update(GridT *d_grid, cudaStream_t stream = 0);

}; // cuda::GridStats
Expand Down Expand Up @@ -285,8 +305,8 @@ __global__ void processRootAndGrid(NodeManager<BuildT> *d_nodeMgr, StatsT *d_sta

//================================================================================================

template<typename BuildT, typename StatsT>
void GridStats<BuildT, StatsT>::update(NanoGrid<BuildT> *d_grid, cudaStream_t stream)
template<typename BuildT, typename StatsT, typename ResourceT>
void GridStats<BuildT, StatsT, ResourceT>::update(NanoGrid<BuildT> *d_grid, cudaStream_t stream)
{
static const uint32_t threadsPerBlock = 128;
auto blocksPerGrid = [&](uint32_t count)->uint32_t{return (count + (threadsPerBlock - 1)) / threadsPerBlock;};
Expand All @@ -298,11 +318,12 @@ void GridStats<BuildT, StatsT>::update(NanoGrid<BuildT> *d_grid, cudaStream_t st
cudaCheck(cudaMemcpyAsync(nodeCount, (char*)d_grid + sizeof(GridData) + 4*sizeof(uint64_t), 3*sizeof(uint32_t), cudaMemcpyDeviceToHost, stream));
//cudaStreamSynchronize(stream);// finish all device tasks in stream

StatsT *d_stats = nullptr;

// One d_stats slot per node (leaves, then lower, then upper) so every node
// has its own slot - see processInternal.
if constexpr(StatsT::hasAverage()) cudaCheck(util::cuda::mallocAsync((void**)&d_stats, (nodeCount[0]+nodeCount[1]+nodeCount[2])*sizeof(StatsT), stream));
// has its own slot - see processInternal. Zero elements when the mode has
// no average, so nothing is allocated.
BufT<StatsT> statsBuf(stream, this->ref(),
StatsT::hasAverage() ? nodeCount[0]+nodeCount[1]+nodeCount[2] : 0, nanovdb::cuda::noInit);
StatsT *d_stats = statsBuf.data();

// warp per leaf (4 warps per 128-thread block); block per internal node
if (nodeCount[0]) processLeaf<BuildT><<<blocksPerGrid(nodeCount[0]*32), threadsPerBlock, 0, stream>>>(d_nodeMgr, d_stats);
Expand All @@ -313,25 +334,25 @@ void GridStats<BuildT, StatsT>::update(NanoGrid<BuildT> *d_grid, cudaStream_t st

processRootAndGrid<BuildT><<<1, 1, 0, stream>>>(d_nodeMgr, d_stats);

if constexpr(StatsT::hasAverage()) cudaCheck(util::cuda::freeAsync(d_stats, stream));
statsBuf.destroy(stream);

} // cuda::GridStats::update( Grid )

//================================================================================================

template<typename BuildT>
template<typename BuildT, typename ResourceT>
void updateGridStats(NanoGrid<BuildT> *d_grid, StatsMode mode, cudaStream_t stream)
{
if (d_grid == nullptr && mode == StatsMode::Disable) {
return;
} else if (mode == StatsMode::BBox || util::is_same<bool, BuildT>::value) {
GridStats<BuildT, NoopStats<BuildT> > stats;
GridStats<BuildT, NoopStats<BuildT>, ResourceT> stats;
stats.update(d_grid, stream);
} else if (mode == StatsMode::MinMax) {
GridStats<BuildT, Extrema<BuildT> > stats;
GridStats<BuildT, Extrema<BuildT>, ResourceT> stats;
stats.update(d_grid, stream);
} else if (mode == StatsMode::All) {
GridStats<BuildT, Stats<BuildT> > stats;
GridStats<BuildT, Stats<BuildT>, ResourceT> stats;
stats.update(d_grid, stream);
} else {
throw std::runtime_error("GridStats: Unsupported statistics mode.");
Expand Down
Loading
Loading