NanoVDB: inject a resource into TopologyBuilder scratch, and align cuda::Buffer with cuda::buffer (CUDA) (#2268)

* NanoVDB: align cuda::Buffer member names with cuda::buffer

Name the free operation destroy, as cuda::buffer does, and add the
stream-taking overload it also provides. clear stays but only as a
transitional delegate, marked as such: it exists because
GridHandle::reset still calls it, and goes away with the legacy dual
buffers that cuda::Buffer replaces.

Rename setStream to set_stream. Member names cannot be aliased, so
matching the standard spelling is the whole reason the resource concept
kept allocate_async; the same argument applies here. Note in passing that
ours deliberately does not synchronize, matching cuda::buffer's
set_stream_unsynchronized rather than its set_stream, whose own
documentation and implementation disagree (NVIDIA/cccl#10649).

Add swap. The generic std::swap already does the right thing through the
move operations, but both std::vector and rmm::device_buffer provide one.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>

* NanoVDB: allocate TopologyBuilder scratch from an injected resource

Eleven of the builder's buffers are device-only -- nothing reads them on
the host -- yet they were cuda::DeviceBuffer, whose host pointer and
per-device array they never use. Move them to the single-space
cuda::Buffer and give the builder a resource parameter, so its scratch is
injectable like PointsToGrid's already is, and freed by scope rather than
by hand. mProcessedRoot and mData are read on the host and stay dual.

This is not a speedup: DeviceBuffer::init allocates host or device memory
but never both, and these buffers always passed a real device id, so no
cudaMallocHost was on this path to begin with. Measured dilate on 1k/20k/
200k points, and the difference is within run-to-run noise.

Hoist the Data struct out of the class. It does not depend on the
resource, and leaving it nested would give every ResourceT its own
incompatible type for the device functors to name.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>

* NanoVDB: note cuda::Buffer and cuda::BufferView in pendingchanges

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>

* NanoVDB: free TopologyBuilder scratch on the caller's stream

Each release site sits in a function that receives the stream, so pass
it to destroy rather than relying on the retained stream matching --
they are the same on every current path, but the explicit form does not
depend on that staying true. Also drop the cudaGetDevice calls whose
result the Buffer conversion left unused.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>

* NanoVDB: set_stream matches cuda::buffer's contract, not a divergence

cuda::buffer's set_stream is deliberately non-synchronizing -- its
documented synchronization is a stale note left behind when the
synchronization was removed -- so our non-synchronizing set_stream
matches it in both name and contract, and the comment claiming a
divergence was wrong.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>

* NanoVDB: mark cuda::Buffer::clear deprecated in the documentation

Documentation-level only: the [[deprecated]] attribute would warn from
GridHandle::reset and NodeManager::reset, template members in our own
headers that must keep calling clear() until every buffer type provides
destroy() -- an unactionable diagnostic for callers of reset(), and a
build break under -Werror. The attribute lands when those callers
migrate.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>

* NanoVDB: exercise the concept-required members of the test resources

The trait checks detect allocate/deallocate through unevaluated contexts,
which never odr-use them, so nvcc warned that every synchronous pair and
stub member was declared but never referenced (#177-D). The attribute
route is closed -- nvcc's front end ignores [[maybe_unused]] for this
diagnostic -- so reference them the honest way: a test that verifies the
synchronous halves of the counting and stream-recording doubles behave
like their stream-ordered halves, and one that pins the trait probes'
stub behavior. The TU now builds with no warnings at all.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>

---------

Signed-off-by: Mark Harris <mharris@nvidia.com>
Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>

Author: harrism

nanovdb/nanovdb/cuda/Buffer.h

@@ -50,7 +50,7 @@ struct StreamHolder<true> { cudaStream_t mStream = 0; };
 ///         is_resource). When @c R provides both interfaces the stream-ordered
 ///         one is used.
 /// @details With a stream-ordered resource the Buffer retains the stream of
-///          the most recent allocation (or the one supplied via setStream)
+///          the most recent allocation (or the one supplied via set_stream)
 ///          and orders its deallocation on that stream. Buffer is move-only.
 template<typename T, typename R = DeviceResource>
 class Buffer : private detail::StreamHolder<is_async_resource<R>::value>
@@ -133,7 +133,7 @@ class Buffer : private detail::StreamHolder<is_async_resource<R>::value>
     Buffer& operator=(Buffer&& other) noexcept
     {
         if (this != &other) {
-            this->clear();
+            this->destroy();
             static_cast<detail::StreamHolder<IsAsync>&>(*this) = other;
             mResource = std::move(other.mResource);
             mData = other.mData;
@@ -168,7 +168,7 @@ class Buffer : private detail::StreamHolder<is_async_resource<R>::value>
 
     /// @brief D-tor. A stream-ordered resource frees on the retained stream;
     ///        a synchronous resource frees immediately.
-    ~Buffer() { this->clear(); }
+    ~Buffer() { this->destroy(); }
 
     /// @brief Returns the retained stream, i.e. the stream the buffer's memory
     ///        will be freed on.
@@ -179,9 +179,11 @@ class Buffer : private detail::StreamHolder<is_async_resource<R>::value>
     ///        deallocation (and destruction) is ordered on @c stream instead.
     /// @param stream cuda stream subsequent deallocation is ordered on
     /// @warning The caller is responsible for ordering @c stream after any
-    ///          in-flight work that uses the buffer's memory.
+    ///          in-flight work that uses the buffer's memory. This deliberately
+    ///          does not synchronize, matching cuda::buffer's set_stream, which
+    ///          avoids implicit synchronization in fundamental primitives.
     template<typename S = R, std::enable_if_t<is_async_resource<S>::value, int> = 0>
-    void setStream(cudaStream_t stream) { this->mStream = stream; }
+    void set_stream(cudaStream_t stream) { this->mStream = stream; }
 
     /// @brief Resizes the buffer to @c count elements, preserving the leading
     ///        min(old, new) elements. Every operation — the new allocation, the
@@ -219,7 +221,7 @@ class Buffer : private detail::StreamHolder<is_async_resource<R>::value>
             mSize = count;
         }
         else {
-            this->mStream = stream; // no reallocation: setStream semantics
+            this->mStream = stream; // no reallocation: set_stream semantics
         }
     }
 
@@ -263,13 +265,50 @@ class Buffer : private detail::StreamHolder<is_async_resource<R>::value>
     bool empty() const { return mSize == 0; }
 
     /// @brief Frees the buffer memory (if any) and resets to the empty state.
-    void clear()
+    ///        A stream-ordered resource frees on the retained stream.
+    /// @note Spelled destroy to match cuda::buffer. This is the name to use.
+    void destroy()
     {
         this->deallocate(mData, mSize);
         mData = nullptr;
         mSize = 0;
     }
 
+    /// @brief Frees the buffer memory (if any) and resets to the empty state.
+    /// @deprecated Use destroy(). Documentation-level only for now: the
+    ///             [[deprecated]] attribute would fire from GridHandle::reset
+    ///             and NodeManager::reset, which must keep calling clear()
+    ///             until every buffer type provides destroy().
+    /// @note Transitional, and not the name to use: it exists only because
+    ///       GridHandle::reset still calls clear() on its buffer. It goes away
+    ///       when the legacy dual buffers do and GridHandle moves to destroy().
+    void clear() { this->destroy(); }
+
+    /// @brief Frees the buffer memory (if any) on @c stream and resets to the
+    ///        empty state. @c stream becomes the retained stream.
+    /// @param stream cuda stream the deallocation is ordered on
+    /// @warning The caller is responsible for ordering @c stream after any
+    ///          in-flight work that uses the buffer's memory.
+    template<typename S = R, std::enable_if_t<is_async_resource<S>::value, int> = 0>
+    void destroy(cudaStream_t stream)
+    {
+        this->mStream = stream;
+        this->destroy();
+    }
+
+    /// @brief Exchanges the contents of this buffer with @c other. Neither
+    ///        buffer allocates, frees, or copies element data.
+    /// @param other buffer to exchange contents with
+    void swap(Buffer& other) noexcept
+    {
+        auto& lhs = static_cast<detail::StreamHolder<IsAsync>&>(*this);
+        auto& rhs = static_cast<detail::StreamHolder<IsAsync>&>(other);
+        std::swap(lhs, rhs);
+        std::swap(mResource, other.mResource);
+        std::swap(mData, other.mData);
+        std::swap(mSize, other.mSize);
+    }
+
 private:
     /// @brief Returns @c count * sizeof(T), throwing std::runtime_error if the
     ///        byte size would overflow size_t instead of silently wrapping into

nanovdb/nanovdb/tools/cuda/CoarsenGrid.cuh

@@ -210,7 +210,7 @@ void CoarsenGrid<BuildT>::coarsenInternalNodes()
     if (auto srcLeafCount = mSrcTreeData.mNodeCount[0]) { // Unless it's an empty grid
         util::cuda::lambdaKernel<<<numBlocks(srcLeafCount), mNumThreads, 0, mStream>>>(
             srcLeafCount, util::morphology::cuda::CoarsenInternalNodesFunctor<BuildT>(),
-            mDeviceSrcGrid, mBuilder.deviceProcessedRoot(), mBuilder.mUpperMasks.deviceData(), mBuilder.mLowerMasks.deviceData() );
+            mDeviceSrcGrid, mBuilder.deviceProcessedRoot(), mBuilder.mUpperMasks.data(), mBuilder.mLowerMasks.data() );
     }
 }// CoarsenGrid<BuildT>::coarsenInternalNodes
 

nanovdb/nanovdb/tools/cuda/PruneGrid.cuh

@@ -213,7 +213,7 @@ void PruneGrid<BuildT>::pruneInternalNodes()
     if (auto srcLeafCount = mSrcTreeData.mNodeCount[0]) { // Unless it's an empty grid
         util::cuda::lambdaKernel<<<numBlocks(srcLeafCount), mNumThreads, 0, mStream>>>(
             srcLeafCount, util::morphology::cuda::PruneInternalNodesFunctor<BuildT>(),
-            mDeviceSrcGrid, mBuilder.deviceProcessedRoot(), mDeviceSrcLeafMask, mBuilder.mUpperMasks.deviceData(), mBuilder.mLowerMasks.deviceData() );
+            mDeviceSrcGrid, mBuilder.deviceProcessedRoot(), mDeviceSrcLeafMask, mBuilder.mUpperMasks.data(), mBuilder.mLowerMasks.data() );
     }
 }// PruneGrid<BuildT>::pruneInternalNodes
 

nanovdb/nanovdb/tools/cuda/RefineGrid.cuh

@@ -225,7 +225,7 @@ void RefineGrid<BuildT>::refineInternalNodes()
     if (auto srcLeafCount = mSrcTreeData.mNodeCount[0]) { // Unless it's an empty grid
         util::cuda::lambdaKernel<<<numBlocks(srcLeafCount), mNumThreads, 0, mStream>>>(
             srcLeafCount, util::morphology::cuda::RefineInternalNodesFunctor<BuildT>(),
-            mDeviceSrcGrid, mBuilder.deviceProcessedRoot(), mBuilder.mUpperMasks.deviceData(), mBuilder.mLowerMasks.deviceData() );
+            mDeviceSrcGrid, mBuilder.deviceProcessedRoot(), mBuilder.mUpperMasks.data(), mBuilder.mLowerMasks.data() );
     }
 }// RefineGrid<BuildT>::refineInternalNodes