diff --git a/nanovdb/nanovdb/cuda/DeviceBuffer.h b/nanovdb/nanovdb/cuda/DeviceBuffer.h index 9ad1d25441..f542bfdf10 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 + 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 @@ -46,6 +47,40 @@ 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 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) { + 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; + } + } + } + public: using PtrT = std::shared_ptr; @@ -122,8 +157,10 @@ class DeviceBuffer , mGpuData(other.mGpuData) , mDeviceCount(other.mDeviceCount) , mManaged(other.mManaged) + , mEvents(other.mEvents) { other.mCpuData = other.mGpuData = nullptr; + other.mEvents = nullptr; other.mSize = other.mDeviceCount = other.mManaged = 0; } @@ -142,6 +179,8 @@ class DeviceBuffer } /// @brief Destructor frees memory on both the host and device + /// @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 @@ -226,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]; @@ -301,6 +365,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));} @@ -310,18 +378,22 @@ class DeviceBuffer inline DeviceBuffer& DeviceBuffer::operator=(DeviceBuffer&& other) noexcept { - if (mManaged) {// first free all the managed data buffers + 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)); - for (int i=0; ifreeDeviceBuffers(cudaStream_t{0}); } delete [] mGpuData; + delete [] mEvents; mSize = other.mSize; mCpuData = other.mCpuData; mGpuData = other.mGpuData; mDeviceCount = other.mDeviceCount; mManaged = other.mManaged; + mEvents = other.mEvents; other.mCpuData = nullptr; other.mGpuData = nullptr; + other.mEvents = nullptr; other.mSize = 0; other.mDeviceCount = 0; other.mManaged = 0; @@ -333,6 +405,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 + 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 @@ -340,6 +413,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"); + 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 @@ -354,7 +428,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->recordUse(device, stream); if (sync) cudaCheck(cudaStreamSynchronize(stream)); } // DeviceBuffer::deviceUpload @@ -374,7 +452,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)); + this->recordUse(device, stream); if (sync) cudaCheck(cudaStreamSynchronize(stream)); } // DeviceBuffer::deviceDownload @@ -387,13 +467,15 @@ inline void DeviceBuffer::deviceDownload(void* stream, bool sync) inline void DeviceBuffer::clear(cudaStream_t stream) { - if (mManaged) {// free all the managed data buffers + 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 [] mEvents; mCpuData = nullptr; mGpuData = nullptr; + mEvents = nullptr; mSize = 0; mDeviceCount = 0; mManaged = 0; diff --git a/nanovdb/nanovdb/tools/cuda/DilateGrid.cuh b/nanovdb/nanovdb/tools/cuda/DilateGrid.cuh index f13ea5df32..01f9da9b67 100644 --- a/nanovdb/nanovdb/tools/cuda/DilateGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/DilateGrid.cuh @@ -279,7 +279,7 @@ void DilateGrid::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..1591749b58 100644 --- a/nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh @@ -20,6 +20,7 @@ #include #include #include +#include namespace nanovdb { @@ -909,11 +910,14 @@ 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). 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()) { + // 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/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..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,6 +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 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)); diff --git a/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh b/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh index d143793dce..81709e39ec 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 #include #include @@ -884,12 +885,15 @@ 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. 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)); + // 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)); } } // MeshToGrid::processGridTreeRoot diff --git a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh index 65b05b116b..4adebafe33 100644 --- a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh @@ -22,6 +22,7 @@ #include #include #include +#include #include #include @@ -792,6 +793,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 +945,18 @@ 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()) { + // 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 @@ -995,9 +1009,6 @@ 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(); }// PointsToGrid::processUpperNodes //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- @@ -1039,9 +1050,6 @@ 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(); }// PointsToGrid::processLowerNodes //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- @@ -1134,10 +1142,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) 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; + 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..69b7aa4e00 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,12 @@ void processRoot(NanoTree *d_tree) using ChildT = RootChild; static const int dim = int(RootT::ChildNodeType::DIM); + // 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)); cudaCheck(cudaMemcpy(uBuffer.data(), d_tree, uBuffer.size(), cudaMemcpyDeviceToHost));// copy Tree and Root (minus tiles) @@ -108,11 +114,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 +214,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..215a8a60be 100644 --- a/nanovdb/nanovdb/unittest/TestNanoVDB.cu +++ b/nanovdb/nanovdb/unittest/TestNanoVDB.cu @@ -3607,6 +3607,289 @@ 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, dev = 0; + cudaCheck(cudaGetDevice(&dev)); + cudaCheck(cudaDeviceGetAttribute(&clockKHz, cudaDevAttrClockRate, dev)); + 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(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 + 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 +// 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)); + // 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 +// 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 + +// 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; + 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 + +__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); +} + +// 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. +// +// 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; + if (nonBlockingUser) { + cudaCheck(cudaStreamCreateWithFlags(&user, cudaStreamNonBlocking)); + } else { + cudaCheck(cudaStreamCreate(&user)); + } + cudaCheck(cudaStreamCreate(&other)); + 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); + 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; + 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 << ")"; +}// 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; diff --git a/pendingchanges/nanovdbcorrectness.txt b/pendingchanges/nanovdbcorrectness.txt new file mode 100644 index 0000000000..631144e67c --- /dev/null +++ b/pendingchanges/nanovdbcorrectness.txt @@ -0,0 +1,16 @@ +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 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.