From c3cde0fea6701931df50177c9bec3dcaec97d678 Mon Sep 17 00:00:00 2001 From: Jonathan Swartz Date: Tue, 14 Jul 2026 02:42:13 +0000 Subject: [PATCH 01/17] NanoVDB CUDA: correctness and determinism fixes Five correctness/determinism fixes to the NanoVDB CUDA tools: * GridChecksum: evalChecksum copied sizeof(GridData)+sizeof(TreeData) (736 B) from a device buffer holding a single uint32 CRC into the 4-byte head/tail of a stack Checksum -- an out-of-bounds device read plus host stack corruption (compute-sanitizer visible; a hard CUDA "invalid argument" on some configs). Copy sizeof(uint32_t) at all four sites, matching validateChecksum. * Stream propagation: the NN_FACE_EDGE_VERTEX leaf-dilation kernel and SignedFloodFill's root pass ran on the default stream instead of the caller's stream, silently corrupting output for callers using a non-blocking stream. Launch and synchronize on the operation's stream. * Bounded grid-name copy: the grid-name copy read a fixed MaxNameSize bytes from a possibly shorter std::string (whose .data() is never null, making the memset branch dead), over-reading up to 255 bytes of host heap into the grid and any written .nvdb file. Zero the field, then copy only min(name.size()+1, MaxNameSize) bytes. Applied to PointsToGrid, DistributedPointsToGrid, MeshToGrid. * Deterministic output initialization: PointsToGrid and IndexToGrid left alignment padding and stats fields carrying recycled pool bytes, making output nondeterministic and leaking host heap into files. Zero the whole grid buffer once after allocation. For PointsToGrid this makes the three dense background-fill passes (upper/lower value tables, and the inactive-leaf-values pass for all but the Point build type) redundant, so they are dropped. * DeviceBuffer stream-ordered lifetime: the destructor, move-assignment and clear() freed device allocations on the default stream regardless of the stream the buffer was last used on -- a latent use-after-free for non-blocking-stream callers, masked today by legacy-stream semantics. Track the allocation stream and free on it, matching the stream-ordered free TempPool now performs. Adds gtest coverage in unittest/TestNanoVDB.cu: cross-stream dilation determinism, grid-name round-trip (empty/short/max-length), PointsToGrid output determinism, and DeviceBuffer non-blocking-stream lifetime. No performance claims. Signed-off-by: Jonathan Swartz Co-Authored-By: Claude Opus 4.8 (1M context) --- nanovdb/nanovdb/cuda/DeviceBuffer.h | 19 ++- nanovdb/nanovdb/tools/cuda/DilateGrid.cuh | 2 +- .../tools/cuda/DistributedPointsToGrid.cuh | 10 +- nanovdb/nanovdb/tools/cuda/GridChecksum.cuh | 8 +- nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh | 4 + nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh | 10 +- nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh | 45 +++++-- .../nanovdb/tools/cuda/SignedFloodFill.cuh | 11 +- nanovdb/nanovdb/unittest/TestNanoVDB.cu | 126 ++++++++++++++++++ 9 files changed, 200 insertions(+), 35 deletions(-) diff --git a/nanovdb/nanovdb/cuda/DeviceBuffer.h b/nanovdb/nanovdb/cuda/DeviceBuffer.h index 9ad1d25441..43e7273e34 100644 --- a/nanovdb/nanovdb/cuda/DeviceBuffer.h +++ b/nanovdb/nanovdb/cuda/DeviceBuffer.h @@ -38,6 +38,7 @@ class DeviceBuffer uint64_t mSize; // total number of bytes managed by this buffer (assumed to be identical for host and device) void *mCpuData, **mGpuData; // raw pointers to the host and device buffers int mDeviceCount, mManaged;// if mManaged is non-zero this class is responsible for allocating and freeing memory buffers. Otherwise this is assumed to be handled externally + cudaStream_t mStream = 0;// stream the managed device allocations are associated with. Frees (destructor, move-assign, clear) are ordered on this stream instead of the default stream 0, otherwise a buffer last used on a non-blocking stream could be freed while that stream's work is still in flight (use-after-free). The caller must keep this stream alive at least until the buffer is destroyed. /// @brief Initialize buffer /// @param size byte size of buffer to be initialized @@ -122,9 +123,11 @@ class DeviceBuffer , mGpuData(other.mGpuData) , mDeviceCount(other.mDeviceCount) , mManaged(other.mManaged) + , mStream(other.mStream) { other.mCpuData = other.mGpuData = nullptr; other.mSize = other.mDeviceCount = other.mManaged = 0; + other.mStream = 0; } /// @brief Copy-constructor from a HostBuffer @@ -142,7 +145,10 @@ class DeviceBuffer } /// @brief Destructor frees memory on both the host and device - ~DeviceBuffer() { this->clear(); }; + /// @note Frees on the stream the buffer is associated with (mStream), not + /// the default stream, so device frees are ordered after the last + /// work issued on that stream. + ~DeviceBuffer() { this->clear(mStream); }; /// @brief Static factory method that return an instance of this buffer /// @param size byte size of buffer to be initialized @@ -310,9 +316,9 @@ class DeviceBuffer inline DeviceBuffer& DeviceBuffer::operator=(DeviceBuffer&& other) noexcept { - if (mManaged) {// first free all the managed data buffers + if (mManaged) {// first free all the managed data buffers on the stream they are associated with cudaCheck(cudaFreeHost(mCpuData)); - for (int i=0; i::dilateLeafNodes() else if (mOp == morphology::NN_FACE_EDGE_VERTEX) { using Op = util::morphology::cuda::DilateLeafNodesFunctor; util::cuda::operatorKernel - <<nodeCount[1],Op::SlicesPerLowerNode,1), Op::MaxThreadsPerBlock>>> + <<nodeCount[1],Op::SlicesPerLowerNode,1), Op::MaxThreadsPerBlock, 0, mStream>>> (mDeviceSrcGrid, static_cast(mBuilder.data()->d_bufferPtr)); } } diff --git a/nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh b/nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh index d36f2a921e..ab2336b624 100644 --- a/nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh @@ -909,11 +909,13 @@ inline void DistributedPointsToGrid::processGridTreeRoot(const PtrT poin util::cuda::lambdaKernel<<<1, 1, 0, stream>>>(1, BuildGridTreeRootFunctor(), mData, mPointType, pointCount);// lambdaKernel cudaCheckError(); + // Zero the name field, then copy only the actual string (if any); see the + // matching fix in PointsToGrid.cuh (heap over-read of the name buffer). char *dst = mData->getGrid().mGridName; - if (const char *src = mGridName.data()) { - cudaCheck(cudaMemcpyAsync(dst, src, GridData::MaxNameSize, cudaMemcpyHostToDevice, stream)); - } else { - cudaCheck(cudaMemsetAsync(dst, 0, GridData::MaxNameSize, stream)); + cudaCheck(cudaMemsetAsync(dst, 0, GridData::MaxNameSize, stream)); + if (!mGridName.empty()) { + const size_t nameSize = std::min(mGridName.size() + 1, GridData::MaxNameSize); + cudaCheck(cudaMemcpyAsync(dst, mGridName.c_str(), nameSize, cudaMemcpyHostToDevice, stream)); } cudaEventRecord(processGridTreeRootEvent); diff --git a/nanovdb/nanovdb/tools/cuda/GridChecksum.cuh b/nanovdb/nanovdb/tools/cuda/GridChecksum.cuh index eef49d5ebc..2d9bc02af5 100644 --- a/nanovdb/nanovdb/tools/cuda/GridChecksum.cuh +++ b/nanovdb/nanovdb/tools/cuda/GridChecksum.cuh @@ -234,7 +234,7 @@ inline Checksum evalChecksum(const GridData *d_gridData, CheckMode mode, cudaStr if (mode != CheckMode::Empty) { auto d_lut = util::cuda::createCrc32Lut(1, stream); crc32Head(d_gridData, d_lut.get(), d_lut.get() + 256, stream); - cudaCheck(cudaMemcpyAsync(&(cs.head()), d_lut.get() + 256, headSize, cudaMemcpyDeviceToHost, stream)); + cudaCheck(cudaMemcpyAsync(&(cs.head()), d_lut.get() + 256, sizeof(uint32_t), cudaMemcpyDeviceToHost, stream)); if (mode == CheckMode::Full) { std::unique_ptr buffer(new char[headSize]); auto *gridData = (GridData*)(buffer.get()); @@ -244,7 +244,7 @@ inline Checksum evalChecksum(const GridData *d_gridData, CheckMode mode, cudaStr } else { callNanoGrid(d_gridData, gridData, d_lut.get(), d_lut.get() + 256, stream); } - cudaCheck(cudaMemcpyAsync(&(cs.tail()), d_lut.get() + 256, headSize, cudaMemcpyDeviceToHost, stream)); + cudaCheck(cudaMemcpyAsync(&(cs.tail()), d_lut.get() + 256, sizeof(uint32_t), cudaMemcpyDeviceToHost, stream)); } } return cs; @@ -265,7 +265,7 @@ Checksum evalChecksum(const NanoGrid *d_grid, CheckMode mode, cudaStream if (mode != CheckMode::Empty) { auto d_lut = util::cuda::createCrc32Lut(1, stream); crc32Head(d_grid, d_lut.get(), d_lut.get() + 256, stream); - cudaCheck(cudaMemcpyAsync(&(cs.head()), d_lut.get() + 256, headSize, cudaMemcpyDeviceToHost, stream)); + cudaCheck(cudaMemcpyAsync(&(cs.head()), d_lut.get() + 256, sizeof(uint32_t), cudaMemcpyDeviceToHost, stream)); if (mode == CheckMode::Full) { std::unique_ptr buffer(new char[headSize]); auto *gridData = (GridData*)(buffer.get()); @@ -275,7 +275,7 @@ Checksum evalChecksum(const NanoGrid *d_grid, CheckMode mode, cudaStream } else { crc32TailOld(d_grid, gridData, d_lut.get(), d_lut.get() + 256, stream); } - cudaCheck(cudaMemcpyAsync(&(cs.tail()), d_lut.get() + 256, headSize, cudaMemcpyDeviceToHost, stream)); + cudaCheck(cudaMemcpyAsync(&(cs.tail()), d_lut.get() + 256, sizeof(uint32_t), cudaMemcpyDeviceToHost, stream)); } } return cs; diff --git a/nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh b/nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh index 81110f31d1..f5b440f9f9 100644 --- a/nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh @@ -373,6 +373,10 @@ inline BufferT IndexToGrid::getBuffer(const BufferT &pool) auto buffer = BufferT::create(mNodeAcc.size, &pool, device, mStream); mNodeAcc.d_dstPtr = buffer.deviceData(); if (mNodeAcc.d_dstPtr == nullptr) throw std::runtime_error("Failed memory allocation on the device"); + // Zero the output buffer: regions the kernels below do not write (e.g. + // padding, unused tile fields) otherwise carry recycled allocator bytes, + // making the output nondeterministic and leaking heap contents to files. + cudaCheck(cudaMemsetAsync(mNodeAcc.d_dstPtr, 0, mNodeAcc.size, mStream)); if (size_t size = mGridName.size()) { cudaCheck(util::cuda::mallocAsync((void**)&mNodeAcc.d_gridName, size, mStream)); diff --git a/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh b/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh index d143793dce..34b036c90e 100644 --- a/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh @@ -884,12 +884,14 @@ void MeshToGrid::processGridTreeRoot() topology::detail::InitGridTreeRootFunctor{mMap}, mBuilder.deviceData()); cudaCheckError(); - // Copy grid name into the output grid's name field + // Copy grid name into the output grid's name field. Zero the field first + // and copy only the actual string; see the matching fix in PointsToGrid.cuh + // (heap over-read of the name buffer). char *dst = mBuilder.data()->getGrid().mGridName; + cudaCheck(cudaMemsetAsync(dst, 0, GridData::MaxNameSize, mStream)); if (!mGridName.empty()) { - cudaCheck(cudaMemcpyAsync(dst, mGridName.data(), GridData::MaxNameSize, cudaMemcpyHostToDevice, mStream)); - } else { - cudaCheck(cudaMemsetAsync(dst, 0, GridData::MaxNameSize, mStream)); + const size_t nameSize = std::min(mGridName.size() + 1, GridData::MaxNameSize); + cudaCheck(cudaMemcpyAsync(dst, mGridName.c_str(), nameSize, cudaMemcpyHostToDevice, mStream)); } } // MeshToGrid::processGridTreeRoot diff --git a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh index 65b05b116b..7ce9755c2b 100644 --- a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh @@ -792,6 +792,12 @@ inline BufferT PointsToGrid::getBuffer(const PtrT, size_t poi mData.d_bufferPtr = buffer.deviceData(); if (mData.d_bufferPtr == nullptr) throw std::runtime_error("Failed to allocate grid buffer on the device"); + // Zero the whole grid buffer up front. This (a) makes the dense background + // fills (upper/lower value tables, inactive leaf values - all zero in this + // builder) redundant, so those kernels are skipped, and (b) makes the + // output bit-deterministic: alignment padding and stats fields no longer + // carry recycled pool bytes (which previously leaked into written files). + cudaCheck(cudaMemsetAsync(mData.d_bufferPtr, 0, mData.size, mStream)); cudaCheck(cudaMemcpyAsync(mDeviceData, &mData, sizeof(PointsToGridData), cudaMemcpyHostToDevice, mStream));// copy Data CPU -> GPU return buffer; }// PointsToGrid::getBuffer @@ -938,11 +944,16 @@ inline void PointsToGrid::processGridTreeRoot(const PtrT poin util::cuda::lambdaKernel<<<1, 1, 0, mStream>>>(1, BuildGridTreeRootFunctor(), mDeviceData, mPointType, pointCount);// lambdaKernel cudaCheckError(); + // Zero the name field, then copy only the actual string (if any). The + // previous code copied MaxNameSize bytes from the std::string buffer + // (whose .data() is never null, so the memset branch was dead), reading + // up to 255 bytes past the allocation and leaking host heap contents + // into the grid - nondeterministic output and a hygiene issue for files. char *dst = mData.getGrid().mGridName; - if (const char *src = mGridName.data()) { - cudaCheck(cudaMemcpyAsync(dst, src, GridData::MaxNameSize, cudaMemcpyHostToDevice, mStream)); - } else { - cudaCheck(cudaMemsetAsync(dst, 0, GridData::MaxNameSize, mStream)); + cudaCheck(cudaMemsetAsync(dst, 0, GridData::MaxNameSize, mStream)); + if (!mGridName.empty()) { + const size_t nameSize = std::min(mGridName.size() + 1, GridData::MaxNameSize); + cudaCheck(cudaMemcpyAsync(dst, mGridName.c_str(), nameSize, cudaMemcpyHostToDevice, mStream)); } }// PointsToGrid::processGridTreeRoot @@ -995,9 +1006,9 @@ inline void PointsToGrid::processUpperNodes() mResource->deallocate_async(mData.d_tile_keys, mData.nodeCount[2]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - const uint64_t valueCount = mData.nodeCount[2] << 15; - util::cuda::lambdaKernel<<>>(valueCount, SetUpperBackgroundValuesFunctor(), mDeviceData); - cudaCheckError(); + // Upper background values are zero and the grid buffer is zero-initialized + // in getBuffer, so the former dense SetUpperBackgroundValuesFunctor pass + // (nodeCount[2]<<15 threads) is unnecessary. }// PointsToGrid::processUpperNodes //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- @@ -1039,9 +1050,9 @@ inline void PointsToGrid::processLowerNodes() util::cuda::lambdaKernel<<>>(mData.nodeCount[1], BuildLowerNodesFunctor(), mDeviceData); cudaCheckError(); - const uint64_t valueCount = mData.nodeCount[1] << 12; - util::cuda::lambdaKernel<<>>(valueCount, SetLowerBackgroundValuesFunctor(), mDeviceData); - cudaCheckError(); + // Lower background values are zero and the grid buffer is zero-initialized + // in getBuffer, so the former dense SetLowerBackgroundValuesFunctor pass + // (nodeCount[1]<<12 threads) is unnecessary. }// PointsToGrid::processLowerNodes //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- @@ -1134,10 +1145,16 @@ inline void PointsToGrid::processLeafNodes(size_t pointCount) 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); - if (mVerbose==2) mTimer.restart("set inactive voxel values"); - const uint64_t denseVoxelCount = mData.nodeCount[0] << 9; - util::cuda::lambdaKernel<<>>(denseVoxelCount, SetLeafInactiveVoxelValuesFunctor(), mDeviceData); - cudaCheckError(); + // Inactive voxel values are zero for every build type except Point (which + // copies the previous active value for rank queries); the zero cases are + // covered by the buffer memset in getBuffer, and for index grids this + // dense pass (nodeCount[0]<<9 threads) was a no-op to begin with. + if constexpr(util::is_same::value) { + if (mVerbose==2) mTimer.restart("set inactive voxel values"); + const uint64_t denseVoxelCount = mData.nodeCount[0] << 9; + util::cuda::lambdaKernel<<>>(denseVoxelCount, SetLeafInactiveVoxelValuesFunctor(), mDeviceData); + cudaCheckError(); + } if constexpr(BuildTraits::is_onindex) { if (mVerbose==2) mTimer.restart("prefix-sum for index grid"); diff --git a/nanovdb/nanovdb/tools/cuda/SignedFloodFill.cuh b/nanovdb/nanovdb/tools/cuda/SignedFloodFill.cuh index f356e0a5d0..f3c5d9a189 100644 --- a/nanovdb/nanovdb/tools/cuda/SignedFloodFill.cuh +++ b/nanovdb/nanovdb/tools/cuda/SignedFloodFill.cuh @@ -78,7 +78,7 @@ struct RootChild { // CPU kernel! template -void processRoot(NanoTree *d_tree) +void processRoot(NanoTree *d_tree, cudaStream_t stream = 0) {// the root needs special care since unlike other nodes it's sparse and not dense! using TreeT = NanoTree; using RootT = NanoRoot; @@ -87,6 +87,9 @@ void processRoot(NanoTree *d_tree) using ChildT = RootChild; static const int dim = int(RootT::ChildNodeType::DIM); + // Ensure the node passes issued on 'stream' have completed before the synchronous copies below read d_tree + cudaCheck(cudaStreamSynchronize(stream)); + // First copy the tree and root and then its tiles, which is of unknown size nanovdb::cuda::UnifiedBuffer uBuffer(sizeof(TreeT) + sizeof(RootT), sizeof(TreeT) + sizeof(RootT) + 64*sizeof(TileT)); cudaCheck(cudaMemcpy(uBuffer.data(), d_tree, uBuffer.size(), cudaMemcpyDeviceToHost));// copy Tree and Root (minus tiles) @@ -108,11 +111,11 @@ void processRoot(NanoTree *d_tree) for (ChildT *a = first, *b = a+1; b!=last; ++a, ++b) {// loop over pairs of adjacent child nodes const Coord d = b->ijk - a->ijk;// coord delta of adjacent child nodes if (d[0]!=0 || d[1]!=0 || d[2]==dim) continue;// not same z-scanline or they are neighbors - util::cuda::lambdaKernel<<<1, 1>>>(1, [=] __device__(size_t) { + util::cuda::lambdaKernel<<<1, 1, 0, stream>>>(1, [=] __device__(size_t) { a->val[1] = root->getChild(root->tile(a->idx))->getLastValue(); b->val[0] = root->getChild(root->tile(b->idx))->getFirstValue(); }); - cudaCheck(cudaDeviceSynchronize());// required for host access to RootChild::val[2] + cudaCheck(cudaStreamSynchronize(stream));// required for host access to RootChild::val[2] if (a->val[1] > 0 || b->val[0] > 0) continue; // scanline is not inside a surface for (Coord c = a->ijk.offsetBy(0,0,dim); c[2] != b->ijk[2]; c[2] += dim) { TileT *tile = root->probeTile(c); @@ -208,7 +211,7 @@ void SignedFloodFill::operator()(NanoGrid *d_grid) cudaCheckError(); if (mVerbose) mTimer.restart("Process root node"); - kernels::processRoot(d_tree); + kernels::processRoot(d_tree, mStream); if (mVerbose) mTimer.stop(); cudaCheckError(); }// SignedFloodFill::operator() diff --git a/nanovdb/nanovdb/unittest/TestNanoVDB.cu b/nanovdb/nanovdb/unittest/TestNanoVDB.cu index 465fad1bce..ad562231b6 100644 --- a/nanovdb/nanovdb/unittest/TestNanoVDB.cu +++ b/nanovdb/nanovdb/unittest/TestNanoVDB.cu @@ -3607,6 +3607,132 @@ TEST(TestNanoVDBCUDA, DilateInjectPrune_ValueOnIndex) EXPECT_EQ(inputHandle.grid()->mChecksum.full(), prunedHandle.grid()->mChecksum.full()); }// DilateInjectPrune_ValueOnIndex +// Busy-wait on whatever stream it is launched on; used to occupy the default +// stream so that a kernel mistakenly launched there (instead of the caller's +// stream) is reordered relative to the rest of the operation. +__global__ void streamBusyWaitKernel(unsigned long long cycles) +{ + const unsigned long long t0 = clock64(); + while (clock64() - t0 < cycles) {} +} + +// Regression test: the NN_FACE_EDGE_VERTEX (26-neighbour) leaf dilation kernel +// must run on the caller's stream. It previously launched on the default stream, +// silently corrupting the result when the caller used a non-blocking stream. +// Dilating on the default stream and on a non-blocking stream (while the default +// stream is occupied) must produce identical output. +TEST(TestNanoVDBCUDA, NonBlockingStreamDilate_ValueOnIndex) +{ + using BuildT = nanovdb::ValueOnIndex; + std::vector pts; + for (int i = 0; i < 64; ++i) pts.emplace_back(i * 3, (i * 7) % 40, (i * 13) % 50); + pts.emplace_back(127, 127, 127); + auto inBuf = nanovdb::cuda::DeviceBuffer::create(pts.size() * sizeof(nanovdb::Coord), nullptr, false); + cudaCheck(cudaMemcpy(inBuf.deviceData(), pts.data(), pts.size() * sizeof(nanovdb::Coord), cudaMemcpyHostToDevice)); + nanovdb::tools::cuda::PointsToGrid conv; + conv.setChecksum(nanovdb::CheckMode::Full); + auto inHandle = conv.getHandle(static_cast(inBuf.deviceData()), pts.size()); + auto* inGrid = inHandle.deviceGrid(); + EXPECT_TRUE(inGrid); + cudaCheck(cudaDeviceSynchronize()); + + auto dilateOn = [&](cudaStream_t stream, bool occupyDefault) { + if (occupyDefault) {// keep the default stream busy for ~20 ms + int clockKHz = 0; + cudaCheck(cudaDeviceGetAttribute(&clockKHz, cudaDevAttrClockRate, 0)); + streamBusyWaitKernel<<<1, 1, 0, 0>>>(static_cast(clockKHz) * 20ull); + cudaCheck(cudaGetLastError()); + } + nanovdb::tools::cuda::DilateGrid dilator(inGrid, stream); + dilator.setOperation(nanovdb::tools::morphology::NN_FACE_EDGE_VERTEX); + dilator.setChecksum(nanovdb::CheckMode::Full); + dilator.setVerbose(0); + auto handle = dilator.getHandle(); + cudaCheck(cudaStreamSynchronize(stream)); + const auto treeData = nanovdb::util::cuda::DeviceGridTraits::getTreeData(handle.deviceGrid()); + handle.deviceDownload(); + return std::make_pair(treeData.mVoxelCount, handle.grid()->mChecksum.full()); + }; + + const auto reference = dilateOn(static_cast(0), false); + cudaCheck(cudaDeviceSynchronize()); + cudaStream_t nb = nullptr; + cudaCheck(cudaStreamCreateWithFlags(&nb, cudaStreamNonBlocking)); + const auto candidate = dilateOn(nb, /*occupyDefault=*/true); + cudaCheck(cudaStreamSynchronize(nb)); + cudaCheck(cudaStreamDestroy(nb)); + + EXPECT_EQ(reference.first, candidate.first); // identical active-voxel count + EXPECT_EQ(reference.second, candidate.second); // identical full checksum +}// NonBlockingStreamDilate_ValueOnIndex + +// Regression test: the grid-name copy must not over-read the source string. +// It previously copied a fixed MaxNameSize bytes from the (possibly shorter) +// std::string, over-reading host heap into the grid; now it copies only +// min(name.size()+1, MaxNameSize) bytes. Empty, short, and maximum-length +// names must round-trip. +TEST(TestNanoVDBCUDA, GridName_CudaPointsToGrid) +{ + using BuildT = nanovdb::ValueOnIndex; + const std::vector pts{ {0, 0, 0}, {1, 2, 3} }; + auto build = [&](const std::string& name) { + auto buf = nanovdb::cuda::DeviceBuffer::create(pts.size() * sizeof(nanovdb::Coord), nullptr, false); + cudaCheck(cudaMemcpy(buf.deviceData(), pts.data(), pts.size() * sizeof(nanovdb::Coord), cudaMemcpyHostToDevice)); + nanovdb::tools::cuda::PointsToGrid conv; + conv.setGridName(name); + conv.setChecksum(nanovdb::CheckMode::Full); + auto h = conv.getHandle(static_cast(buf.deviceData()), pts.size()); + h.deviceDownload(); + return std::string(h.template grid()->gridName()); + }; + EXPECT_EQ(std::string(""), build("")); // empty name + EXPECT_EQ(std::string("rho"), build("rho")); // short name + // Maximum name that fits with a null terminator (MaxNameSize includes it). + const std::string maxName(nanovdb::GridData::MaxNameSize - 1, 'x'); + EXPECT_EQ(maxName, build(maxName)); +}// GridName_CudaPointsToGrid + +// Regression test: PointsToGrid output must be byte-deterministic. Alignment +// padding and stats fields previously carried recycled pool bytes, so the output +// (and its full checksum) was nondeterministic; the buffer-wide zero-init fixes +// this. Two independent builds of the same input must have identical full +// checksums. +TEST(TestNanoVDBCUDA, DeterministicOutput_CudaPointsToGrid) +{ + using BuildT = nanovdb::ValueOnIndex; + std::vector pts; + for (int i = 0; i < 256; ++i) pts.emplace_back((i * 5) % 97, (i * 11) % 53, (i * 17) % 71); + auto buildChecksum = [&]() { + auto buf = nanovdb::cuda::DeviceBuffer::create(pts.size() * sizeof(nanovdb::Coord), nullptr, false); + cudaCheck(cudaMemcpy(buf.deviceData(), pts.data(), pts.size() * sizeof(nanovdb::Coord), cudaMemcpyHostToDevice)); + nanovdb::tools::cuda::PointsToGrid conv; + conv.setChecksum(nanovdb::CheckMode::Full); + auto h = conv.getHandle(static_cast(buf.deviceData()), pts.size()); + h.deviceDownload(); + return h.template grid()->mChecksum.full(); + }; + EXPECT_EQ(buildChecksum(), buildChecksum()); +}// DeterministicOutput_CudaPointsToGrid + +// Regression test: DeviceBuffer must free its device allocation on the stream it +// was allocated on, not the default stream. Repeated allocate/use/free on a +// non-blocking stream must complete without error (compute-sanitizer memcheck on +// this path is the stronger gate). +TEST(TestNanoVDBCUDA, DeviceBufferNonBlockingStreamLifetime) +{ + cudaStream_t stream = nullptr; + cudaCheck(cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking)); + for (int i = 0; i < 16; ++i) { + auto buf = nanovdb::cuda::DeviceBuffer::create(size_t(1) << 20, nullptr, 0, stream);// 1 MiB on stream + EXPECT_TRUE(buf.deviceData()); + cudaCheck(cudaMemsetAsync(buf.deviceData(), i & 0xff, buf.size(), stream)); + // buf is destroyed here and must free on 'stream' (still alive). + } + cudaCheck(cudaStreamSynchronize(stream)); + cudaCheck(cudaStreamDestroy(stream)); + EXPECT_EQ(cudaSuccess, cudaGetLastError()); +}// DeviceBufferNonBlockingStreamLifetime + TEST(TestNanoVDBCUDA, RefineCoarsen_ValueOnIndex) { using BuildT = nanovdb::ValueOnIndex; From 42afea5d407312805f12d1cbb68499ea7998b7c1 Mon Sep 17 00:00:00 2001 From: Jonathan Swartz Date: Tue, 14 Jul 2026 03:02:43 +0000 Subject: [PATCH 02/17] Fix up comments Signed-off-by: Jonathan Swartz --- nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh | 3 +-- nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh | 3 +-- nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh | 6 ------ 3 files changed, 2 insertions(+), 10 deletions(-) diff --git a/nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh b/nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh index ab2336b624..ca3a797d60 100644 --- a/nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh @@ -909,8 +909,7 @@ inline void DistributedPointsToGrid::processGridTreeRoot(const PtrT poin util::cuda::lambdaKernel<<<1, 1, 0, stream>>>(1, BuildGridTreeRootFunctor(), mData, mPointType, pointCount);// lambdaKernel cudaCheckError(); - // Zero the name field, then copy only the actual string (if any); see the - // matching fix in PointsToGrid.cuh (heap over-read of the name buffer). + // Zero the name field, then copy only the actual string (if any). char *dst = mData->getGrid().mGridName; cudaCheck(cudaMemsetAsync(dst, 0, GridData::MaxNameSize, stream)); if (!mGridName.empty()) { diff --git a/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh b/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh index 34b036c90e..42704206cd 100644 --- a/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh @@ -885,8 +885,7 @@ void MeshToGrid::processGridTreeRoot() cudaCheckError(); // Copy grid name into the output grid's name field. Zero the field first - // and copy only the actual string; see the matching fix in PointsToGrid.cuh - // (heap over-read of the name buffer). + // and copy only the actual string. char *dst = mBuilder.data()->getGrid().mGridName; cudaCheck(cudaMemsetAsync(dst, 0, GridData::MaxNameSize, mStream)); if (!mGridName.empty()) { diff --git a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh index 7ce9755c2b..5d79c28006 100644 --- a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh @@ -1006,9 +1006,6 @@ inline void PointsToGrid::processUpperNodes() mResource->deallocate_async(mData.d_tile_keys, mData.nodeCount[2]*sizeof(uint64_t), ResourceT::DEFAULT_ALIGNMENT, mStream); - // Upper background values are zero and the grid buffer is zero-initialized - // in getBuffer, so the former dense SetUpperBackgroundValuesFunctor pass - // (nodeCount[2]<<15 threads) is unnecessary. }// PointsToGrid::processUpperNodes //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- @@ -1050,9 +1047,6 @@ inline void PointsToGrid::processLowerNodes() util::cuda::lambdaKernel<<>>(mData.nodeCount[1], BuildLowerNodesFunctor(), mDeviceData); cudaCheckError(); - // Lower background values are zero and the grid buffer is zero-initialized - // in getBuffer, so the former dense SetLowerBackgroundValuesFunctor pass - // (nodeCount[1]<<12 threads) is unnecessary. }// PointsToGrid::processLowerNodes //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- From fb1c692d051d57760fcfd917bc90d8071ccb9f3d Mon Sep 17 00:00:00 2001 From: Jonathan Swartz Date: Tue, 14 Jul 2026 03:18:17 +0000 Subject: [PATCH 03/17] comment fix Signed-off-by: Jonathan Swartz --- 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 5d79c28006..3f4946005b 100644 --- a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh @@ -1142,7 +1142,7 @@ inline void PointsToGrid::processLeafNodes(size_t pointCount) // Inactive voxel values are zero for every build type except Point (which // copies the previous active value for rank queries); the zero cases are // covered by the buffer memset in getBuffer, and for index grids this - // dense pass (nodeCount[0]<<9 threads) was a no-op to begin with. + // dense pass (nodeCount[0]<<9 threads) is a no-op. if constexpr(util::is_same::value) { if (mVerbose==2) mTimer.restart("set inactive voxel values"); const uint64_t denseVoxelCount = mData.nodeCount[0] << 9; From 992aace6cfd2b61d28208382f95e8298233ce57d Mon Sep 17 00:00:00 2001 From: Jonathan Swartz Date: Tue, 14 Jul 2026 04:08:04 +0000 Subject: [PATCH 04/17] NanoVDB CUDA: scope IndexToGrid zero-init to the output it actually needs The Tranche-1 determinism fix zero-initialized the *entire* IndexToGrid output buffer so that bytes the build kernels never write - stats fields absent from the source grid, and struct alignment padding - are deterministic instead of carrying recycled pool bytes (which otherwise leak into written .nvdb files). Measured against master this cost 10-22% on indextogrid across dragon/emu/ crawler/wdas_cloud: the leaf array dominates the buffer and was memset to zero and then immediately overwritten value-by-value - a redundant multi-GB pass. Scope the initialization to the bytes that are genuinely left unwritten: * getBuffer() memsets only the non-leaf region [0, node[0]) - grid, tree, root, root tiles and the internal nodes - a small fraction of the buffer. * processLeafsKernel zeroes each leaf's own gap [&mMinimum, mValues): the stats fields (skipped when the source has no stats) and any padding before the 32-aligned mValues array. Every mValues[i] is written anyway, so this reproduces the former full zero-init byte-for-byte. Output is bit-identical to the full zero-init - all indextogrid goldens pass on dragon/emu/crawler/wdas_cloud, compute-sanitizer memcheck is clean and initcheck shows no new reports - and the regression vs master is gone (NEUTRAL on all four datasets, RTX PRO 6000 / sm_120). Co-Authored-By: Claude Opus 4.8 (1M context) Signed-off-by: Jonathan Swartz --- nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh | 20 ++++++++++++++++---- 1 file changed, 16 insertions(+), 4 deletions(-) diff --git a/nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh b/nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh index f5b440f9f9..6d8bde11e6 100644 --- a/nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh @@ -266,6 +266,13 @@ __global__ void processLeafsKernel(typename IndexToGrid::NodeAccessor for (int i=0; i<3; ++i) dstLeaf.mBBoxDif[i] = srcLeaf.mBBoxDif[i]; dstLeaf.mFlags = srcLeaf.mFlags; dstLeaf.mValueMask = srcLeaf.mValueMask; + // The leaf array is excluded from the buffer zero-init in getBuffer (it + // is the bulk of the grid and every mValues[i] is written below), so + // make the only otherwise-unwritten leaf bytes deterministic here: the + // stats fields (absent when the source has no stats) and any alignment + // padding before the 32-aligned mValues array. Real stats, if present, + // overwrite the zeros just below. Byte-identical to a full zero-init. + for (uint8_t *p = (uint8_t*)&dstLeaf.mMinimum, *e = (uint8_t*)dstLeaf.mValues; p < e; ++p) *p = 0; /// auto &srcGrid = nodeAcc->srcGrid(); if (srcGrid.hasMinMax()) { @@ -373,10 +380,15 @@ inline BufferT IndexToGrid::getBuffer(const BufferT &pool) auto buffer = BufferT::create(mNodeAcc.size, &pool, device, mStream); mNodeAcc.d_dstPtr = buffer.deviceData(); if (mNodeAcc.d_dstPtr == nullptr) throw std::runtime_error("Failed memory allocation on the device"); - // Zero the output buffer: regions the kernels below do not write (e.g. - // padding, unused tile fields) otherwise carry recycled allocator bytes, - // making the output nondeterministic and leaking heap contents to files. - cudaCheck(cudaMemsetAsync(mNodeAcc.d_dstPtr, 0, mNodeAcc.size, mStream)); + // Zero the non-leaf region: grid, tree, root, root tiles and the internal + // nodes. Bytes the kernels below do not explicitly write - stats fields + // absent from the source and struct alignment padding - would otherwise + // carry recycled allocator bytes, making the output nondeterministic and + // leaking heap contents into written files. The leaf array [node[0], size) + // is the bulk of the buffer and is fully overwritten by processLeafsKernel + // (values + header, which zeroes its own stats/padding gap), so it is + // excluded here to avoid a redundant multi-GB memset. + cudaCheck(cudaMemsetAsync(mNodeAcc.d_dstPtr, 0, mNodeAcc.node[0], mStream)); if (size_t size = mGridName.size()) { cudaCheck(util::cuda::mallocAsync((void**)&mNodeAcc.d_gridName, size, mStream)); From 465599ddabceb9e430b786256b98c551edf0dd7e Mon Sep 17 00:00:00 2001 From: Jonathan Swartz Date: Tue, 14 Jul 2026 04:23:13 +0000 Subject: [PATCH 05/17] NanoVDB CUDA: null-terminate over-long grid names; query current device in test Addresses PR review: - The grid-name copy could leave GridData::mGridName without a null terminator when the name length >= MaxNameSize: min(size+1, MaxNameSize) == MaxNameSize overwrote the entire zeroed field, and gridName() (a C-string accessor) would then read past it. Copy at most MaxNameSize-1 bytes and rely on the preceding memset for termination, truncating over-long names instead. Fixed in PointsToGrid, MeshToGrid and DistributedPointsToGrid. - GridName_CudaPointsToGrid gains an over-long (>= MaxNameSize) case asserting the result is truncated to MaxNameSize-1 chars and stays null-terminated. - NonBlockingStreamDilate_ValueOnIndex queried device 0 for the clock rate; query the current device via cudaGetDevice() instead. Co-Authored-By: Claude Opus 4.8 (1M context) Signed-off-by: Jonathan Swartz --- .../nanovdb/tools/cuda/DistributedPointsToGrid.cuh | 4 +++- nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh | 4 +++- nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh | 4 +++- nanovdb/nanovdb/unittest/TestNanoVDB.cu | 13 +++++++++++-- 4 files changed, 20 insertions(+), 5 deletions(-) diff --git a/nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh b/nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh index ca3a797d60..bce6614a87 100644 --- a/nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh @@ -913,7 +913,9 @@ inline void DistributedPointsToGrid::processGridTreeRoot(const PtrT poin char *dst = mData->getGrid().mGridName; cudaCheck(cudaMemsetAsync(dst, 0, GridData::MaxNameSize, stream)); if (!mGridName.empty()) { - const size_t nameSize = std::min(mGridName.size() + 1, GridData::MaxNameSize); + // Copy at most MaxNameSize-1 bytes so the memset's trailing '\0' always + // survives; a name >= MaxNameSize is truncated, never left unterminated. + const size_t nameSize = std::min(mGridName.size(), GridData::MaxNameSize - 1); cudaCheck(cudaMemcpyAsync(dst, mGridName.c_str(), nameSize, cudaMemcpyHostToDevice, stream)); } cudaEventRecord(processGridTreeRootEvent); diff --git a/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh b/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh index 42704206cd..8c65a3a0c3 100644 --- a/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh @@ -889,7 +889,9 @@ void MeshToGrid::processGridTreeRoot() char *dst = mBuilder.data()->getGrid().mGridName; cudaCheck(cudaMemsetAsync(dst, 0, GridData::MaxNameSize, mStream)); if (!mGridName.empty()) { - const size_t nameSize = std::min(mGridName.size() + 1, GridData::MaxNameSize); + // Copy at most MaxNameSize-1 bytes so the memset's trailing '\0' always + // survives; a name >= MaxNameSize is truncated, never left unterminated. + const size_t nameSize = std::min(mGridName.size(), GridData::MaxNameSize - 1); cudaCheck(cudaMemcpyAsync(dst, mGridName.c_str(), nameSize, cudaMemcpyHostToDevice, mStream)); } diff --git a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh index 3f4946005b..94844670fe 100644 --- a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh @@ -952,7 +952,9 @@ inline void PointsToGrid::processGridTreeRoot(const PtrT poin char *dst = mData.getGrid().mGridName; cudaCheck(cudaMemsetAsync(dst, 0, GridData::MaxNameSize, mStream)); if (!mGridName.empty()) { - const size_t nameSize = std::min(mGridName.size() + 1, GridData::MaxNameSize); + // Copy at most MaxNameSize-1 bytes so the memset's trailing '\0' always + // survives; a name >= MaxNameSize is truncated, never left unterminated. + const size_t nameSize = std::min(mGridName.size(), GridData::MaxNameSize - 1); cudaCheck(cudaMemcpyAsync(dst, mGridName.c_str(), nameSize, cudaMemcpyHostToDevice, mStream)); } }// PointsToGrid::processGridTreeRoot diff --git a/nanovdb/nanovdb/unittest/TestNanoVDB.cu b/nanovdb/nanovdb/unittest/TestNanoVDB.cu index ad562231b6..7a0c8dba77 100644 --- a/nanovdb/nanovdb/unittest/TestNanoVDB.cu +++ b/nanovdb/nanovdb/unittest/TestNanoVDB.cu @@ -3638,8 +3638,9 @@ TEST(TestNanoVDBCUDA, NonBlockingStreamDilate_ValueOnIndex) auto dilateOn = [&](cudaStream_t stream, bool occupyDefault) { if (occupyDefault) {// keep the default stream busy for ~20 ms - int clockKHz = 0; - cudaCheck(cudaDeviceGetAttribute(&clockKHz, cudaDevAttrClockRate, 0)); + int clockKHz = 0, dev = 0; + cudaCheck(cudaGetDevice(&dev)); + cudaCheck(cudaDeviceGetAttribute(&clockKHz, cudaDevAttrClockRate, dev)); streamBusyWaitKernel<<<1, 1, 0, 0>>>(static_cast(clockKHz) * 20ull); cudaCheck(cudaGetLastError()); } @@ -3690,6 +3691,14 @@ TEST(TestNanoVDBCUDA, GridName_CudaPointsToGrid) // Maximum name that fits with a null terminator (MaxNameSize includes it). const std::string maxName(nanovdb::GridData::MaxNameSize - 1, 'x'); EXPECT_EQ(maxName, build(maxName)); + // Over-long name (>= MaxNameSize) must be truncated to MaxNameSize-1 chars + // and stay null-terminated - never a non-terminated field that gridName() + // would over-read. The result is the max-length string, and reading it back + // as a C-string does not exceed MaxNameSize. + const std::string tooLong(nanovdb::GridData::MaxNameSize + 37, 'y'); + const std::string got = build(tooLong); + EXPECT_EQ(size_t(nanovdb::GridData::MaxNameSize - 1), got.size()); + EXPECT_EQ(std::string(nanovdb::GridData::MaxNameSize - 1, 'y'), got); }// GridName_CudaPointsToGrid // Regression test: PointsToGrid output must be byte-deterministic. Alignment From 0e5fbf3c19c1ba239cd01cb0c64de7ddd96cf2ed Mon Sep 17 00:00:00 2001 From: Jonathan Swartz Date: Tue, 14 Jul 2026 04:23:13 +0000 Subject: [PATCH 06/17] NanoVDB CUDA: only sync SignedFloodFill root pass on a non-default stream The Tranche-1 stream-correctness fix added an unconditional cudaStreamSynchronize(stream) before processRoot's synchronous default-stream cudaMemcpy of the tree. That sync is only needed when the caller's node passes ran on a non-default stream; on the default stream (0) the blocking copy already orders after prior stream-0 work, so the extra sync there was pure overhead - a measured ~4-11% regression on floodfill/dragon vs master. Guard it with `if (stream != 0)`. Correctness for non-blocking-stream callers is unchanged. floodfill is now NEUTRAL vs master on dragon/emu/crawler (0 regressions), and byte-exact validation is unaffected. Co-Authored-By: Claude Opus 4.8 (1M context) Signed-off-by: Jonathan Swartz --- nanovdb/nanovdb/tools/cuda/SignedFloodFill.cuh | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) diff --git a/nanovdb/nanovdb/tools/cuda/SignedFloodFill.cuh b/nanovdb/nanovdb/tools/cuda/SignedFloodFill.cuh index f3c5d9a189..69b7aa4e00 100644 --- a/nanovdb/nanovdb/tools/cuda/SignedFloodFill.cuh +++ b/nanovdb/nanovdb/tools/cuda/SignedFloodFill.cuh @@ -87,8 +87,11 @@ void processRoot(NanoTree *d_tree, cudaStream_t stream = 0) using ChildT = RootChild; static const int dim = int(RootT::ChildNodeType::DIM); - // Ensure the node passes issued on 'stream' have completed before the synchronous copies below read d_tree - cudaCheck(cudaStreamSynchronize(stream)); + // Ensure node passes issued on a non-default 'stream' have completed before + // the synchronous default-stream cudaMemcpy below reads d_tree. For the + // default stream (0) that blocking copy already orders after prior stream-0 + // work, so the extra sync there is pure overhead - skip it. + if (stream != cudaStream_t{0}) cudaCheck(cudaStreamSynchronize(stream)); // First copy the tree and root and then its tiles, which is of unknown size nanovdb::cuda::UnifiedBuffer uBuffer(sizeof(TreeT) + sizeof(RootT), sizeof(TreeT) + sizeof(RootT) + 64*sizeof(TileT)); From 5bb2eba69defc9f214ae750e6c10666623cbd77f Mon Sep 17 00:00:00 2001 From: Jonathan Swartz Date: Tue, 14 Jul 2026 04:44:24 +0000 Subject: [PATCH 07/17] NanoVDB CUDA: track DeviceBuffer free stream per device (multi-GPU fix) Addresses PR review: the Tranche-1 stream-ordered-free change stored a single mStream, overwritten by every DeviceBuffer::deviceUpload. A DeviceBuffer can hold an allocation on each device (mGpuData is per-device), and the multi-GPU example (ex_make_mgpu_nanovdb) uploads one handle to every device with that device's own stream. The destructor/move-assign then freed *every* device's allocation on the last device's stream - a cross-device cudaFreeAsync(ptr@devA, stream@devB), which is invalid (the free stream must belong to the allocation's device). Pre-Tranche-1 code freed on stream 0 uniformly, so this was a regression for the multi-GPU path. Track the allocation stream per device in a cudaStream_t array parallel to mGpuData, set in init()/deviceUpload() at mStreams[device], and free each mGpuData[i] on its own mStreams[i] in clear()/move-assignment. The single-GPU path is unchanged (one device, one stream); the multi-GPU path now frees each device's allocation on its own device-bound stream. Single-GPU verified (DeviceBuffer lifetime gtest + byte-exact validation); the multi-GPU free path is correct by construction but not runtime-tested here (single-GPU environment). Co-Authored-By: Claude Opus 4.8 (1M context) Signed-off-by: Jonathan Swartz --- nanovdb/nanovdb/cuda/DeviceBuffer.h | 35 ++++++++++++++++------------- 1 file changed, 19 insertions(+), 16 deletions(-) diff --git a/nanovdb/nanovdb/cuda/DeviceBuffer.h b/nanovdb/nanovdb/cuda/DeviceBuffer.h index 43e7273e34..83f34394c2 100644 --- a/nanovdb/nanovdb/cuda/DeviceBuffer.h +++ b/nanovdb/nanovdb/cuda/DeviceBuffer.h @@ -38,7 +38,7 @@ class DeviceBuffer uint64_t mSize; // total number of bytes managed by this buffer (assumed to be identical for host and device) void *mCpuData, **mGpuData; // raw pointers to the host and device buffers int mDeviceCount, mManaged;// if mManaged is non-zero this class is responsible for allocating and freeing memory buffers. Otherwise this is assumed to be handled externally - cudaStream_t mStream = 0;// stream the managed device allocations are associated with. Frees (destructor, move-assign, clear) are ordered on this stream instead of the default stream 0, otherwise a buffer last used on a non-blocking stream could be freed while that stream's work is still in flight (use-after-free). The caller must keep this stream alive at least until the buffer is destroyed. + cudaStream_t *mStreams = nullptr;// per-device stream each managed allocation was made on (parallel to mGpuData, length mDeviceCount). Frees (destructor, move-assign, clear) are ordered on the owning device's stream instead of the default stream 0, otherwise a buffer last used on a non-blocking stream could be freed while that stream's work is still in flight (use-after-free). One stream per device is required because a single DeviceBuffer can hold allocations on several devices (multi-GPU), each on its own device-bound stream. The caller must keep these streams alive until the buffer is destroyed. /// @brief Initialize buffer /// @param size byte size of buffer to be initialized @@ -123,11 +123,11 @@ class DeviceBuffer , mGpuData(other.mGpuData) , mDeviceCount(other.mDeviceCount) , mManaged(other.mManaged) - , mStream(other.mStream) + , mStreams(other.mStreams) { other.mCpuData = other.mGpuData = nullptr; + other.mStreams = nullptr; other.mSize = other.mDeviceCount = other.mManaged = 0; - other.mStream = 0; } /// @brief Copy-constructor from a HostBuffer @@ -145,10 +145,10 @@ class DeviceBuffer } /// @brief Destructor frees memory on both the host and device - /// @note Frees on the stream the buffer is associated with (mStream), not - /// the default stream, so device frees are ordered after the last - /// work issued on that stream. - ~DeviceBuffer() { this->clear(mStream); }; + /// @note Each managed device allocation is freed on the stream it was + /// allocated on (mStreams[device]), not the default stream, so device + /// frees are ordered after the last work issued on that stream. + ~DeviceBuffer() { this->clear(); }; /// @brief Static factory method that return an instance of this buffer /// @param size byte size of buffer to be initialized @@ -316,23 +316,24 @@ class DeviceBuffer inline DeviceBuffer& DeviceBuffer::operator=(DeviceBuffer&& other) noexcept { - if (mManaged) {// first free all the managed data buffers on the stream they are associated with + if (mManaged) {// first free all the managed data buffers, each on the stream its device was allocated on cudaCheck(cudaFreeHost(mCpuData)); - for (int i=0; i= cudaCpuDeviceId && device < mDeviceCount); if (device == cudaCpuDeviceId) { cudaCheck(cudaMallocHost((void**)&mCpuData, size)); // un-managed pinned memory on the host (can be slow to access!). Always 32B aligned @@ -348,10 +350,10 @@ inline void DeviceBuffer::init(uint64_t size, int device, cudaStream_t stream) } else { cudaCheck(util::cuda::mallocAsync(mGpuData+device, size, stream)); // un-managed memory on the device, always 32B aligned! checkPtr(mGpuData[device], "cuda::DeviceBuffer::init: failed to allocate device buffer"); + mStreams[device] = stream;// remember this device's allocation stream so its free is ordered on it, not stream 0 } mSize = size; mManaged = 1;// i.e. this instance is responsible for allocating and delete memory - mStream = stream;// remember the allocation stream so managed frees are ordered on it, not stream 0 } // DeviceBuffer::init inline void DeviceBuffer::deviceUpload(int device, cudaStream_t stream, bool sync) @@ -361,7 +363,7 @@ inline void DeviceBuffer::deviceUpload(int device, cudaStream_t stream, bool syn if (mGpuData[device] == nullptr) { if (mManaged==0) throw std::runtime_error("DeviceBuffer::deviceUpload called on externally managed memory that wasn\'t allocated."); cudaCheck(util::cuda::mallocAsync(mGpuData+device, mSize, stream)); // un-managed memory on the device, always 32B aligned! - mStream = stream;// remember the allocation stream so this managed device buffer is freed on it, not stream 0 + mStreams[device] = stream;// remember this device's allocation stream so it is freed on it, not stream 0 } checkPtr(mGpuData[device], "uninitialized gpu destination data"); cudaCheck(cudaMemcpyAsync(mGpuData[device], mCpuData, mSize, cudaMemcpyHostToDevice, stream)); @@ -397,17 +399,18 @@ inline void DeviceBuffer::deviceDownload(void* stream, bool sync) inline void DeviceBuffer::clear(cudaStream_t stream) { - if (mManaged) {// free all the managed data buffers on the requested stream + if (mManaged) {// free all the managed data buffers, each on the stream its device was allocated on cudaCheck(cudaFreeHost(mCpuData)); - for (int i=0; i Date: Tue, 14 Jul 2026 05:52:51 +0000 Subject: [PATCH 08/17] NanoVDB CUDA: drain default stream after NonBlockingStreamDilate busy-wait (test) Addresses PR review: the test launches streamBusyWaitKernel on the default stream to occupy it during the candidate dilation, but only synchronized the non-blocking stream afterward. The busy-wait could remain queued on stream 0 and leak into subsequent tests, causing timing-dependent flakes. Drain it with a cudaDeviceSynchronize() after the candidate run. Co-Authored-By: Claude Opus 4.8 (1M context) Signed-off-by: Jonathan Swartz --- nanovdb/nanovdb/unittest/TestNanoVDB.cu | 1 + 1 file changed, 1 insertion(+) diff --git a/nanovdb/nanovdb/unittest/TestNanoVDB.cu b/nanovdb/nanovdb/unittest/TestNanoVDB.cu index 7a0c8dba77..d25823eaaa 100644 --- a/nanovdb/nanovdb/unittest/TestNanoVDB.cu +++ b/nanovdb/nanovdb/unittest/TestNanoVDB.cu @@ -3661,6 +3661,7 @@ TEST(TestNanoVDBCUDA, NonBlockingStreamDilate_ValueOnIndex) cudaCheck(cudaStreamCreateWithFlags(&nb, cudaStreamNonBlocking)); const auto candidate = dilateOn(nb, /*occupyDefault=*/true); cudaCheck(cudaStreamSynchronize(nb)); + cudaCheck(cudaDeviceSynchronize());// drain the default-stream busy-wait so it can't leak into later tests cudaCheck(cudaStreamDestroy(nb)); EXPECT_EQ(reference.first, candidate.first); // identical active-voxel count From 04c16f94ade1a47c90fd35aeaec165271a837dcc Mon Sep 17 00:00:00 2001 From: Jonathan Swartz Date: Tue, 14 Jul 2026 22:54:59 +0000 Subject: [PATCH 09/17] NanoVDB CUDA: update DeviceBuffer free-stream on every up/download Addresses PR review: mStreams[device] was set only when the device allocation was first created. If the buffer is later re-uploaded (or downloaded) on a different stream, the destructor/move-assign/clear would still free on the old stream, which does not necessarily order after the most recent work on the new stream (notably for cudaStreamNonBlocking). Update mStreams[device] to the current stream on every host<->device copy in deviceUpload/deviceDownload, not just on first allocation. Guarded on mStreams (null for externally-managed buffers, which are never freed here). Co-Authored-By: Claude Opus 4.8 (1M context) Signed-off-by: Jonathan Swartz --- nanovdb/nanovdb/cuda/DeviceBuffer.h | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/nanovdb/nanovdb/cuda/DeviceBuffer.h b/nanovdb/nanovdb/cuda/DeviceBuffer.h index 83f34394c2..68e63cab79 100644 --- a/nanovdb/nanovdb/cuda/DeviceBuffer.h +++ b/nanovdb/nanovdb/cuda/DeviceBuffer.h @@ -363,10 +363,13 @@ inline void DeviceBuffer::deviceUpload(int device, cudaStream_t stream, bool syn if (mGpuData[device] == nullptr) { if (mManaged==0) throw std::runtime_error("DeviceBuffer::deviceUpload called on externally managed memory that wasn\'t allocated."); cudaCheck(util::cuda::mallocAsync(mGpuData+device, mSize, stream)); // un-managed memory on the device, always 32B aligned! - mStreams[device] = stream;// remember this device's allocation stream so it is freed on it, not stream 0 } checkPtr(mGpuData[device], "uninitialized gpu destination data"); cudaCheck(cudaMemcpyAsync(mGpuData[device], mCpuData, mSize, cudaMemcpyHostToDevice, stream)); + // This device buffer was last used on 'stream', so free it there (not on + // stream 0). Updated on every upload - not just first allocation - so a + // re-upload on a different stream still orders the free after its work. + if (mStreams) mStreams[device] = stream; if (sync) cudaCheck(cudaStreamSynchronize(stream)); } // DeviceBuffer::deviceUpload @@ -387,6 +390,7 @@ inline void DeviceBuffer::deviceDownload(int device, cudaStream_t stream, bool s } checkPtr(mCpuData, "uninitialized cpu destination data"); cudaCheck(cudaMemcpyAsync(mCpuData, mGpuData[device], mSize, cudaMemcpyDeviceToHost, stream)); + if (mStreams) mStreams[device] = stream;// last used on 'stream'; free this device buffer there so the free orders after this read if (sync) cudaCheck(cudaStreamSynchronize(stream)); } // DeviceBuffer::deviceDownload From bbb19a9c274d2bbe43033378d1bf3267aa25849d Mon Sep 17 00:00:00 2001 From: Jonathan Swartz Date: Wed, 15 Jul 2026 00:02:23 +0000 Subject: [PATCH 10/17] NanoVDB CUDA: include for std::min; fix DeviceBuffer stream comments Addresses PR review: - PointsToGrid/MeshToGrid/DistributedPointsToGrid use std::min in the grid-name copy but did not include (it compiled only via transitive includes). Add the include so each header is self-sufficient. - DeviceBuffer comments still described mStreams as the "allocation stream", but it is now updated on every deviceUpload/deviceDownload; reword them to "last-used stream" to match the actual behavior. Co-Authored-By: Claude Opus 4.8 (1M context) Signed-off-by: Jonathan Swartz --- nanovdb/nanovdb/cuda/DeviceBuffer.h | 12 ++++++------ .../nanovdb/tools/cuda/DistributedPointsToGrid.cuh | 1 + nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh | 1 + nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh | 1 + 4 files changed, 9 insertions(+), 6 deletions(-) diff --git a/nanovdb/nanovdb/cuda/DeviceBuffer.h b/nanovdb/nanovdb/cuda/DeviceBuffer.h index 68e63cab79..ccad461207 100644 --- a/nanovdb/nanovdb/cuda/DeviceBuffer.h +++ b/nanovdb/nanovdb/cuda/DeviceBuffer.h @@ -38,7 +38,7 @@ class DeviceBuffer uint64_t mSize; // total number of bytes managed by this buffer (assumed to be identical for host and device) void *mCpuData, **mGpuData; // raw pointers to the host and device buffers int mDeviceCount, mManaged;// if mManaged is non-zero this class is responsible for allocating and freeing memory buffers. Otherwise this is assumed to be handled externally - cudaStream_t *mStreams = nullptr;// per-device stream each managed allocation was made on (parallel to mGpuData, length mDeviceCount). Frees (destructor, move-assign, clear) are ordered on the owning device's stream instead of the default stream 0, otherwise a buffer last used on a non-blocking stream could be freed while that stream's work is still in flight (use-after-free). One stream per device is required because a single DeviceBuffer can hold allocations on several devices (multi-GPU), each on its own device-bound stream. The caller must keep these streams alive until the buffer is destroyed. + cudaStream_t *mStreams = nullptr;// per-device stream each managed device buffer was last used on - set at allocation, then updated by every deviceUpload/deviceDownload (parallel to mGpuData, length mDeviceCount). Frees (destructor, move-assign, clear) are ordered on the owning device's stream instead of the default stream 0, otherwise a buffer last used on a non-blocking stream could be freed while that stream's work is still in flight (use-after-free). One stream per device is required because a single DeviceBuffer can hold allocations on several devices (multi-GPU), each on its own device-bound stream. The caller must keep these streams alive until the buffer is destroyed. /// @brief Initialize buffer /// @param size byte size of buffer to be initialized @@ -145,8 +145,8 @@ class DeviceBuffer } /// @brief Destructor frees memory on both the host and device - /// @note Each managed device allocation is freed on the stream it was - /// allocated on (mStreams[device]), not the default stream, so device + /// @note Each managed device allocation is freed on the stream it was last + /// used on (mStreams[device]), not the default stream, so device /// frees are ordered after the last work issued on that stream. ~DeviceBuffer() { this->clear(); }; @@ -316,7 +316,7 @@ class DeviceBuffer inline DeviceBuffer& DeviceBuffer::operator=(DeviceBuffer&& other) noexcept { - if (mManaged) {// first free all the managed data buffers, each on the stream its device was allocated on + if (mManaged) {// first free all the managed data buffers, each on the stream its device was last used on cudaCheck(cudaFreeHost(mCpuData)); for (int i=0; i #include #include +#include // std::min in processGridTreeRoot's grid-name copy namespace nanovdb { diff --git a/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh b/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh index 8c65a3a0c3..c9521b8001 100644 --- a/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh @@ -17,6 +17,7 @@ #define NVIDIA_TOOLS_CUDA_MESHTOGRID_CUH_HAS_BEEN_INCLUDED #include +#include // std::min in processGridTreeRoot's grid-name copy #include #include diff --git a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh index 94844670fe..d2aae62da4 100644 --- a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh @@ -22,6 +22,7 @@ #include #include #include +#include // std::min in processGridTreeRoot's grid-name copy #include #include From f20dd2a50eacc50977a626ecbcd50eaf9ec5433c Mon Sep 17 00:00:00 2001 From: Jonathan Swartz Date: Thu, 16 Jul 2026 23:00:35 +0000 Subject: [PATCH 11/17] NanoVDB: add pending CHANGES entries for the CUDA correctness fixes Co-Authored-By: Claude Opus 4.8 (1M context) Signed-off-by: Jonathan Swartz --- pendingchanges/nanovdbcorrectness.txt | 13 +++++++++++++ 1 file changed, 13 insertions(+) create mode 100644 pendingchanges/nanovdbcorrectness.txt diff --git a/pendingchanges/nanovdbcorrectness.txt b/pendingchanges/nanovdbcorrectness.txt new file mode 100644 index 0000000000..368c7c2ada --- /dev/null +++ b/pendingchanges/nanovdbcorrectness.txt @@ -0,0 +1,13 @@ +NanoVDB: + Bug Fixes: + - Fixed an out-of-bounds device read and host stack corruption in + tools::cuda::evalChecksum. + - Fixed tools::cuda::DilateGrid and signedFloodFill launching kernels on the + default stream instead of the caller's stream, corrupting output on + non-blocking streams. + - Fixed a host-heap over-read when copying grid names in the CUDA + PointsToGrid, DistributedPointsToGrid and MeshToGrid tools. + - Fixed nondeterministic output from the CUDA PointsToGrid and IndexToGrid tools. + - Fixed cuda::DeviceBuffer freeing device memory on the default stream rather + than the stream it was last used on, a latent use-after-free for + non-blocking-stream callers. From 59fcab450afb412ebb89a6b081e1458050692e27 Mon Sep 17 00:00:00 2001 From: Jonathan Swartz Date: Mon, 20 Jul 2026 00:43:38 +0000 Subject: [PATCH 12/17] Comment fix Signed-off-by: Jonathan Swartz --- nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh | 2 +- nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh | 2 +- nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh | 2 +- 3 files changed, 3 insertions(+), 3 deletions(-) diff --git a/nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh b/nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh index b3165b27ac..1591749b58 100644 --- a/nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh @@ -20,7 +20,7 @@ #include #include #include -#include // std::min in processGridTreeRoot's grid-name copy +#include namespace nanovdb { diff --git a/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh b/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh index c9521b8001..81709e39ec 100644 --- a/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh @@ -17,7 +17,7 @@ #define NVIDIA_TOOLS_CUDA_MESHTOGRID_CUH_HAS_BEEN_INCLUDED #include -#include // std::min in processGridTreeRoot's grid-name copy +#include #include #include diff --git a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh index d2aae62da4..4adebafe33 100644 --- a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh @@ -22,7 +22,7 @@ #include #include #include -#include // std::min in processGridTreeRoot's grid-name copy +#include #include #include From e3351803dce85f99bba37bd4b41580f748e63561 Mon Sep 17 00:00:00 2001 From: Jonathan Swartz Date: Fri, 7 Aug 2026 21:55:45 +0000 Subject: [PATCH 13/17] NanoVDB: order DeviceBuffer device frees after every stream that used the buffer Tracking only the stream a device buffer was LAST used on is not sufficient: two streams can both have work in flight, and freeing on the later one leaves the earlier one's work running against memory the allocator has already recycled. Reported in review by matthewdcong. Reproduced: a device-only DeviceBuffer plus a user kernel on a second blocking stream, with a warm-up pass so the launches actually overlap. The freed block is recycled and all 64 MB of the next allocation are overwritten by the late kernel. Freeing on the default stream (the previous behaviour) passes, because the legacy stream is implicitly ordered after every blocking stream; freeing on the last-used stream fails. Replace the per-device stream with a per-device event. Each use waits on the event before issuing work and re-records it afterwards, so one event transitively covers every stream the buffer has been used on, and the frees wait on it. This is correct for any mix of blocking and non-blocking streams and for any number of them, and it removes the requirement that callers keep their streams alive until the buffer is destroyed, since the event is owned by the buffer. Verified: both the multi-blocking-stream case and the non-blocking case pass, and the VBM, topology and maintenance goldens are unchanged (46/46 on dragon and emu). Signed-off-by: Jonathan Swartz --- nanovdb/nanovdb/cuda/DeviceBuffer.h | 82 +++++++++++++++++++++-------- 1 file changed, 60 insertions(+), 22 deletions(-) diff --git a/nanovdb/nanovdb/cuda/DeviceBuffer.h b/nanovdb/nanovdb/cuda/DeviceBuffer.h index ccad461207..8735859683 100644 --- a/nanovdb/nanovdb/cuda/DeviceBuffer.h +++ b/nanovdb/nanovdb/cuda/DeviceBuffer.h @@ -38,7 +38,7 @@ class DeviceBuffer uint64_t mSize; // total number of bytes managed by this buffer (assumed to be identical for host and device) void *mCpuData, **mGpuData; // raw pointers to the host and device buffers int mDeviceCount, mManaged;// if mManaged is non-zero this class is responsible for allocating and freeing memory buffers. Otherwise this is assumed to be handled externally - cudaStream_t *mStreams = nullptr;// per-device stream each managed device buffer was last used on - set at allocation, then updated by every deviceUpload/deviceDownload (parallel to mGpuData, length mDeviceCount). Frees (destructor, move-assign, clear) are ordered on the owning device's stream instead of the default stream 0, otherwise a buffer last used on a non-blocking stream could be freed while that stream's work is still in flight (use-after-free). One stream per device is required because a single DeviceBuffer can hold allocations on several devices (multi-GPU), each on its own device-bound stream. The caller must keep these streams alive until the buffer is destroyed. + cudaEvent_t *mEvents = nullptr;// per-device event marking the last use of each managed device buffer (parallel to mGpuData, length mDeviceCount). Every use waits on this event before issuing work and re-records it afterwards, so the single event transitively covers EVERY stream the buffer has been used on. Frees then wait on it, which orders them after all outstanding work: freeing on the default stream alone is only safe for blocking streams, and freeing on the last-used stream alone is only safe when just one stream was used. /// @brief Initialize buffer /// @param size byte size of buffer to be initialized @@ -47,6 +47,44 @@ class DeviceBuffer /// @warning size is expected to be non-zero. Use clear() clear buffer! void init(uint64_t size, int device, cudaStream_t stream); + /// @brief Order work subsequently issued on @a stream after every prior use of this + /// device buffer, whichever stream those uses were issued on. + void orderAfterPriorUses(int device, cudaStream_t stream) const + { + if (mEvents && mEvents[device]) cudaCheck(cudaStreamWaitEvent(stream, mEvents[device], 0)); + } + + /// @brief Record that this device buffer has just been used on @a stream. Paired with + /// orderAfterPriorUses this keeps a single event that covers all prior uses. + void recordUse(int device, cudaStream_t stream) + { + if (!mEvents) return; + if (mEvents[device] == nullptr) {// events are per-device, so create it on the right one + int current = 0; + cudaCheck(cudaGetDevice(¤t)); + if (current != device) cudaCheck(cudaSetDevice(device)); + cudaCheck(cudaEventCreateWithFlags(&mEvents[device], cudaEventDisableTiming)); + if (current != device) cudaCheck(cudaSetDevice(current)); + } + cudaCheck(cudaEventRecord(mEvents[device], stream)); + } + + /// @brief Free every managed device allocation on @a stream, ordered after all outstanding + /// work on the buffer, and destroy the tracking events. + /// @note Destroying an event with a pending wait is safe: CUDA releases it once the device + /// has completed it. + void freeDeviceBuffers(cudaStream_t stream) + { + for (int i = 0; i < mDeviceCount; ++i) { + this->orderAfterPriorUses(i, stream); + cudaCheck(util::cuda::freeAsync(mGpuData[i], stream)); + if (mEvents && mEvents[i]) { + cudaCheck(cudaEventDestroy(mEvents[i])); + mEvents[i] = nullptr; + } + } + } + public: using PtrT = std::shared_ptr; @@ -123,10 +161,10 @@ class DeviceBuffer , mGpuData(other.mGpuData) , mDeviceCount(other.mDeviceCount) , mManaged(other.mManaged) - , mStreams(other.mStreams) + , mEvents(other.mEvents) { other.mCpuData = other.mGpuData = nullptr; - other.mStreams = nullptr; + other.mEvents = nullptr; other.mSize = other.mDeviceCount = other.mManaged = 0; } @@ -145,9 +183,8 @@ class DeviceBuffer } /// @brief Destructor frees memory on both the host and device - /// @note Each managed device allocation is freed on the stream it was last - /// used on (mStreams[device]), not the default stream, so device - /// frees are ordered after the last work issued on that stream. + /// @note Each managed device free waits on that device's tracking event first, so it is + /// ordered after every stream the buffer was used on, not just the most recent one. ~DeviceBuffer() { this->clear(); }; /// @brief Static factory method that return an instance of this buffer @@ -316,21 +353,21 @@ class DeviceBuffer inline DeviceBuffer& DeviceBuffer::operator=(DeviceBuffer&& other) noexcept { - if (mManaged) {// first free all the managed data buffers, each on the stream its device was last used on + if (mManaged) {// first free all the managed data buffers, ordered after every use of each cudaCheck(cudaFreeHost(mCpuData)); - for (int i=0; ifreeDeviceBuffers(cudaStream_t{0}); } delete [] mGpuData; - delete [] mStreams; + delete [] mEvents; mSize = other.mSize; mCpuData = other.mCpuData; mGpuData = other.mGpuData; mDeviceCount = other.mDeviceCount; mManaged = other.mManaged; - mStreams = other.mStreams; + mEvents = other.mEvents; other.mCpuData = nullptr; other.mGpuData = nullptr; - other.mStreams = nullptr; + other.mEvents = nullptr; other.mSize = 0; other.mDeviceCount = 0; other.mManaged = 0; @@ -342,7 +379,7 @@ inline void DeviceBuffer::init(uint64_t size, int device, cudaStream_t stream) if (size==0) return; cudaCheck(cudaGetDeviceCount(&mDeviceCount)); mGpuData = new void*[mDeviceCount]();// NULL initialization - mStreams = new cudaStream_t[mDeviceCount]();// zero (default-stream) initialization + mEvents = new cudaEvent_t[mDeviceCount]();// NULL initialization; created lazily on first use NANOVDB_ASSERT(device >= cudaCpuDeviceId && device < mDeviceCount); if (device == cudaCpuDeviceId) { cudaCheck(cudaMallocHost((void**)&mCpuData, size)); // un-managed pinned memory on the host (can be slow to access!). Always 32B aligned @@ -350,7 +387,7 @@ inline void DeviceBuffer::init(uint64_t size, int device, cudaStream_t stream) } else { cudaCheck(util::cuda::mallocAsync(mGpuData+device, size, stream)); // un-managed memory on the device, always 32B aligned! checkPtr(mGpuData[device], "cuda::DeviceBuffer::init: failed to allocate device buffer"); - mStreams[device] = stream;// record the stream this device buffer is used on so its free is ordered on it, not stream 0 + this->recordUse(device, stream);// the free must be ordered after this allocation } mSize = size; mManaged = 1;// i.e. this instance is responsible for allocating and delete memory @@ -365,11 +402,11 @@ inline void DeviceBuffer::deviceUpload(int device, cudaStream_t stream, bool syn cudaCheck(util::cuda::mallocAsync(mGpuData+device, mSize, stream)); // un-managed memory on the device, always 32B aligned! } checkPtr(mGpuData[device], "uninitialized gpu destination data"); + // Order this transfer after any use of the buffer on another stream, then mark it as the + // latest use, so the tracking event keeps covering every stream the buffer has seen. + this->orderAfterPriorUses(device, stream); cudaCheck(cudaMemcpyAsync(mGpuData[device], mCpuData, mSize, cudaMemcpyHostToDevice, stream)); - // This device buffer was last used on 'stream', so free it there (not on - // stream 0). Updated on every upload - not just first allocation - so a - // re-upload on a different stream still orders the free after its work. - if (mStreams) mStreams[device] = stream; + this->recordUse(device, stream); if (sync) cudaCheck(cudaStreamSynchronize(stream)); } // DeviceBuffer::deviceUpload @@ -389,8 +426,9 @@ inline void DeviceBuffer::deviceDownload(int device, cudaStream_t stream, bool s cudaCheck(cudaMallocHost((void**)&mCpuData, mSize)); // un-managed pinned memory on the host (can be slow to access!). Always 32B aligned } checkPtr(mCpuData, "uninitialized cpu destination data"); + this->orderAfterPriorUses(device, stream); cudaCheck(cudaMemcpyAsync(mCpuData, mGpuData[device], mSize, cudaMemcpyDeviceToHost, stream)); - if (mStreams) mStreams[device] = stream;// last used on 'stream'; free this device buffer there so the free orders after this read + this->recordUse(device, stream); if (sync) cudaCheck(cudaStreamSynchronize(stream)); } // DeviceBuffer::deviceDownload @@ -403,15 +441,15 @@ inline void DeviceBuffer::deviceDownload(void* stream, bool sync) inline void DeviceBuffer::clear(cudaStream_t stream) { - if (mManaged) {// free all the managed data buffers, each on the stream its device was last used on + if (mManaged) {// free all the managed data buffers, ordered after every use of each cudaCheck(cudaFreeHost(mCpuData)); - for (int i=0; ifreeDeviceBuffers(stream); } delete [] mGpuData; - delete [] mStreams; + delete [] mEvents; mCpuData = nullptr; mGpuData = nullptr; - mStreams = nullptr; + mEvents = nullptr; mSize = 0; mDeviceCount = 0; mManaged = 0; From 5b5e936723361bed669637fa9e9bbde1a344db1b Mon Sep 17 00:00:00 2001 From: Jonathan Swartz Date: Fri, 7 Aug 2026 22:04:07 +0000 Subject: [PATCH 14/17] NanoVDB: update the DeviceBuffer entry in pending changes The entry described the mechanism as freeing on the stream the buffer was last used on, which was superseded: frees are now ordered after every stream the buffer was used on via a per-device event. Signed-off-by: Jonathan Swartz --- pendingchanges/nanovdbcorrectness.txt | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/pendingchanges/nanovdbcorrectness.txt b/pendingchanges/nanovdbcorrectness.txt index 368c7c2ada..631144e67c 100644 --- a/pendingchanges/nanovdbcorrectness.txt +++ b/pendingchanges/nanovdbcorrectness.txt @@ -8,6 +8,9 @@ NanoVDB: - Fixed a host-heap over-read when copying grid names in the CUDA PointsToGrid, DistributedPointsToGrid and MeshToGrid tools. - Fixed nondeterministic output from the CUDA PointsToGrid and IndexToGrid tools. - - Fixed cuda::DeviceBuffer freeing device memory on the default stream rather - than the stream it was last used on, a latent use-after-free for - non-blocking-stream callers. + - Fixed cuda::DeviceBuffer freeing device memory without ordering the free + after work still outstanding on the buffer, a latent use-after-free for + callers using non-blocking streams. Each managed device allocation now + tracks its uses with an event, so a free is ordered after every stream the + buffer was used on rather than only those the freeing stream happens to + synchronize with. From f95115304efb903db16e3aa3ea82e24eed410cd4 Mon Sep 17 00:00:00 2001 From: Jonathan Swartz Date: Fri, 7 Aug 2026 22:20:16 +0000 Subject: [PATCH 15/17] NanoVDB: add a regression test for DeviceBuffer multi-stream free ordering The existing DeviceBufferNonBlockingStreamLifetime test only asserts that repeated allocate/use/free returns cudaSuccess, so it passes with or without a correctly ordered free. Its comment also described the superseded 'free on the stream it was allocated on' mechanism. Reword it as the smoke test it is, and add a test that actually discriminates. DeviceBufferMultiStreamFreeOrdering uses a device-only buffer on purpose - it owns no pinned host memory, so clear()'s cudaFreeHost (which implicitly synchronizes) cannot mask the problem. A second stream is parked behind a busy-wait with a write to the buffer queued behind it, the buffer is destroyed, and the next allocation is checked for corruption. Verified to fail against the pre-fix DeviceBuffer (64 MB clobbered) and pass with the fix. A warm-up pass is required: on a cold context the first launches serialize and the race never forms. The recycled/still-pending state is reported in the failure message but deliberately not used as a preconditon - with a correctly ordered free the allocator cannot hand the block out again until the other stream drains, so those observations are the fix working rather than a reason to skip. Signed-off-by: Jonathan Swartz --- nanovdb/nanovdb/unittest/TestNanoVDB.cu | 85 +++++++++++++++++++++++-- 1 file changed, 81 insertions(+), 4 deletions(-) diff --git a/nanovdb/nanovdb/unittest/TestNanoVDB.cu b/nanovdb/nanovdb/unittest/TestNanoVDB.cu index d25823eaaa..a69cb384be 100644 --- a/nanovdb/nanovdb/unittest/TestNanoVDB.cu +++ b/nanovdb/nanovdb/unittest/TestNanoVDB.cu @@ -3724,10 +3724,9 @@ TEST(TestNanoVDBCUDA, DeterministicOutput_CudaPointsToGrid) EXPECT_EQ(buildChecksum(), buildChecksum()); }// DeterministicOutput_CudaPointsToGrid -// Regression test: DeviceBuffer must free its device allocation on the stream it -// was allocated on, not the default stream. Repeated allocate/use/free on a -// non-blocking stream must complete without error (compute-sanitizer memcheck on -// this path is the stronger gate). +// Smoke test: repeated allocate/use/free on a non-blocking stream must complete without error. +// Note this is only a smoke test - DeviceBufferMultiStreamFreeOrdering below is what actually +// discriminates a premature free (compute-sanitizer memcheck on this path is a further gate). TEST(TestNanoVDBCUDA, DeviceBufferNonBlockingStreamLifetime) { cudaStream_t stream = nullptr; @@ -3743,6 +3742,84 @@ TEST(TestNanoVDBCUDA, DeviceBufferNonBlockingStreamLifetime) EXPECT_EQ(cudaSuccess, cudaGetLastError()); }// DeviceBufferNonBlockingStreamLifetime +__global__ void deviceBufferFillKernel(unsigned char *p, size_t n, unsigned char v) +{ + for (size_t i = blockIdx.x*(size_t)blockDim.x + threadIdx.x; i < n; i += (size_t)gridDim.x*blockDim.x) p[i] = v; +} +__global__ void deviceBufferCountKernel(const unsigned char *p, size_t n, unsigned char v, unsigned long long *bad) +{ + for (size_t i = blockIdx.x*(size_t)blockDim.x + threadIdx.x; i < n; i += (size_t)gridDim.x*blockDim.x) + if (p[i] != v) atomicAdd(bad, 1ull); +} + +// Regression test: a DeviceBuffer can be used on more than one stream, so ordering its device +// free on any single stream is not sufficient - it has to be ordered after every stream that +// touched the buffer. A device-only buffer is used here on purpose: it owns no pinned host +// memory, so clear()'s cudaFreeHost (which implicitly synchronizes) cannot mask the problem. +// +// 'user' is parked behind a spin kernel with a write to the buffer queued behind it, then the +// buffer is destroyed. If the free is not ordered after 'user', the allocator recycles the block +// into the next allocation and the late write corrupts it. +TEST(TestNanoVDBCUDA, DeviceBufferMultiStreamFreeOrdering) +{ + const size_t N = size_t(64) << 20;// large enough that a recycled block is the same address + const unsigned char LATE = 0xAA, VICTIM = 0x55; + const unsigned long long CYCLES = 400000000ull;// parks 'user' for O(100 ms) + + cudaStream_t user = nullptr, other = nullptr; + cudaCheck(cudaStreamCreate(&user));// blocking streams, i.e. the ones the legacy + cudaCheck(cudaStreamCreate(&other));// default stream implicitly synchronizes with + unsigned long long *bad = nullptr; + cudaCheck(cudaMallocManaged(&bad, sizeof(*bad))); + + {// warm-up: on a cold context the first launches serialize, which would hide the race + unsigned char *w = nullptr; + cudaCheck(cudaMallocAsync((void**)&w, N, other)); + streamBusyWaitKernel<<<1,1,0,user>>>(CYCLES/10); + deviceBufferFillKernel<<<1024,256,0,other>>>(w, N, 0); + deviceBufferCountKernel<<<1024,256,0,other>>>(w, N, 0, bad); + cudaCheck(cudaFreeAsync(w, other)); + cudaCheck(cudaDeviceSynchronize()); + } + + void *devPtr = nullptr; + { + auto buf = nanovdb::cuda::DeviceBuffer::create(N, nullptr, 0, other);// device-only + devPtr = buf.deviceData(0); + ASSERT_TRUE(devPtr); + streamBusyWaitKernel<<<1,1,0,user>>>(CYCLES);// park 'user' + deviceBufferFillKernel<<<1024,256,0,user>>>((unsigned char*)devPtr, N, LATE); + }// destroyed here; the free must be ordered after 'user' as well as 'other' + + unsigned char *victim = nullptr; + cudaCheck(cudaMallocAsync((void**)&victim, N, other)); + deviceBufferFillKernel<<<1024,256,0,other>>>(victim, N, VICTIM); + cudaCheck(cudaStreamSynchronize(other)); + + // Diagnostics only. These are deliberately NOT preconditions: with a correctly ordered free + // the allocator cannot hand this block out again until 'user' has drained, so observing + // "not recycled" or "no longer pending" here is the fix working, not a reason to skip. + const bool stillPending = (cudaStreamQuery(user) == cudaErrorNotReady); + cudaGetLastError();// clear the cudaErrorNotReady left by the query above + const bool recycled = (victim == devPtr); + + cudaCheck(cudaStreamSynchronize(user));// let the late write land + *bad = 0; + deviceBufferCountKernel<<<1024,256>>>(victim, N, VICTIM, bad); + cudaCheck(cudaDeviceSynchronize()); + const unsigned long long clobbered = *bad; + + cudaCheck(cudaFreeAsync(victim, other)); + cudaCheck(cudaStreamSynchronize(other)); + cudaCheck(cudaFree(bad)); + cudaCheck(cudaStreamDestroy(user)); + cudaCheck(cudaStreamDestroy(other)); + + EXPECT_EQ(0u, clobbered) << "device memory was freed while another stream still had work in " + "flight (block recycled: " << recycled + << ", work still pending when it was reused: " << stillPending << ")"; +}// DeviceBufferMultiStreamFreeOrdering + TEST(TestNanoVDBCUDA, RefineCoarsen_ValueOnIndex) { using BuildT = nanovdb::ValueOnIndex; From 111b12f9c10bcf19d78415c8ed3d81cc334b414e Mon Sep 17 00:00:00 2001 From: Jonathan Swartz Date: Fri, 7 Aug 2026 22:31:52 +0000 Subject: [PATCH 16/17] NanoVDB: guard DeviceBuffer self-move and document clear()'s stream argument Two non-blocking review points from harrism. Self-move assignment freed the buffers and then read its own (now dangling) members. Verified: without the guard 'buf = std::move(buf)' aborts on the deviceData assertion because mDeviceCount has been zeroed; with it the buffer is left intact. Note UnifiedBuffer::operator=(&&) has the same flaw and is left for a separate change. Also document clear()'s stream argument. The review noted it had become dead, which was true of the revision reviewed - the per-device stream overrode it. With frees now ordered by the tracking event the argument is honoured again and simply selects where the free is enqueued, so any stream is safe to pass regardless of where the buffer was used. Signed-off-by: Jonathan Swartz --- nanovdb/nanovdb/cuda/DeviceBuffer.h | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/nanovdb/nanovdb/cuda/DeviceBuffer.h b/nanovdb/nanovdb/cuda/DeviceBuffer.h index 8735859683..ba52ba6056 100644 --- a/nanovdb/nanovdb/cuda/DeviceBuffer.h +++ b/nanovdb/nanovdb/cuda/DeviceBuffer.h @@ -344,6 +344,10 @@ class DeviceBuffer /// @} /// @brief De-allocate all memory managed by this allocator and set all pointers to NULL + /// @param stream Stream the device frees are issued on. The frees are additionally ordered + /// after every stream the buffer was used on (via the per-device tracking event), so + /// @a stream selects where the free is enqueued, not what it is ordered against - any + /// stream is safe to pass here regardless of where the buffer was used. void clear(cudaStream_t stream = 0); void clear(void* stream){this->clear(cudaStream_t(stream));} @@ -353,6 +357,7 @@ class DeviceBuffer inline DeviceBuffer& DeviceBuffer::operator=(DeviceBuffer&& other) noexcept { + if (this == &other) return *this;// self-move would free our buffers and then read them back if (mManaged) {// first free all the managed data buffers, ordered after every use of each cudaCheck(cudaFreeHost(mCpuData)); this->freeDeviceBuffers(cudaStream_t{0}); From 457b671c34b1e047c0aec880a4a1715591a26f41 Mon Sep 17 00:00:00 2001 From: Jonathan Swartz Date: Fri, 7 Aug 2026 23:35:23 +0000 Subject: [PATCH 17/17] NanoVDB: public DeviceBuffer::recordUse, per-device frees, and stronger stream tests Review follow-ups from the Copilot pass on this PR. The lifetime event only covers uses issued through deviceUpload/deviceDownload; work enqueued against the raw pointer from deviceData() is invisible to it, so on a non-blocking stream such work could still outlive the buffer (as on master, where no API could express the dependency). Make recordUse public so raw-pointer callers can register their stream, and document the contract on deviceData(). freeDeviceBuffers issued every device's free on one stream, but a stream belongs to a single device; free each allocation on its own device instead (caller's stream for the current device, the owning device's default stream otherwise, switching devices as needed). Same pre-existing shape as master's loop; correctness on multi-GPU is by inspection - this machine has one GPU. Split the free-ordering regression test into the two scenarios that discriminate the two historical bugs: a blocking stream with an unregistered write (fails if frees move off the default stream: the reviewed revision) and a non-blocking stream with a recordUse-registered write (fails without the event: master). Verified per revision: current passes both, the reviewed revision fails both, master fails the non-blocking one. Add NonBlockingStreamSignedFloodFill, mirroring the dilation regression: identical output on the default stream and on a non-blocking stream while the default stream is occupied, with the readback on the producing stream. This locks the cross-stream contract; it cannot discriminate processRoot's internal ordering on constructible inputs (verified: the pre-fix header passes it) because that path only does work when interior root-level tiles exist and pre-fix code was accidentally host-synchronous otherwise - noted in the test. Signed-off-by: Jonathan Swartz --- nanovdb/nanovdb/cuda/DeviceBuffer.h | 59 ++++++++++------ nanovdb/nanovdb/unittest/TestNanoVDB.cu | 90 ++++++++++++++++++++++--- 2 files changed, 120 insertions(+), 29 deletions(-) diff --git a/nanovdb/nanovdb/cuda/DeviceBuffer.h b/nanovdb/nanovdb/cuda/DeviceBuffer.h index ba52ba6056..f542bfdf10 100644 --- a/nanovdb/nanovdb/cuda/DeviceBuffer.h +++ b/nanovdb/nanovdb/cuda/DeviceBuffer.h @@ -54,30 +54,26 @@ class DeviceBuffer if (mEvents && mEvents[device]) cudaCheck(cudaStreamWaitEvent(stream, mEvents[device], 0)); } - /// @brief Record that this device buffer has just been used on @a stream. Paired with - /// orderAfterPriorUses this keeps a single event that covers all prior uses. - void recordUse(int device, cudaStream_t stream) - { - if (!mEvents) return; - if (mEvents[device] == nullptr) {// events are per-device, so create it on the right one - int current = 0; - cudaCheck(cudaGetDevice(¤t)); - if (current != device) cudaCheck(cudaSetDevice(device)); - cudaCheck(cudaEventCreateWithFlags(&mEvents[device], cudaEventDisableTiming)); - if (current != device) cudaCheck(cudaSetDevice(current)); - } - cudaCheck(cudaEventRecord(mEvents[device], stream)); - } - - /// @brief Free every managed device allocation on @a stream, ordered after all outstanding - /// work on the buffer, and destroy the tracking events. + /// @brief Free every managed device allocation, each ordered after all tracked uses of the + /// buffer, and destroy the tracking events. + /// @param stream Stream the frees are issued on for allocations owned by the CURRENT device. + /// A stream belongs to a single device, so it cannot carry frees for other devices' + /// memory pools; allocations on other devices are freed on their own device's default + /// stream, after switching to that device. /// @note Destroying an event with a pending wait is safe: CUDA releases it once the device /// has completed it. void freeDeviceBuffers(cudaStream_t stream) { + int current = 0; + cudaCheck(cudaGetDevice(¤t)); for (int i = 0; i < mDeviceCount; ++i) { - this->orderAfterPriorUses(i, stream); - cudaCheck(util::cuda::freeAsync(mGpuData[i], stream)); + if (mGpuData[i]) { + const cudaStream_t freeStream = (i == current) ? stream : cudaStream_t{0}; + if (i != current) cudaCheck(cudaSetDevice(i)); + this->orderAfterPriorUses(i, freeStream); + cudaCheck(util::cuda::freeAsync(mGpuData[i], freeStream)); + if (i != current) cudaCheck(cudaSetDevice(current)); + } if (mEvents && mEvents[i]) { cudaCheck(cudaEventDestroy(mEvents[i])); mEvents[i] = nullptr; @@ -269,8 +265,33 @@ class DeviceBuffer /////////////////////////////////////////////////////////////////////// + /// @brief Record that this buffer's device data was just used on @a stream, so that the + /// buffer's device frees (destructor, move-assignment, clear) are ordered after that + /// work. Uses issued through deviceUpload/deviceDownload are recorded automatically; + /// callers that enqueue their own kernels or copies against the raw pointer returned + /// by deviceData() should call this afterwards. Without it, such work is only safe if + /// it is on a blocking stream (which the free, issued on the default stream, waits on + /// implicitly) or if the caller synchronizes before the buffer is cleared/destroyed. + /// @param device Device whose buffer was used + /// @param stream Stream the work was issued on + void recordUse(int device, cudaStream_t stream) + { + if (!mEvents) return; + if (mEvents[device] == nullptr) {// events are per-device, so create it on the right one + int current = 0; + cudaCheck(cudaGetDevice(¤t)); + if (current != device) cudaCheck(cudaSetDevice(device)); + cudaCheck(cudaEventCreateWithFlags(&mEvents[device], cudaEventDisableTiming)); + if (current != device) cudaCheck(cudaSetDevice(current)); + } + cudaCheck(cudaEventRecord(mEvents[device], stream)); + } + /// @brief Retuns a raw pointer to the specified device/GPU buffer managed by this allocator. /// @warning Note that the pointer can be NULL! + /// @note Work enqueued against this raw pointer is invisible to the buffer's lifetime + /// tracking: on a non-blocking stream, call recordUse afterwards (or synchronize + /// before the buffer is cleared/destroyed) so the device free is ordered after it. void* deviceData(int device) const { NANOVDB_ASSERT(device >= 0 && device < mDeviceCount); return mGpuData[device]; diff --git a/nanovdb/nanovdb/unittest/TestNanoVDB.cu b/nanovdb/nanovdb/unittest/TestNanoVDB.cu index a69cb384be..215a8a60be 100644 --- a/nanovdb/nanovdb/unittest/TestNanoVDB.cu +++ b/nanovdb/nanovdb/unittest/TestNanoVDB.cu @@ -3668,6 +3668,55 @@ TEST(TestNanoVDBCUDA, NonBlockingStreamDilate_ValueOnIndex) EXPECT_EQ(reference.second, candidate.second); // identical full checksum }// NonBlockingStreamDilate_ValueOnIndex +// Cross-stream determinism coverage: signedFloodFill must run entirely on the caller's stream. +// Mirrors the dilation test above: repair the same corrupted level set once on the default +// stream and once on a non-blocking stream while the default stream is occupied, reading each +// result back on the stream that produced it (no default-stream rescue), and require identical +// output. Node passes escaping to the occupied default stream show up as an unrepaired or +// partially repaired grid. +// +// Scope note: this locks the observable contract but cannot discriminate processRoot's internal +// ordering on constructible inputs - its root-tile scanline repair only does work when interior +// root-level tiles exist, which needs a level set thousands of voxels across, and the pre-fix +// code was accidentally host-synchronous for other topologies (a blocking legacy-stream memcpy). +TEST(TestNanoVDBCUDA, NonBlockingStreamSignedFloodFill) +{ + using BufferT = nanovdb::cuda::DeviceBuffer; + auto runOn = [](cudaStream_t stream, bool occupyDefault) { + auto hdl = nanovdb::tools::createLevelSetSphere(100); + auto* grid = hdl.grid(); + auto acc = grid->getAccessor(); + using OpT = nanovdb::SetVoxel; + acc.set(nanovdb::Coord(103,0,0), -1.0f);// flip sign and value of an inactive voxel + acc.set(nanovdb::Coord( 97,0,0), 1.0f);// (the corruption CudaSignedFloodFill uses) + hdl.deviceUpload(0, stream, true); + if (occupyDefault) {// keep the default stream busy for ~20 ms + int clockKHz = 0, dev = 0; + cudaCheck(cudaGetDevice(&dev)); + cudaCheck(cudaDeviceGetAttribute(&clockKHz, cudaDevAttrClockRate, dev)); + streamBusyWaitKernel<<<1, 1, 0, 0>>>(static_cast(clockKHz) * 20ull); + cudaCheck(cudaGetLastError()); + } + nanovdb::tools::cuda::signedFloodFill(hdl.deviceGrid(), false, stream); + hdl.deviceDownload(0, stream, true); + auto* out = hdl.grid(); + auto outAcc = out->getAccessor(); + EXPECT_EQ( 3.0f, outAcc(103,0,0)); + EXPECT_EQ( 0.0f, outAcc(100,0,0)); + EXPECT_EQ(-3.0f, outAcc( 97,0,0)); + return out->mChecksum.full(); + }; + const uint64_t reference = runOn(static_cast(0), false); + cudaCheck(cudaDeviceSynchronize()); + cudaStream_t nb = nullptr; + cudaCheck(cudaStreamCreateWithFlags(&nb, cudaStreamNonBlocking)); + const uint64_t candidate = runOn(nb, /*occupyDefault=*/true); + cudaCheck(cudaStreamSynchronize(nb)); + cudaCheck(cudaDeviceSynchronize());// drain the default-stream busy-wait so it can't leak into later tests + cudaCheck(cudaStreamDestroy(nb)); + EXPECT_EQ(reference, candidate); +}// NonBlockingStreamSignedFloodFill + // Regression test: the grid-name copy must not over-read the source string. // It previously copied a fixed MaxNameSize bytes from the (possibly shorter) // std::string, over-reading host heap into the grid; now it copies only @@ -3752,23 +3801,33 @@ __global__ void deviceBufferCountKernel(const unsigned char *p, size_t n, unsign if (p[i] != v) atomicAdd(bad, 1ull); } -// Regression test: a DeviceBuffer can be used on more than one stream, so ordering its device -// free on any single stream is not sufficient - it has to be ordered after every stream that -// touched the buffer. A device-only buffer is used here on purpose: it owns no pinned host -// memory, so clear()'s cudaFreeHost (which implicitly synchronizes) cannot mask the problem. +// Shared body of the two free-ordering regression tests below. A device-only buffer is used on +// purpose: it owns no pinned host memory, so clear()'s cudaFreeHost (which implicitly +// synchronizes) cannot mask a premature free. A second stream 'user' is parked behind a +// busy-wait with a write to the buffer's raw pointer queued behind it, then the buffer is +// destroyed. If the free is not ordered after 'user', the allocator recycles the block into the +// next allocation and the late write corrupts it. // -// 'user' is parked behind a spin kernel with a write to the buffer queued behind it, then the -// buffer is destroyed. If the free is not ordered after 'user', the allocator recycles the block -// into the next allocation and the late write corrupts it. -TEST(TestNanoVDBCUDA, DeviceBufferMultiStreamFreeOrdering) +// The two scenarios cover the two distinct ways the free can be ordered: +// - blocking 'user', write NOT registered: the free, issued on the legacy default stream, is +// implicitly ordered after all blocking streams. Freeing on any other single stream (e.g. +// the buffer's most recently used one) would drop that and fail here. +// - non-blocking 'user', write registered via recordUse: no implicit ordering exists for +// non-blocking streams, so only the recorded event orders the free. Freeing on the default +// stream without the event fails here. +static void testDeviceBufferFreeOrdering(bool nonBlockingUser, bool registerUse) { const size_t N = size_t(64) << 20;// large enough that a recycled block is the same address const unsigned char LATE = 0xAA, VICTIM = 0x55; const unsigned long long CYCLES = 400000000ull;// parks 'user' for O(100 ms) cudaStream_t user = nullptr, other = nullptr; - cudaCheck(cudaStreamCreate(&user));// blocking streams, i.e. the ones the legacy - cudaCheck(cudaStreamCreate(&other));// default stream implicitly synchronizes with + if (nonBlockingUser) { + cudaCheck(cudaStreamCreateWithFlags(&user, cudaStreamNonBlocking)); + } else { + cudaCheck(cudaStreamCreate(&user)); + } + cudaCheck(cudaStreamCreate(&other)); unsigned long long *bad = nullptr; cudaCheck(cudaMallocManaged(&bad, sizeof(*bad))); @@ -3789,6 +3848,7 @@ TEST(TestNanoVDBCUDA, DeviceBufferMultiStreamFreeOrdering) ASSERT_TRUE(devPtr); streamBusyWaitKernel<<<1,1,0,user>>>(CYCLES);// park 'user' deviceBufferFillKernel<<<1024,256,0,user>>>((unsigned char*)devPtr, N, LATE); + if (registerUse) buf.recordUse(0, user);// raw-pointer use: tell the buffer about it }// destroyed here; the free must be ordered after 'user' as well as 'other' unsigned char *victim = nullptr; @@ -3818,8 +3878,18 @@ TEST(TestNanoVDBCUDA, DeviceBufferMultiStreamFreeOrdering) EXPECT_EQ(0u, clobbered) << "device memory was freed while another stream still had work in " "flight (block recycled: " << recycled << ", work still pending when it was reused: " << stillPending << ")"; +}// testDeviceBufferFreeOrdering + +TEST(TestNanoVDBCUDA, DeviceBufferMultiStreamFreeOrdering) +{ + testDeviceBufferFreeOrdering(/*nonBlockingUser=*/false, /*registerUse=*/false); }// DeviceBufferMultiStreamFreeOrdering +TEST(TestNanoVDBCUDA, DeviceBufferNonBlockingFreeOrdering) +{ + testDeviceBufferFreeOrdering(/*nonBlockingUser=*/true, /*registerUse=*/true); +}// DeviceBufferNonBlockingFreeOrdering + TEST(TestNanoVDBCUDA, RefineCoarsen_ValueOnIndex) { using BuildT = nanovdb::ValueOnIndex;