diff --git a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh index 4ac4ed386d..4f81a51a54 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 @@ -303,10 +305,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 +339,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 +403,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}; @@ -610,126 +634,165 @@ 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 - - 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)); - - 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 - - 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.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)); + // 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; + // 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, + // then either converges or frees this iteration's buffers and retries. + while (true) { + + 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"); + 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(); - 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 + 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"); + 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)); + 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"); + 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]); + 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(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); + } 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)); + 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); + 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); + indxScratch.destroy(); d_indx = nullptr;// tile-order copy, superseded by the voxel sort - 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)); + 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_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; + voxelCountScratch.destroy(); d_voxel_count = nullptr; + + if (util::is_same::value) { + if (mVerbose==2) mTimer.restart("Count max points per voxel"); + 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)); + 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) { + 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); + // 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; } - 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"); 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)); @@ -737,27 +800,26 @@ jump:// this marks the beginning of the actual algorithm 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)); //printf("Leaf count = %u, lower count = %u, upper count = %u\n", mData.nodeCount[0], mData.nodeCount[1], mData.nodeCount[2]); }// PointsToGrid::countNodes @@ -1000,7 +1062,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); @@ -1135,11 +1197,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; @@ -1148,15 +1210,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(); @@ -1179,7 +1242,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 @@ -1315,7 +1378,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 @@ -1324,8 +1387,8 @@ 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); - cudaCheckError() + mLowerKeysBuf.destroy(); mData.d_lower_keys = nullptr; + cudaCheckError(); // propagate bbox from upper -> root/parent node util::cuda::lambdaKernel<<>>(mData.nodeCount[2], PropagateUpperBBoxFunctor(), mDeviceData); diff --git a/pendingchanges/nanovdb.txt b/pendingchanges/nanovdb.txt index bf0e59fce6..906e6fdaeb 100644 --- a/pendingchanges/nanovdb.txt +++ b/pendingchanges/nanovdb.txt @@ -5,7 +5,7 @@ NanoVDB: - 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 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: