NanoVDB: migrate off the dual-space buffers and deprecate the DeviceBuffer name (CUDA)

nanovdb/nanovdb/CMakeLists.txt

@@ -199,12 +199,16 @@ set(NANOVDB_INCLUDE_FILES
 
 # NanoVDB cuda header files
 set(NANOVDB_INCLUDE_CUDA_FILES
+  cuda/Buffer.h
   cuda/DeviceBuffer.h
   cuda/DeviceMesh.h
   cuda/DeviceResource.h
   cuda/DeviceStreamMap.h
   cuda/GridHandle.cuh
+  cuda/HandleStorage.h
+  cuda/ManagedResource.h
   cuda/NodeManager.cuh
+  cuda/PinnedResource.h
   cuda/TempPool.h
   cuda/UnifiedBuffer.h
 )

nanovdb/nanovdb/GridHandle.h

@@ -31,6 +31,13 @@ namespace nanovdb {
 
 struct GridHandleMetaData {uint64_t offset, size; GridType gridType;};
 
+namespace cuda { namespace detail {
+// Defined in nanovdb/cuda/HandleStorage.h: the one gateway to constructing a
+// handle from a buffer plus already-validated metadata (handle-to-handle
+// transfers), so the trust boundary stays visible in a single place.
+struct HandleFactory;
+}}// namespace cuda::detail
+
 namespace detail {
 
 /// @brief Allocates @c bytes of host-readable storage for a GridHandle:
@@ -88,6 +95,8 @@ class GridHandle
         : mMetaData(std::move(meta))
         , mBuffer(std::move(buffer)) {}
 
+    friend struct cuda::detail::HandleFactory;
+
 public:
     using BufferType = BufferT;
 
@@ -171,9 +180,9 @@ class GridHandle
     /// @brief Returns a pointer to the host data; not available for a
     ///        single-space device buffer, which has no host-readable bytes.
     /// @warning Note that the return pointer can be NULL if the GridHandle was not initialized
-    template<typename U = BufferT, typename util::disable_if<BufferHasDeviceSingle<U>::value, int>::type = 0>
+    template<typename U = BufferT, typename util::disable_if<BufferIsDeviceOnly<U>::value, int>::type = 0>
     void* data() { return mBuffer.data(); }
-    template<typename U = BufferT, typename util::disable_if<BufferHasDeviceSingle<U>::value, int>::type = 0>
+    template<typename U = BufferT, typename util::disable_if<BufferIsDeviceOnly<U>::value, int>::type = 0>
     const void* data() const { return mBuffer.data(); }
     //@}
 
@@ -220,15 +229,15 @@ class GridHandle
     /// @param n Index of the (host) grid pointer to be returned
     /// @warning Note that the return pointer can be NULL if the GridHandle no host grid, @a n is invalid
     ///          or if the template parameter does not match the specified grid!
-    template<typename ValueT, typename U = BufferT, typename util::disable_if<BufferHasDeviceSingle<U>::value, int>::type = 0>
+    template<typename ValueT, typename U = BufferT, typename util::disable_if<BufferIsDeviceOnly<U>::value, int>::type = 0>
     const NanoGrid<ValueT>* grid(uint32_t n = 0) const;
 
     /// @brief Returns a host pointer to the @a n'th  NanoVDB grid encoded in this GridHandle.
     /// @tparam ValueT Value type of the grid point to be returned
     /// @param n Index of the (host) grid pointer to be returned
     /// @warning Note that the return pointer can be NULL if the GridHandle no host grid, @a n is invalid
     ///          or if the template parameter does not match the specified grid!
-    template<typename ValueT, typename U = BufferT, typename util::disable_if<BufferHasDeviceSingle<U>::value, int>::type = 0>
+    template<typename ValueT, typename U = BufferT, typename util::disable_if<BufferIsDeviceOnly<U>::value, int>::type = 0>
     NanoGrid<ValueT>* grid(uint32_t n = 0) {return const_cast<NanoGrid<ValueT>*>(static_cast<const GridHandle*>(this)->template grid<ValueT>(n));}
 
     /// @brief Return a const pointer to the @a n'th grid encoded in this GridHandle on the device, e.g. GPU
@@ -322,20 +331,20 @@ class GridHandle
     /// @brief Access to the GridData of the n'th grid in the current handle
     /// @param n zero-based ID of the grid
     /// @return Const pointer to the n'th GridData in the current handle
-    template<typename U = BufferT, typename util::disable_if<BufferHasDeviceSingle<U>::value, int>::type = 0>
+    template<typename U = BufferT, typename util::disable_if<BufferIsDeviceOnly<U>::value, int>::type = 0>
     const GridData* gridData(uint32_t n = 0) const;
 
     /// @brief Returns a const point to the @a n'th grid meta data
     /// @param n zero-based ID of the grid
     /// @warning Note that the return pointer can be NULL if the GridHandle was not initialized
-    template<typename U = BufferT, typename util::disable_if<BufferHasDeviceSingle<U>::value, int>::type = 0>
+    template<typename U = BufferT, typename util::disable_if<BufferIsDeviceOnly<U>::value, int>::type = 0>
     const GridMetaData* gridMetaData(uint32_t n = 0) const;
 
     /// @brief Write a specific grid in this buffer to an output stream
     /// @param os  output stream that the buffer will be written to
     /// @param n zero-based index of the grid to be written to stream
     void write(std::ostream& os, uint32_t n) const {
-        static_assert(!BufferHasDeviceSingle<BufferT>::value,
+        static_assert(!(BufferIsDeviceOnly<BufferT>::value),
                       "GridHandle::write requires host-accessible grids: cuda::copyTo a host-readable handle first");
         if (const GridData* data = this->gridData(n)) {
             os.write((const char*)data, data->mGridSize);
@@ -347,7 +356,7 @@ class GridHandle
     /// @brief Write the entire grid buffer to an output stream
     /// @param os output stream that the buffer will be written to
     void write(std::ostream& os) const {
-        static_assert(!BufferHasDeviceSingle<BufferT>::value,
+        static_assert(!(BufferIsDeviceOnly<BufferT>::value),
                       "GridHandle::write requires host-accessible grids: cuda::copyTo a host-readable handle first");
 
         for (uint32_t n=0; n<this->gridCount(); ++n) this->write(os, n);
@@ -427,7 +436,7 @@ class GridHandle
 // --------------------------> Implementation of private methods in GridHandle <------------------------------------
 
 template<typename BufferT>
-template<typename U, typename util::disable_if<BufferHasDeviceSingle<U>::value, int>::type>
+template<typename U, typename util::disable_if<BufferIsDeviceOnly<U>::value, int>::type>
 inline const GridData* GridHandle<BufferT>::gridData(uint32_t n) const
 {
     const void *data = this->data();
@@ -436,7 +445,7 @@ inline const GridData* GridHandle<BufferT>::gridData(uint32_t n) const
 }// const GridData* GridHandle<BufferT>::gridData(uint32_t n) const
 
 template<typename BufferT>
-template<typename U, typename util::disable_if<BufferHasDeviceSingle<U>::value, int>::type>
+template<typename U, typename util::disable_if<BufferIsDeviceOnly<U>::value, int>::type>
 inline const GridMetaData* GridHandle<BufferT>::gridMetaData(uint32_t n) const
 {
     const auto *data = this->data();
@@ -474,7 +483,8 @@ inline GridHandle<OtherBufferT> GridHandle<BufferT>::copy(const OtherBufferT& ot
 {
     static_assert(!(BufferHasDeviceSingle<BufferT>::value || BufferHasDeviceSingle<OtherBufferT>::value),
                   "GridHandle::copy(pool) cannot honor a pool argument for a single-space device buffer, "
-                  "whose copy allocates through the source buffer's resource: use the no-argument copy()");
+                  "whose copy allocates through the source buffer's resource: use the no-argument copy() "
+                  "for a same-space deep copy, or cuda::copyTo (cuda/HandleStorage.h) to cross address spaces");
     if (mBuffer.size() == 0) return GridHandle<OtherBufferT>();// return an empty handle
     auto buffer = detail::createHostStorage<OtherBufferT>(mBuffer.size(), other);
     std::memcpy(buffer.data(), mBuffer.data(), mBuffer.size());// deep copy of buffer
@@ -488,19 +498,22 @@ inline GridHandle<OtherBufferT> GridHandle<BufferT>::copy() const
     if constexpr (BufferHasDeviceSingle<BufferT>::value || BufferHasDeviceSingle<OtherBufferT>::value) {
         static_assert(util::is_same<OtherBufferT, BufferT>::value && BufferHasDeviceSingle<BufferT>::value,
                       "GridHandle::copy is same-space only: a single-space device handle copies to its own "
-                      "buffer type; use cuda::copyTo (cuda/GridHandle.cuh) to move grids across address spaces");
+                      "buffer type; use cuda::copyTo (cuda/HandleStorage.h) to move grids across address spaces");
         // Device-to-device deep copy; for a stream-ordered resource it is
         // ordered on the source's retained stream, so synchronize that stream
         // before reading the result. Metadata is host-resident, so the copy
         // adopts it directly with no device re-parse.
         return GridHandle(mBuffer.copy(), mMetaData);
     } else {
+        static_assert(BufferIsDefaultConstructible<OtherBufferT>::value,
+                      "GridHandle::copy() without arguments default-constructs the target pool buffer: "
+                      "pass a prototype to copy(other) for a buffer over a non-default-constructible resource");
         return this->copy(OtherBufferT());
     }
 }// GridHandle<OtherBufferT> GridHandle<BufferT>::copy() const
 
 template<typename BufferT>
-template<typename ValueT, typename U, typename util::disable_if<BufferHasDeviceSingle<U>::value, int>::type>
+template<typename ValueT, typename U, typename util::disable_if<BufferIsDeviceOnly<U>::value, int>::type>
 inline const NanoGrid<ValueT>* GridHandle<BufferT>::grid(uint32_t n) const
 {
     return this->template gridAt<ValueT>(mBuffer.data(), n);
@@ -517,7 +530,7 @@ GridHandle<BufferT>::deviceGrid(uint32_t n) const
 template<typename BufferT>
 void GridHandle<BufferT>::read(std::istream& is, const BufferT& pool)
 {
-    static_assert(!BufferHasDeviceSingle<BufferT>::value,
+    static_assert(!(BufferIsDeviceOnly<BufferT>::value),
                   "GridHandle::read requires a host-accessible buffer: read into a host-readable handle, then cuda::copyTo");
     const std::streampos start = is.tellg();// remember where the raw buffer begins
     GridData data;
@@ -542,7 +555,7 @@ void GridHandle<BufferT>::read(std::istream& is, const BufferT& pool)
 template<typename BufferT>
 void GridHandle<BufferT>::read(std::istream& is, uint32_t n, const BufferT& pool)
 {
-    static_assert(!BufferHasDeviceSingle<BufferT>::value,
+    static_assert(!(BufferIsDeviceOnly<BufferT>::value),
                   "GridHandle::read requires a host-accessible buffer: read into a host-readable handle, then cuda::copyTo");
     GridData data;
     is.read((char*)&data, sizeof(GridData));
@@ -566,7 +579,7 @@ void GridHandle<BufferT>::read(std::istream& is, uint32_t n, const BufferT& pool
 template<typename BufferT>
 void GridHandle<BufferT>::read(std::istream& is, const std::string &gridName, const BufferT& pool)
 {
-    static_assert(!BufferHasDeviceSingle<BufferT>::value,
+    static_assert(!(BufferIsDeviceOnly<BufferT>::value),
                   "GridHandle::read requires a host-accessible buffer: read into a host-readable handle, then cuda::copyTo");
     static const std::streamsize byteSize = sizeof(GridData);
     GridData data;
@@ -600,7 +613,7 @@ template<typename BufferT, template <class, class...> class VectorT = std::vecto
 inline VectorT<GridHandle<BufferT>>
 splitGrids(const GridHandle<BufferT> &handle, const BufferT* other = nullptr)
 {
-    static_assert(!BufferHasDeviceSingle<BufferT>::value,
+    static_assert(!(BufferIsDeviceOnly<BufferT>::value),
                   "splitGrids requires a buffer type providing create(): cuda::copyTo a HostBuffer handle first");
     static_assert(!BufferHasHostSingle<BufferT>::value,
                   "splitGrids requires a buffer type providing create(): copy the handle to a HostBuffer first");
@@ -634,7 +647,7 @@ template<typename BufferT>
 inline GridHandle<BufferT>
 mergeGrids(const GridHandle<BufferT>* const* handles, size_t count, const BufferT* pool = nullptr)
 {
-    static_assert(!BufferHasDeviceSingle<BufferT>::value,
+    static_assert(!(BufferIsDeviceOnly<BufferT>::value),
                   "mergeGrids requires a buffer type providing create(): cuda::copyTo HostBuffer handles first");
     static_assert(!BufferHasHostSingle<BufferT>::value,
                   "mergeGrids requires a buffer type providing create(): copy the handles to HostBuffer first");

nanovdb/nanovdb/HostBuffer.h

@@ -123,6 +123,14 @@ template<typename BufferT>
 struct BufferHasHostSingle<BufferT, std::void_t<decltype(BufferTraits<BufferT>::hasHostSingle)>>
 { static constexpr bool value = BufferTraits<BufferT>::hasHostSingle; };
 
+/// @brief A single-space buffer whose storage the host cannot read: the
+///        device-single family minus its host-accessible members (managed or
+///        pinned resources). This is the predicate that gates the handles'
+///        host accessors off.
+template<typename BufferT>
+struct BufferIsDeviceOnly
+{ static constexpr bool value = BufferHasDeviceSingle<BufferT>::value && !BufferHasHostSingle<BufferT>::value; };
+
 /// @brief Detects whether a buffer exposes a retained stream (a stream()
 ///        member), i.e. whether its resource is stream-ordered. Used to pick
 ///        the buffer's stream-taking constructor without naming CUDA types.
@@ -141,6 +149,15 @@ template<typename BufferT>
 struct BufferHasByteElements<BufferT, std::void_t<typename BufferT::ElementType>>
 { static constexpr bool value = sizeof(typename BufferT::ElementType) == 1; };
 
+/// @brief Detects whether a buffer type is default-constructible, so
+///        consumers can name that requirement in a static_assert instead of
+///        failing wherever the default construction happens to occur.
+template<typename BufferT, typename = void>
+struct BufferIsDefaultConstructible { static constexpr bool value = false; };
+template<typename BufferT>
+struct BufferIsDefaultConstructible<BufferT, std::void_t<decltype(BufferT())>>
+{ static constexpr bool value = true; };
+
 /// @brief Detects whether a buffer provides destroy(), the cuda::Buffer
 ///        spelling for releasing its storage. Handle reset() dispatches to it
 ///        when present and falls back to the legacy clear() otherwise.

nanovdb/nanovdb/NodeManager.h

@@ -60,7 +60,7 @@ class NodeManagerHandle
     GridType mGridType{GridType::Unknown};
     BufferT  mBuffer;
 
-    template<typename BuildT, typename U = BufferT, typename util::disable_if<BufferHasDeviceSingle<U>::value, int>::type = 0>
+    template<typename BuildT, typename U = BufferT, typename util::disable_if<BufferIsDeviceOnly<U>::value, int>::type = 0>
     const NodeManager<BuildT>* getMgr() const {
         return mGridType == toGridType<BuildT>() ? (const NodeManager<BuildT>*)mBuffer.data() : nullptr;
     }
@@ -117,9 +117,9 @@ class NodeManagerHandle
     /// @brief Returns a pointer to the host data; not available for a
     ///        single-space device buffer, which has no host-readable bytes.
     /// @warning Note that the return pointer can be NULL if the NodeManagerHandle was not initialized
-    template<typename U = BufferT, typename util::disable_if<BufferHasDeviceSingle<U>::value, int>::type = 0>
+    template<typename U = BufferT, typename util::disable_if<BufferIsDeviceOnly<U>::value, int>::type = 0>
     void* data() { return mBuffer.data(); }
-    template<typename U = BufferT, typename util::disable_if<BufferHasDeviceSingle<U>::value, int>::type = 0>
+    template<typename U = BufferT, typename util::disable_if<BufferIsDeviceOnly<U>::value, int>::type = 0>
     const void* data() const { return mBuffer.data(); }
     //@}
 

nanovdb/nanovdb/cuda/Buffer.h

@@ -52,6 +52,18 @@ struct StreamHolder<true> { cudaStream_t mStream = 0; };
 /// @details With a stream-ordered resource the Buffer retains the stream of
 ///          the most recent allocation (or the one supplied via set_stream)
 ///          and orders its deallocation on that stream. Buffer is move-only.
+/// @note Cross-stream ordering is the caller's, expressed with ordinary CUDA
+///       events -- the buffer deliberately tracks nothing. To hand a buffer's
+///       contents to work on another stream (a consumer library, a wrapped
+///       tensor), record after the last write and make the consumer wait:
+/// @code
+///     cudaEvent_t ready;
+///     cudaEventCreateWithFlags(&ready, cudaEventDisableTiming);
+///     cudaEventRecord(ready, producerStream);     // after the last write
+///     cudaStreamWaitEvent(consumerStream, ready); // before the first read
+/// @endcode
+///       and order the buffer's destruction (which frees on its retained
+///       stream) after all consumers the same way, or synchronize.
 template<typename T, typename R = DeviceResource>
 class Buffer : private detail::StreamHolder<is_async_resource<R>::value>
 {
@@ -452,11 +464,15 @@ struct BufferTraits<cuda::Buffer<T, R>>
     // Device-resident storage; the byte-addressed requirement is enforced by
     // the single-space GridHandle constructor, so trait queries stay
     // answerable for any element type.
-    static constexpr bool hasDeviceSingle = !cuda::is_host_accessible_resource<R>::value;
+    static constexpr bool hasDeviceSingle = !cuda::is_host_accessible_resource<R>::value
+                                            || cuda::is_device_accessible_resource<R>::value;
     // A buffer over a host-accessible resource (e.g. PinnedResource) is
     // host-readable single-space storage: GridHandle parses its metadata on
     // the host, exposes the host accessors, and allocates reads and copies
-    // through the buffer's resource.
+    // through the buffer's resource. A resource that is host- AND
+    // device-accessible (ManagedResource) sets both members: the handle
+    // parses metadata through the device (a host parse could race producer
+    // kernels) and exposes both accessor families.
     static constexpr bool hasHostSingle   = cuda::is_host_accessible_resource<R>::value;
 };