From 5b192c996776feb82cd983ba02c13c3bdb2b4d75 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 5 Aug 2026 03:53:01 +0000 Subject: [PATCH 1/4] NanoVDB: express PointsToGrid's density search as a loop The bisection search over voxel size was written as a backward goto, which made the lifetimes of the buffers it retries over non-lexical. Rewrite it as while(true) with continue on retry and break on convergence; the six hand-written frees before the jump are unchanged. The change is easiest to review with whitespace ignored, since the loop body re-indents: git diff -w shows 27 changed lines. d_keys and d_node_count carry results past the loop, so their declarations move above it, as does the copy event, which was created inside the retried region on every iteration but destroyed only once at the end -- each retry leaked the previous handle. Co-Authored-By: Claude Opus 5 (1M context) Signed-off-by: Mark Harris --- nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh | 199 +++++++++++--------- 1 file changed, 105 insertions(+), 94 deletions(-) diff --git a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh index 4ac4ed386d..3fb555851a 100644 --- a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh @@ -610,112 +610,123 @@ void PointsToGrid::countNodes(const PtrT points, size_t point bool operator<(const Foo &rhs) const {return density < rhs.density || (density == rhs.density && dx < rhs.dx);} } min{0.0, 1}, max{0.0, 0};// min: as dx -> 0 density -> 1 point per voxel, max: density is 0 i.e. undefined -jump:// this marks the beginning of the actual algorithm - - mData.d_keys = static_cast(mResource->allocate_async(pointCount*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); - mData.d_indx = static_cast(mResource->allocate_async(pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream));// uint32_t can index 4.29 billion Coords, corresponding to 48 GB - cudaCheck(cudaMemcpyAsync(mDeviceData, &mData, sizeof(PointsToGridData), cudaMemcpyHostToDevice, mStream));// copy mData from CPU -> GPU + // Declared ahead of the search loop below: d_keys and d_node_count carry + // results past it, and the event is recorded and re-used across iterations + // (previously it was re-created per iteration, leaking the prior handle). + uint64_t *d_keys = nullptr; + uint32_t *d_indx = nullptr, *d_points_per_tile = nullptr, *d_node_count = nullptr; + cudaEvent_t copyEvent; + cudaCheck(cudaEventCreate(©Event)); - if (mVerbose==2) mTimer.start("\nAllocating arrays for keys and indices"); - auto *d_keys = static_cast(mResource->allocate_async(pointCount*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); - auto *d_indx = static_cast(mResource->allocate_async(pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + // Bisection search for the voxel size dx that yields the target point + // density: each iteration builds tile and voxel keys at the current dx, + // then either converges or frees this iteration's buffers and retries. + while (true) { - if (mVerbose==2) mTimer.restart("Generate tile keys"); - util::cuda::lambdaKernel<<>>(pointCount, TileKeyFunctor(), mDeviceData, points, d_keys, d_indx); - cudaCheckError(); - if (mVerbose==2) mTimer.restart("DeviceRadixSort of "+std::to_string(pointCount)+" tile keys"); - CALL_CUBS(DeviceRadixSort::SortPairs, d_keys, mData.d_keys, d_indx, mData.d_indx, pointCount, 0, 63);// 21 bits per coord - std::swap(d_indx, mData.d_indx);// sorted indices are now in d_indx + mData.d_keys = static_cast(mResource->allocate_async(pointCount*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + mData.d_indx = static_cast(mResource->allocate_async(pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream));// uint32_t can index 4.29 billion Coords, corresponding to 48 GB + cudaCheck(cudaMemcpyAsync(mDeviceData, &mData, sizeof(PointsToGridData), cudaMemcpyHostToDevice, mStream));// copy mData from CPU -> GPU - if (mVerbose==2) mTimer.restart("Allocate runs"); - auto *d_points_per_tile = static_cast(mResource->allocate_async(pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); - uint32_t *d_node_count = static_cast(mResource->allocate_async(3*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + if (mVerbose==2) mTimer.start("\nAllocating arrays for keys and indices"); + d_keys = static_cast(mResource->allocate_async(pointCount*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + d_indx = static_cast(mResource->allocate_async(pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); - if (mVerbose==2) mTimer.restart("DeviceRunLengthEncode tile keys"); - CALL_CUBS(DeviceRunLengthEncode::Encode, mData.d_keys, d_keys, d_points_per_tile, d_node_count+2, pointCount); - cudaCheck(cudaMemcpyAsync(mData.nodeCount+2, d_node_count+2, sizeof(uint32_t), cudaMemcpyDeviceToHost, mStream)); - cudaCheck(cudaStreamSynchronize(mStream)); - mData.d_tile_keys = static_cast(mResource->allocate_async(mData.nodeCount[2]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); - cudaCheck(cudaMemcpyAsync(mData.d_tile_keys, d_keys, mData.nodeCount[2]*sizeof(uint64_t), cudaMemcpyDeviceToDevice, mStream)); - - static constexpr uint32_t SEGMENTED_SORT_TILE_THRESHOLD = 32; - if (mData.nodeCount[2] >= SEGMENTED_SORT_TILE_THRESHOLD) { - // Bulk segmented sort: one kernel launch + one segmented radix sort (faster for many tiles) - if (mVerbose==2) mTimer.restart("Segmented radix sort of " + std::to_string(pointCount) + " voxel keys in " + std::to_string(mData.nodeCount[2]) + " tiles"); - auto *d_tile_offsets = static_cast(mResource->allocate_async((mData.nodeCount[2]+1)*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); - cudaCheck(cudaMemsetAsync(d_tile_offsets, 0, sizeof(uint32_t), mStream)); - CALL_CUBS(DeviceScan::InclusiveSum, d_points_per_tile, d_tile_offsets + 1, mData.nodeCount[2]); - mResource->deallocate_async(d_points_per_tile, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - - util::cuda::lambdaKernel<<>>(pointCount, BulkVoxelKeyFunctor(), mDeviceData, points, d_tile_offsets, mData.nodeCount[2], d_keys, d_indx, uint32_t(0)); + if (mVerbose==2) mTimer.restart("Generate tile keys"); + util::cuda::lambdaKernel<<>>(pointCount, TileKeyFunctor(), mDeviceData, points, d_keys, d_indx); cudaCheckError(); - CALL_CUBS(DeviceSegmentedRadixSort::SortPairs, d_keys, mData.d_keys, d_indx, mData.d_indx, (int)pointCount, (int)mData.nodeCount[2], d_tile_offsets, d_tile_offsets + 1, 0, 36); - mResource->deallocate_async(d_tile_offsets, (mData.nodeCount[2]+1)*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - } else { - // Serial per-tile sort: individual kernel + sort per tile (lower overhead for few tiles) - if (mVerbose==2) mTimer.restart("DeviceRadixSort of " + std::to_string(pointCount) + " voxel keys in " + std::to_string(mData.nodeCount[2]) + " tiles"); - uint32_t *points_per_tile = new uint32_t[mData.nodeCount[2]]; - cudaCheck(cudaMemcpyAsync(points_per_tile, d_points_per_tile, mData.nodeCount[2]*sizeof(uint32_t), cudaMemcpyDeviceToHost, mStream)); - mResource->deallocate_async(d_points_per_tile, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - for (uint32_t id = 0, offset = 0; id < mData.nodeCount[2]; ++id) { - const uint32_t count = points_per_tile[id]; - util::cuda::offsetLambdaKernel<<>>(count, offset, VoxelKeyFunctor(), mDeviceData, points, id, d_keys, d_indx); + if (mVerbose==2) mTimer.restart("DeviceRadixSort of "+std::to_string(pointCount)+" tile keys"); + CALL_CUBS(DeviceRadixSort::SortPairs, d_keys, mData.d_keys, d_indx, mData.d_indx, pointCount, 0, 63);// 21 bits per coord + std::swap(d_indx, mData.d_indx);// sorted indices are now in d_indx + + if (mVerbose==2) mTimer.restart("Allocate runs"); + d_points_per_tile = static_cast(mResource->allocate_async(pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + d_node_count = static_cast(mResource->allocate_async(3*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + + if (mVerbose==2) mTimer.restart("DeviceRunLengthEncode tile keys"); + CALL_CUBS(DeviceRunLengthEncode::Encode, mData.d_keys, d_keys, d_points_per_tile, d_node_count+2, pointCount); + cudaCheck(cudaMemcpyAsync(mData.nodeCount+2, d_node_count+2, sizeof(uint32_t), cudaMemcpyDeviceToHost, mStream)); + cudaCheck(cudaStreamSynchronize(mStream)); + mData.d_tile_keys = static_cast(mResource->allocate_async(mData.nodeCount[2]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + cudaCheck(cudaMemcpyAsync(mData.d_tile_keys, d_keys, mData.nodeCount[2]*sizeof(uint64_t), cudaMemcpyDeviceToDevice, mStream)); + + static constexpr uint32_t SEGMENTED_SORT_TILE_THRESHOLD = 32; + if (mData.nodeCount[2] >= SEGMENTED_SORT_TILE_THRESHOLD) { + // Bulk segmented sort: one kernel launch + one segmented radix sort (faster for many tiles) + if (mVerbose==2) mTimer.restart("Segmented radix sort of " + std::to_string(pointCount) + " voxel keys in " + std::to_string(mData.nodeCount[2]) + " tiles"); + auto *d_tile_offsets = static_cast(mResource->allocate_async((mData.nodeCount[2]+1)*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + cudaCheck(cudaMemsetAsync(d_tile_offsets, 0, sizeof(uint32_t), mStream)); + CALL_CUBS(DeviceScan::InclusiveSum, d_points_per_tile, d_tile_offsets + 1, mData.nodeCount[2]); + mResource->deallocate_async(d_points_per_tile, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + + util::cuda::lambdaKernel<<>>(pointCount, BulkVoxelKeyFunctor(), mDeviceData, points, d_tile_offsets, mData.nodeCount[2], d_keys, d_indx, uint32_t(0)); cudaCheckError(); - CALL_CUBS(DeviceRadixSort::SortPairs, d_keys + offset, mData.d_keys + offset, d_indx + offset, mData.d_indx + offset, count, 0, 36); - offset += count; + CALL_CUBS(DeviceSegmentedRadixSort::SortPairs, d_keys, mData.d_keys, d_indx, mData.d_indx, (int)pointCount, (int)mData.nodeCount[2], d_tile_offsets, d_tile_offsets + 1, 0, 36); + mResource->deallocate_async(d_tile_offsets, (mData.nodeCount[2]+1)*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + } else { + // Serial per-tile sort: individual kernel + sort per tile (lower overhead for few tiles) + if (mVerbose==2) mTimer.restart("DeviceRadixSort of " + std::to_string(pointCount) + " voxel keys in " + std::to_string(mData.nodeCount[2]) + " tiles"); + uint32_t *points_per_tile = new uint32_t[mData.nodeCount[2]]; + cudaCheck(cudaMemcpyAsync(points_per_tile, d_points_per_tile, mData.nodeCount[2]*sizeof(uint32_t), cudaMemcpyDeviceToHost, mStream)); + mResource->deallocate_async(d_points_per_tile, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + for (uint32_t id = 0, offset = 0; id < mData.nodeCount[2]; ++id) { + const uint32_t count = points_per_tile[id]; + util::cuda::offsetLambdaKernel<<>>(count, offset, VoxelKeyFunctor(), mDeviceData, points, id, d_keys, d_indx); + cudaCheckError(); + CALL_CUBS(DeviceRadixSort::SortPairs, d_keys + offset, mData.d_keys + offset, d_indx + offset, mData.d_indx + offset, count, 0, 36); + offset += count; + } + delete [] points_per_tile; } - delete [] points_per_tile; - } - mResource->deallocate_async(d_indx, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + mResource->deallocate_async(d_indx, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - if (mVerbose==2) mTimer.restart("Count points per voxel"); + if (mVerbose==2) mTimer.restart("Count points per voxel"); - cudaEvent_t copyEvent; - cudaCheck(cudaEventCreate(©Event)); - mData.pointsPerVoxel = static_cast(mResource->allocate_async(pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); - uint32_t *d_voxel_count = static_cast(mResource->allocate_async(sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); - CALL_CUBS(DeviceRunLengthEncode::Encode, mData.d_keys, d_keys, mData.pointsPerVoxel, d_voxel_count, pointCount); - cudaCheck(cudaMemcpyAsync(&mData.voxelCount, d_voxel_count, sizeof(uint32_t), cudaMemcpyDeviceToHost, mStream)); - cudaCheck(cudaEventRecord(copyEvent, mStream)); - mResource->deallocate_async(d_voxel_count, sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - - if (util::is_same::value) { - if (mVerbose==2) mTimer.restart("Count max points per voxel"); - uint32_t *d_maxPointsPerVoxel = static_cast(mResource->allocate_async(sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)), maxPointsPerVoxel; - cudaCheck(cudaEventSynchronize(copyEvent)); - CALL_CUBS(DeviceReduce::Max, mData.pointsPerVoxel, d_maxPointsPerVoxel, mData.voxelCount); - cudaCheck(cudaMemcpyAsync(&maxPointsPerVoxel, d_maxPointsPerVoxel, sizeof(uint32_t), cudaMemcpyDeviceToHost, mStream)); + mData.pointsPerVoxel = static_cast(mResource->allocate_async(pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + uint32_t *d_voxel_count = static_cast(mResource->allocate_async(sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + CALL_CUBS(DeviceRunLengthEncode::Encode, mData.d_keys, d_keys, mData.pointsPerVoxel, d_voxel_count, pointCount); + cudaCheck(cudaMemcpyAsync(&mData.voxelCount, d_voxel_count, sizeof(uint32_t), cudaMemcpyDeviceToHost, mStream)); cudaCheck(cudaEventRecord(copyEvent, mStream)); - mResource->deallocate_async(d_maxPointsPerVoxel, sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - double dx = mData.map.getVoxelSize()[0]; - cudaCheck(cudaEventSynchronize(copyEvent)); - if (++iterCounter >= mMaxIterations || pointCount == 1u || math::Abs((int)maxPointsPerVoxel - (int)mMaxPointsPerVoxel) <= mTolerance) { - mMaxPointsPerVoxel = maxPointsPerVoxel; - } else { - const Foo tmp{dx, maxPointsPerVoxel}; - if (maxPointsPerVoxel < mMaxPointsPerVoxel) { - if (min < tmp) min = tmp; - } else if (max.density == 0 || tmp < max) { - max = tmp; - } - if (max.density) { - dx = (min.dx*(max.density - mMaxPointsPerVoxel) + max.dx*(mMaxPointsPerVoxel-min.density))/double(max.density-min.density); - } else if (maxPointsPerVoxel > 1u) { - dx *= (mMaxPointsPerVoxel-1.0)/(maxPointsPerVoxel-1.0); - } else {// maxPointsPerVoxel = 1 so increase dx significantly - dx *= 10.0; + mResource->deallocate_async(d_voxel_count, sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + + if (util::is_same::value) { + if (mVerbose==2) mTimer.restart("Count max points per voxel"); + uint32_t *d_maxPointsPerVoxel = static_cast(mResource->allocate_async(sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)), maxPointsPerVoxel; + cudaCheck(cudaEventSynchronize(copyEvent)); + CALL_CUBS(DeviceReduce::Max, mData.pointsPerVoxel, d_maxPointsPerVoxel, mData.voxelCount); + cudaCheck(cudaMemcpyAsync(&maxPointsPerVoxel, d_maxPointsPerVoxel, sizeof(uint32_t), cudaMemcpyDeviceToHost, mStream)); + cudaCheck(cudaEventRecord(copyEvent, mStream)); + mResource->deallocate_async(d_maxPointsPerVoxel, sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + double dx = mData.map.getVoxelSize()[0]; + cudaCheck(cudaEventSynchronize(copyEvent)); + if (++iterCounter >= mMaxIterations || pointCount == 1u || math::Abs((int)maxPointsPerVoxel - (int)mMaxPointsPerVoxel) <= mTolerance) { + mMaxPointsPerVoxel = maxPointsPerVoxel; + } else { + const Foo tmp{dx, maxPointsPerVoxel}; + if (maxPointsPerVoxel < mMaxPointsPerVoxel) { + if (min < tmp) min = tmp; + } else if (max.density == 0 || tmp < max) { + max = tmp; + } + if (max.density) { + dx = (min.dx*(max.density - mMaxPointsPerVoxel) + max.dx*(mMaxPointsPerVoxel-min.density))/double(max.density-min.density); + } else if (maxPointsPerVoxel > 1u) { + dx *= (mMaxPointsPerVoxel-1.0)/(maxPointsPerVoxel-1.0); + } else {// maxPointsPerVoxel = 1 so increase dx significantly + dx *= 10.0; + } + if (mVerbose==2) printf("\ntarget density = %" PRIu32 ", current density = %" PRIu32 ", current dx = %f, next dx = %f\n", mMaxPointsPerVoxel, maxPointsPerVoxel, tmp.dx, dx); + mData.map = Map(dx); + mResource->deallocate_async(mData.d_keys, pointCount*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + mResource->deallocate_async(mData.d_indx, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + mResource->deallocate_async(d_keys, pointCount*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + mResource->deallocate_async(mData.d_tile_keys, mData.nodeCount[2]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + mResource->deallocate_async(d_node_count, 3*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + mResource->deallocate_async(mData.pointsPerVoxel, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + continue; } - if (mVerbose==2) printf("\ntarget density = %" PRIu32 ", current density = %" PRIu32 ", current dx = %f, next dx = %f\n", mMaxPointsPerVoxel, maxPointsPerVoxel, tmp.dx, dx); - mData.map = Map(dx); - mResource->deallocate_async(mData.d_keys, pointCount*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - mResource->deallocate_async(mData.d_indx, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - mResource->deallocate_async(d_keys, pointCount*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - mResource->deallocate_async(mData.d_tile_keys, mData.nodeCount[2]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - mResource->deallocate_async(d_node_count, 3*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - mResource->deallocate_async(mData.pointsPerVoxel, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - goto jump; } - } + break; + }// while (true) if (iterCounter>1 && mVerbose) std::cerr << "Used " << iterCounter << " attempts to determine dx that produces a target dpoint denisty\n\n"; if (mVerbose==2) mTimer.restart("Compute prefix sum of points per voxel"); From a12201c9856e20ed81837b28f07aa765a43b1667 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 5 Aug 2026 04:02:36 +0000 Subject: [PATCH 2/4] NanoVDB: own PointsToGrid's device arrays with cuda::Buffer Every device array PointsToGrid allocates is now owned by a Buffer> borrowing the injected resource -- members where the pipeline frees them in a later member function than the one that allocated them (countNodes allocates; processUpperNodes, processLeafNodes, processPoints and processBBox release), locals where the lifetime is contained. The raw pointers survive only as views: the device-visible fields inside mData, and the working pointers the cub and kernel calls take. Every owner event -- assignment, swap, destroy -- immediately refreshes its view, and released views are nulled so a stale use faults instead of reading a freed block. The index ping-pong becomes a swap of owners across the member/local boundary, replacing the bare pointer swap whose safety depended on nothing reading the device copy of d_indx between the two uploads. The density-search retry keeps its free-before-reallocate order via explicit destroy calls, so peak device memory is unchanged. The hand-matched byte sizes at every free site disappear, and the arrays released one line before scope exit now just leave scope. One behavior change worth naming: the too-many-points-per-leaf throw previously leaked the reduction scratch; ownership now releases it during unwind. Verified against the previous commit with a counting resource over three shapes (bulk segmented-sort branch, serial per-tile branch, and the bisection retry engaged): allocation count, free count, and total bytes are identical. Full CUDA and memory-resource suites unchanged. Co-Authored-By: Claude Opus 5 (1M context) Signed-off-by: Mark Harris --- nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh | 141 +++++++++++++------- pendingchanges/nanovdb.txt | 2 +- 2 files changed, 93 insertions(+), 50 deletions(-) diff --git a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh index 3fb555851a..5fd5640d2f 100644 --- a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh @@ -303,10 +303,20 @@ public: , mResource(&resource) , mTimer(stream) , mPointType(util::is_same::value ? PointType::Default : PointType::Disable) + , mDeviceDataBuf(stream, nanovdb::cuda::ResourceRef(resource), 1, nanovdb::cuda::noInit) + , mKeysBuf(stream, nanovdb::cuda::ResourceRef(resource), 0, nanovdb::cuda::noInit) + , mTileKeysBuf(stream, nanovdb::cuda::ResourceRef(resource), 0, nanovdb::cuda::noInit) + , mLeafKeysBuf(stream, nanovdb::cuda::ResourceRef(resource), 0, nanovdb::cuda::noInit) + , mLowerKeysBuf(stream, nanovdb::cuda::ResourceRef(resource), 0, nanovdb::cuda::noInit) + , mIndxBuf(stream, nanovdb::cuda::ResourceRef(resource), 0, nanovdb::cuda::noInit) + , mPointsPerVoxelBuf(stream, nanovdb::cuda::ResourceRef(resource), 0, nanovdb::cuda::noInit) + , mPointsPerVoxelPrefixBuf(stream, nanovdb::cuda::ResourceRef(resource), 0, nanovdb::cuda::noInit) + , mPointsPerLeafBuf(stream, nanovdb::cuda::ResourceRef(resource), 0, nanovdb::cuda::noInit) + , mPointsPerLeafPrefixBuf(stream, nanovdb::cuda::ResourceRef(resource), 0, nanovdb::cuda::noInit) , mTempDevicePool(resource) { mData.map = map; - mDeviceData = static_cast*>(mResource->allocate_async(sizeof(PointsToGridData), ResourceT::DEFAULT_ALIGNMENT, mStream)); + mDeviceData = mDeviceDataBuf.data(); } /// @brief Default constructor that calls the Map constructor defined above @@ -327,7 +337,6 @@ public: mMaxIterations = maxIterations; } - ~PointsToGrid(){ mResource->deallocate_async(mDeviceData, sizeof(PointsToGridData), ResourceT::DEFAULT_ALIGNMENT, mStream); } /// @brief Toggle on and off verbose mode /// @param level Verbose level: 0=quiet, 1=timing, 2=benchmarking @@ -392,13 +401,26 @@ private: static constexpr unsigned int mNumThreads = 128;// seems faster than the old value of 256! static unsigned int numBlocks(unsigned int n) {return (n + mNumThreads - 1) / mNumThreads;} + template + using BufT = nanovdb::cuda::Buffer>; + nanovdb::cuda::ResourceRef ref() { return nanovdb::cuda::ResourceRef(*mResource); } + cudaStream_t mStream{0}; ResourceT* mResource;// non-owning; all device allocations (mDeviceData + scratch) route through this resource instance util::cuda::Timer mTimer; PointType mPointType; std::string mGridName; int mVerbose{0}; - PointsToGridData mData, *mDeviceData; + PointsToGridData mData, *mDeviceData;// mDeviceData views mDeviceDataBuf + // Owners of the device arrays that mData's raw pointers view. Raw views are + // refreshed immediately after every owner event (assign, swap, destroy), so + // the views -- including the device-visible fields inside mData -- are + // never stale. Members rather than locals because the pipeline frees them + // across member functions (countNodes allocates; processUpperNodes, + // processLeafNodes, processPoints and processBBox release). + BufT> mDeviceDataBuf; + BufT mKeysBuf, mTileKeysBuf, mLeafKeysBuf, mLowerKeysBuf; + BufT mIndxBuf, mPointsPerVoxelBuf, mPointsPerVoxelPrefixBuf, mPointsPerLeafBuf, mPointsPerLeafPrefixBuf; uint32_t mMaxPointsPerVoxel{0u}, mMaxPointsPerLeaf{0u}; int mTolerance{1}, mMaxIterations{1}; CheckMode mChecksum{CheckMode::Disable}; @@ -613,6 +635,9 @@ void PointsToGrid::countNodes(const PtrT points, size_t point // Declared ahead of the search loop below: d_keys and d_node_count carry // results past it, and the event is recorded and re-used across iterations // (previously it was re-created per iteration, leaking the prior handle). + BufT keysScratch(mStream, this->ref(), 0, nanovdb::cuda::noInit); + BufT indxScratch(mStream, this->ref(), 0, nanovdb::cuda::noInit); + BufT nodeCountScratch(mStream, this->ref(), 0, nanovdb::cuda::noInit); uint64_t *d_keys = nullptr; uint32_t *d_indx = nullptr, *d_points_per_tile = nullptr, *d_node_count = nullptr; cudaEvent_t copyEvent; @@ -623,51 +648,60 @@ void PointsToGrid::countNodes(const PtrT points, size_t point // then either converges or frees this iteration's buffers and retries. while (true) { - mData.d_keys = static_cast(mResource->allocate_async(pointCount*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); - mData.d_indx = static_cast(mResource->allocate_async(pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream));// uint32_t can index 4.29 billion Coords, corresponding to 48 GB + mKeysBuf = BufT(mStream, this->ref(), pointCount, nanovdb::cuda::noInit); + mIndxBuf = BufT(mStream, this->ref(), pointCount, nanovdb::cuda::noInit);// uint32_t can index 4.29 billion Coords, corresponding to 48 GB + mData.d_keys = mKeysBuf.data(); + mData.d_indx = mIndxBuf.data(); cudaCheck(cudaMemcpyAsync(mDeviceData, &mData, sizeof(PointsToGridData), cudaMemcpyHostToDevice, mStream));// copy mData from CPU -> GPU if (mVerbose==2) mTimer.start("\nAllocating arrays for keys and indices"); - d_keys = static_cast(mResource->allocate_async(pointCount*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); - d_indx = static_cast(mResource->allocate_async(pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + keysScratch = BufT(mStream, this->ref(), pointCount, nanovdb::cuda::noInit); + indxScratch = BufT(mStream, this->ref(), pointCount, nanovdb::cuda::noInit); + d_keys = keysScratch.data(); + d_indx = indxScratch.data(); if (mVerbose==2) mTimer.restart("Generate tile keys"); util::cuda::lambdaKernel<<>>(pointCount, TileKeyFunctor(), mDeviceData, points, d_keys, d_indx); cudaCheckError(); if (mVerbose==2) mTimer.restart("DeviceRadixSort of "+std::to_string(pointCount)+" tile keys"); CALL_CUBS(DeviceRadixSort::SortPairs, d_keys, mData.d_keys, d_indx, mData.d_indx, pointCount, 0, 63);// 21 bits per coord - std::swap(d_indx, mData.d_indx);// sorted indices are now in d_indx + mIndxBuf.swap(indxScratch);// the sorted indices' owner is now indxScratch + d_indx = indxScratch.data();// sorted indices + mData.d_indx = mIndxBuf.data();// receives the voxel-sorted indices below if (mVerbose==2) mTimer.restart("Allocate runs"); - d_points_per_tile = static_cast(mResource->allocate_async(pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); - d_node_count = static_cast(mResource->allocate_async(3*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + BufT pointsPerTileScratch(mStream, this->ref(), pointCount, nanovdb::cuda::noInit); + d_points_per_tile = pointsPerTileScratch.data(); + nodeCountScratch = BufT(mStream, this->ref(), 3, nanovdb::cuda::noInit); + d_node_count = nodeCountScratch.data(); if (mVerbose==2) mTimer.restart("DeviceRunLengthEncode tile keys"); CALL_CUBS(DeviceRunLengthEncode::Encode, mData.d_keys, d_keys, d_points_per_tile, d_node_count+2, pointCount); cudaCheck(cudaMemcpyAsync(mData.nodeCount+2, d_node_count+2, sizeof(uint32_t), cudaMemcpyDeviceToHost, mStream)); cudaCheck(cudaStreamSynchronize(mStream)); - mData.d_tile_keys = static_cast(mResource->allocate_async(mData.nodeCount[2]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + mTileKeysBuf = BufT(mStream, this->ref(), mData.nodeCount[2], nanovdb::cuda::noInit); + mData.d_tile_keys = mTileKeysBuf.data(); cudaCheck(cudaMemcpyAsync(mData.d_tile_keys, d_keys, mData.nodeCount[2]*sizeof(uint64_t), cudaMemcpyDeviceToDevice, mStream)); static constexpr uint32_t SEGMENTED_SORT_TILE_THRESHOLD = 32; if (mData.nodeCount[2] >= SEGMENTED_SORT_TILE_THRESHOLD) { // Bulk segmented sort: one kernel launch + one segmented radix sort (faster for many tiles) if (mVerbose==2) mTimer.restart("Segmented radix sort of " + std::to_string(pointCount) + " voxel keys in " + std::to_string(mData.nodeCount[2]) + " tiles"); - auto *d_tile_offsets = static_cast(mResource->allocate_async((mData.nodeCount[2]+1)*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + BufT tileOffsetsScratch(mStream, this->ref(), mData.nodeCount[2]+1, nanovdb::cuda::noInit); + auto *d_tile_offsets = tileOffsetsScratch.data(); cudaCheck(cudaMemsetAsync(d_tile_offsets, 0, sizeof(uint32_t), mStream)); CALL_CUBS(DeviceScan::InclusiveSum, d_points_per_tile, d_tile_offsets + 1, mData.nodeCount[2]); - mResource->deallocate_async(d_points_per_tile, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + pointsPerTileScratch.destroy(); d_points_per_tile = nullptr; util::cuda::lambdaKernel<<>>(pointCount, BulkVoxelKeyFunctor(), mDeviceData, points, d_tile_offsets, mData.nodeCount[2], d_keys, d_indx, uint32_t(0)); cudaCheckError(); CALL_CUBS(DeviceSegmentedRadixSort::SortPairs, d_keys, mData.d_keys, d_indx, mData.d_indx, (int)pointCount, (int)mData.nodeCount[2], d_tile_offsets, d_tile_offsets + 1, 0, 36); - mResource->deallocate_async(d_tile_offsets, (mData.nodeCount[2]+1)*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); } else { // Serial per-tile sort: individual kernel + sort per tile (lower overhead for few tiles) if (mVerbose==2) mTimer.restart("DeviceRadixSort of " + std::to_string(pointCount) + " voxel keys in " + std::to_string(mData.nodeCount[2]) + " tiles"); uint32_t *points_per_tile = new uint32_t[mData.nodeCount[2]]; cudaCheck(cudaMemcpyAsync(points_per_tile, d_points_per_tile, mData.nodeCount[2]*sizeof(uint32_t), cudaMemcpyDeviceToHost, mStream)); - mResource->deallocate_async(d_points_per_tile, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + pointsPerTileScratch.destroy(); d_points_per_tile = nullptr; for (uint32_t id = 0, offset = 0; id < mData.nodeCount[2]; ++id) { const uint32_t count = points_per_tile[id]; util::cuda::offsetLambdaKernel<<>>(count, offset, VoxelKeyFunctor(), mDeviceData, points, id, d_keys, d_indx); @@ -677,25 +711,28 @@ void PointsToGrid::countNodes(const PtrT points, size_t point } delete [] points_per_tile; } - mResource->deallocate_async(d_indx, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + indxScratch.destroy(); d_indx = nullptr;// tile-order copy, superseded by the voxel sort if (mVerbose==2) mTimer.restart("Count points per voxel"); - mData.pointsPerVoxel = static_cast(mResource->allocate_async(pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); - uint32_t *d_voxel_count = static_cast(mResource->allocate_async(sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + mPointsPerVoxelBuf = BufT(mStream, this->ref(), pointCount, nanovdb::cuda::noInit); + mData.pointsPerVoxel = mPointsPerVoxelBuf.data(); + BufT voxelCountScratch(mStream, this->ref(), 1, nanovdb::cuda::noInit); + uint32_t *d_voxel_count = voxelCountScratch.data(); CALL_CUBS(DeviceRunLengthEncode::Encode, mData.d_keys, d_keys, mData.pointsPerVoxel, d_voxel_count, pointCount); cudaCheck(cudaMemcpyAsync(&mData.voxelCount, d_voxel_count, sizeof(uint32_t), cudaMemcpyDeviceToHost, mStream)); cudaCheck(cudaEventRecord(copyEvent, mStream)); - mResource->deallocate_async(d_voxel_count, sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + voxelCountScratch.destroy(); d_voxel_count = nullptr; if (util::is_same::value) { if (mVerbose==2) mTimer.restart("Count max points per voxel"); - uint32_t *d_maxPointsPerVoxel = static_cast(mResource->allocate_async(sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)), maxPointsPerVoxel; + BufT maxPointsPerVoxelScratch(mStream, this->ref(), 1, nanovdb::cuda::noInit); + uint32_t *d_maxPointsPerVoxel = maxPointsPerVoxelScratch.data(), maxPointsPerVoxel; cudaCheck(cudaEventSynchronize(copyEvent)); CALL_CUBS(DeviceReduce::Max, mData.pointsPerVoxel, d_maxPointsPerVoxel, mData.voxelCount); cudaCheck(cudaMemcpyAsync(&maxPointsPerVoxel, d_maxPointsPerVoxel, sizeof(uint32_t), cudaMemcpyDeviceToHost, mStream)); cudaCheck(cudaEventRecord(copyEvent, mStream)); - mResource->deallocate_async(d_maxPointsPerVoxel, sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + maxPointsPerVoxelScratch.destroy(); d_maxPointsPerVoxel = nullptr; double dx = mData.map.getVoxelSize()[0]; cudaCheck(cudaEventSynchronize(copyEvent)); if (++iterCounter >= mMaxIterations || pointCount == 1u || math::Abs((int)maxPointsPerVoxel - (int)mMaxPointsPerVoxel) <= mTolerance) { @@ -716,12 +753,14 @@ void PointsToGrid::countNodes(const PtrT points, size_t point } if (mVerbose==2) printf("\ntarget density = %" PRIu32 ", current density = %" PRIu32 ", current dx = %f, next dx = %f\n", mMaxPointsPerVoxel, maxPointsPerVoxel, tmp.dx, dx); mData.map = Map(dx); - mResource->deallocate_async(mData.d_keys, pointCount*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - mResource->deallocate_async(mData.d_indx, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - mResource->deallocate_async(d_keys, pointCount*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - mResource->deallocate_async(mData.d_tile_keys, mData.nodeCount[2]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - mResource->deallocate_async(d_node_count, 3*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - mResource->deallocate_async(mData.pointsPerVoxel, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + // free before the next iteration reallocates, so peak device + // memory matches the pre-loop behavior + mKeysBuf.destroy(); mData.d_keys = nullptr; + mIndxBuf.destroy(); mData.d_indx = nullptr; + keysScratch.destroy(); d_keys = nullptr; + mTileKeysBuf.destroy(); mData.d_tile_keys = nullptr; + nodeCountScratch.destroy(); d_node_count = nullptr; + mPointsPerVoxelBuf.destroy();mData.pointsPerVoxel = nullptr; continue; } } @@ -731,16 +770,19 @@ void PointsToGrid::countNodes(const PtrT points, size_t point if (mVerbose==2) mTimer.restart("Compute prefix sum of points per voxel"); cudaCheck(cudaEventSynchronize(copyEvent)); - mData.pointsPerVoxelPrefix = static_cast(mResource->allocate_async(mData.voxelCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + mPointsPerVoxelPrefixBuf = BufT(mStream, this->ref(), mData.voxelCount, nanovdb::cuda::noInit); + mData.pointsPerVoxelPrefix = mPointsPerVoxelPrefixBuf.data(); CALL_CUBS(DeviceScan::ExclusiveSum, mData.pointsPerVoxel, mData.pointsPerVoxelPrefix, mData.voxelCount); - mData.pointsPerLeaf = static_cast(mResource->allocate_async(pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + mPointsPerLeafBuf = BufT(mStream, this->ref(), pointCount, nanovdb::cuda::noInit); + mData.pointsPerLeaf = mPointsPerLeafBuf.data(); CALL_CUBS(DeviceRunLengthEncode::Encode, thrust::make_transform_iterator(mData.d_keys, ShiftRight<9>()), d_keys, mData.pointsPerLeaf, d_node_count, pointCount); cudaCheck(cudaMemcpyAsync(mData.nodeCount, d_node_count, sizeof(uint32_t), cudaMemcpyDeviceToHost, mStream)); cudaCheck(cudaEventRecord(copyEvent, mStream)); if constexpr(util::is_same::value) { - uint32_t *d_maxPointsPerLeaf = static_cast(mResource->allocate_async(sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + BufT maxPointsPerLeafScratch(mStream, this->ref(), 1, nanovdb::cuda::noInit); + uint32_t *d_maxPointsPerLeaf = maxPointsPerLeafScratch.data(); cudaCheck(cudaEventSynchronize(copyEvent)); CALL_CUBS(DeviceReduce::Max, mData.pointsPerLeaf, d_maxPointsPerLeaf, mData.nodeCount[0]); cudaCheck(cudaMemcpyAsync(&mMaxPointsPerLeaf, d_maxPointsPerLeaf, sizeof(uint32_t), cudaMemcpyDeviceToHost, mStream)); @@ -748,25 +790,25 @@ void PointsToGrid::countNodes(const PtrT points, size_t point if (mMaxPointsPerLeaf > std::numeric_limits::max()) { throw std::runtime_error("Too many points per leaf: "+std::to_string(mMaxPointsPerLeaf)); } - mResource->deallocate_async(d_maxPointsPerLeaf, sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); } cudaCheck(cudaEventSynchronize(copyEvent)); - mData.pointsPerLeafPrefix = static_cast(mResource->allocate_async(mData.nodeCount[0]*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + mPointsPerLeafPrefixBuf = BufT(mStream, this->ref(), mData.nodeCount[0], nanovdb::cuda::noInit); + mData.pointsPerLeafPrefix = mPointsPerLeafPrefixBuf.data(); CALL_CUBS(DeviceScan::ExclusiveSum, mData.pointsPerLeaf, mData.pointsPerLeafPrefix, mData.nodeCount[0]); cudaCheck(cudaStreamSynchronize(mStream)); - mData.d_leaf_keys = static_cast(mResource->allocate_async(mData.nodeCount[0]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + mLeafKeysBuf = BufT(mStream, this->ref(), mData.nodeCount[0], nanovdb::cuda::noInit); + mData.d_leaf_keys = mLeafKeysBuf.data(); cudaCheck(cudaMemcpyAsync(mData.d_leaf_keys, d_keys, mData.nodeCount[0]*sizeof(uint64_t), cudaMemcpyDeviceToDevice, mStream)); CALL_CUBS(DeviceSelect::Unique, thrust::make_transform_iterator(mData.d_leaf_keys, ShiftRight<12>()), d_keys, d_node_count+1, mData.nodeCount[0]);// count lower nodes cudaCheck(cudaMemcpyAsync(mData.nodeCount+1, d_node_count+1, sizeof(uint32_t), cudaMemcpyDeviceToHost, mStream)); cudaCheck(cudaStreamSynchronize(mStream)); - mData.d_lower_keys = static_cast(mResource->allocate_async(mData.nodeCount[1]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + mLowerKeysBuf = BufT(mStream, this->ref(), mData.nodeCount[1], nanovdb::cuda::noInit); + mData.d_lower_keys = mLowerKeysBuf.data(); cudaCheck(cudaMemcpyAsync(mData.d_lower_keys, d_keys, mData.nodeCount[1]*sizeof(uint64_t), cudaMemcpyDeviceToDevice, mStream)); - mResource->deallocate_async(d_keys, pointCount*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - mResource->deallocate_async(d_node_count, 3*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); if (mVerbose==2) mTimer.stop(); cudaCheck(cudaEventDestroy(copyEvent)); @@ -1011,7 +1053,7 @@ inline void PointsToGrid::processUpperNodes() util::cuda::lambdaKernel<<>>(mData.nodeCount[2], BuildUpperNodesFunctor(), mDeviceData); cudaCheckError(); - mResource->deallocate_async(mData.d_tile_keys, mData.nodeCount[2]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + mTileKeysBuf.destroy(); mData.d_tile_keys = nullptr; const uint64_t valueCount = mData.nodeCount[2] << 15; util::cuda::lambdaKernel<<>>(valueCount, SetUpperBackgroundValuesFunctor(), mDeviceData); @@ -1146,11 +1188,11 @@ inline void PointsToGrid::processLeafNodes(size_t pointCount) util::cuda::lambdaKernel<<>>(mData.voxelCount, SetLeafActiveVoxelStateAndValuesFunctor(), mDeviceData); cudaCheckError(); - mResource->deallocate_async(mData.d_keys, pointCount*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - mResource->deallocate_async(mData.pointsPerVoxel, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - mResource->deallocate_async(mData.pointsPerVoxelPrefix, mData.voxelCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - mResource->deallocate_async(mData.pointsPerLeafPrefix, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - mResource->deallocate_async(mData.pointsPerLeaf,mData.nodeCount[0]*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + mKeysBuf.destroy(); mData.d_keys = nullptr; + mPointsPerVoxelBuf.destroy(); mData.pointsPerVoxel = nullptr; + mPointsPerVoxelPrefixBuf.destroy();mData.pointsPerVoxelPrefix = nullptr; + mPointsPerLeafPrefixBuf.destroy(); mData.pointsPerLeafPrefix = nullptr; + mPointsPerLeafBuf.destroy(); mData.pointsPerLeaf = nullptr; if (mVerbose==2) mTimer.restart("set inactive voxel values"); const uint64_t denseVoxelCount = mData.nodeCount[0] << 9; @@ -1159,15 +1201,16 @@ inline void PointsToGrid::processLeafNodes(size_t pointCount) if constexpr(BuildTraits::is_onindex) { if (mVerbose==2) mTimer.restart("prefix-sum for index grid"); - auto devValueIndex = static_cast(mResource->allocate_async(mData.nodeCount[0]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); - auto devValueIndexPrefix = static_cast(mResource->allocate_async(mData.nodeCount[0]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream)); + BufT valueIndexScratch(mStream, this->ref(), mData.nodeCount[0], nanovdb::cuda::noInit); + BufT valueIndexPrefixScratch(mStream, this->ref(), mData.nodeCount[0], nanovdb::cuda::noInit); + auto devValueIndex = valueIndexScratch.data(); + auto devValueIndexPrefix = valueIndexPrefixScratch.data(); kernels::fillValueIndexKernel<<>>(mData.nodeCount[0], 0, devValueIndex, mDeviceData); cudaCheckError(); CALL_CUBS(DeviceScan::InclusiveSum, devValueIndex, devValueIndexPrefix, mData.nodeCount[0]); - mResource->deallocate_async(devValueIndex, mData.nodeCount[0]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + valueIndexScratch.destroy(); devValueIndex = nullptr; kernels::leafPrefixSumKernel<<>>(mData.nodeCount[0], 0, devValueIndexPrefix, mDeviceData); cudaCheckError(); - mResource->deallocate_async(devValueIndexPrefix, mData.nodeCount[0]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); } if (mVerbose==2) mTimer.stop(); @@ -1190,7 +1233,7 @@ template inline void PointsToGrid::processPoints(const PtrT points, size_t pointCount) { if constexpr(util::is_same::value) this->encodePoints(points, pointCount); - mResource->deallocate_async(mData.d_indx, pointCount*sizeof(uint32_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + mIndxBuf.destroy(); mData.d_indx = nullptr; }// PointsToGrid::processPoints template @@ -1326,7 +1369,7 @@ inline void PointsToGrid::processBBox() // update and propagate bbox from leaf -> lower/parent nodes util::cuda::lambdaKernel<<>>(mData.nodeCount[0], UpdateAndPropagateLeafBBoxFunctor(), mDeviceData); - mResource->deallocate_async(mData.d_leaf_keys, mData.nodeCount[0]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + mLeafKeysBuf.destroy(); mData.d_leaf_keys = nullptr; cudaCheckError(); // reset bbox in upper nodes @@ -1335,7 +1378,7 @@ inline void PointsToGrid::processBBox() // propagate bbox from lower -> upper/parent node util::cuda::lambdaKernel<<>>(mData.nodeCount[1], PropagateLowerBBoxFunctor(), mDeviceData); - mResource->deallocate_async(mData.d_lower_keys, mData.nodeCount[1]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); + mLowerKeysBuf.destroy(); mData.d_lower_keys = nullptr; cudaCheckError() // propagate bbox from upper -> root/parent node diff --git a/pendingchanges/nanovdb.txt b/pendingchanges/nanovdb.txt index c676c3be27..8ddfebaf28 100644 --- a/pendingchanges/nanovdb.txt +++ b/pendingchanges/nanovdb.txt @@ -4,7 +4,7 @@ NanoVDB: - Added new _hostdev_ function named nanovdb::math::isoCrossing, which intersects a ray against a user-defined iso-surface. Improvements: - - The GPU builders now allocate all scratch through an injectable memory resource: tools::cuda::TopologyBuilder and tools::cuda::MeshToGrid gained a ResourceT template parameter (defaulted, so existing code is unaffected), joining PointsToGrid. Added nanovdb::cuda::SyncFromAsync, a CRTP base that derives the synchronous half of the resource concept from the stream-ordered half, and nanovdb::cuda::ResourceRef, a non-owning reference to a resource that is itself a resource, for containers that hold their resource by value. + - The GPU builders now allocate all scratch through an injectable memory resource: tools::cuda::TopologyBuilder and tools::cuda::MeshToGrid gained a ResourceT template parameter (defaulted, so existing code is unaffected), joining PointsToGrid. Added nanovdb::cuda::SyncFromAsync, a CRTP base that derives the synchronous half of the resource concept from the stream-ordered half, and nanovdb::cuda::ResourceRef, a non-owning reference to a resource that is itself a resource, for containers that hold their resource by value. PointsToGrid's device scratch and intermediate arrays are now owned by cuda::Buffer as well, replacing all of its hand-paired allocate/free calls. - The bug-fix to the nanovdb::ReadAccessor (see below) improves random-access performance in some use-cases (especially on the CPU). Fixes: From 49575823678d3c9aec0455b371ec86baa0a5b596 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 5 Aug 2026 05:07:28 +0000 Subject: [PATCH 3/4] NanoVDB: guard PointsToGrid's copy event with a scope owner The too-many-points-per-leaf throw unwinds past the event's manual destroy, leaking the handle. Own it with a small guard so unwinding releases it, consistent with the buffer ownership in this function. Also assert the stream-ordered resource requirement on the class, so a synchronous-only resource fails naming PointsToGrid rather than the pool inside it. Co-Authored-By: Claude Opus 5 (1M context) Signed-off-by: Mark Harris --- nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh | 15 ++++++++++++--- 1 file changed, 12 insertions(+), 3 deletions(-) diff --git a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh index 5fd5640d2f..542efcc5a4 100644 --- a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh @@ -294,6 +294,8 @@ struct PointsToGridData { template class PointsToGrid { + static_assert(nanovdb::cuda::is_async_resource::value, + "PointsToGrid allocates stream-ordered scratch and requires an AsyncResource"); public: /// @brief Map constructor, which other constructors might call /// @param map Map to be used for the output device grid @@ -640,8 +642,16 @@ void PointsToGrid::countNodes(const PtrT points, size_t point BufT nodeCountScratch(mStream, this->ref(), 0, nanovdb::cuda::noInit); uint64_t *d_keys = nullptr; uint32_t *d_indx = nullptr, *d_points_per_tile = nullptr, *d_node_count = nullptr; - cudaEvent_t copyEvent; - cudaCheck(cudaEventCreate(©Event)); + // Owns the copy event so the too-many-points-per-leaf throw below cannot + // leak the handle. + struct EventGuard { + cudaEvent_t event; + EventGuard() { cudaCheck(cudaEventCreate(&event)); } + ~EventGuard() { cudaCheck(cudaEventDestroy(event)); } + EventGuard(const EventGuard&) = delete; + EventGuard& operator=(const EventGuard&) = delete; + } eventGuard; + cudaEvent_t copyEvent = eventGuard.event; // Bisection search for the voxel size dx that yields the target point // density: each iteration builds tile and voxel keys at the current dx, @@ -810,7 +820,6 @@ void PointsToGrid::countNodes(const PtrT points, size_t point cudaCheck(cudaMemcpyAsync(mData.d_lower_keys, d_keys, mData.nodeCount[1]*sizeof(uint64_t), cudaMemcpyDeviceToDevice, mStream)); if (mVerbose==2) mTimer.stop(); - cudaCheck(cudaEventDestroy(copyEvent)); //printf("Leaf count = %u, lower count = %u, upper count = %u\n", mData.nodeCount[0], mData.nodeCount[1], mData.nodeCount[2]); }// PointsToGrid::countNodes From 2e1e237518b047f0ad32b5ee0f61fa8490555a23 Mon Sep 17 00:00:00 2001 From: Mark Harris Date: Wed, 5 Aug 2026 05:29:58 +0000 Subject: [PATCH 4/4] NanoVDB: terminate a bare cudaCheckError with a semicolon Legal without one -- the macro expands to a braced block -- but every other use in the file spells it as a statement. Co-Authored-By: Claude Opus 5 (1M context) Signed-off-by: Mark Harris --- nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh index 542efcc5a4..4f81a51a54 100644 --- a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh @@ -1388,7 +1388,7 @@ inline void PointsToGrid::processBBox() // propagate bbox from lower -> upper/parent node util::cuda::lambdaKernel<<>>(mData.nodeCount[1], PropagateLowerBBoxFunctor(), mDeviceData); mLowerKeysBuf.destroy(); mData.d_lower_keys = nullptr; - cudaCheckError() + cudaCheckError(); // propagate bbox from upper -> root/parent node util::cuda::lambdaKernel<<>>(mData.nodeCount[2], PropagateUpperBBoxFunctor(), mDeviceData);