diff --git a/nanovdb/nanovdb/CMakeLists.txt b/nanovdb/nanovdb/CMakeLists.txt index 8f4fa07c56..69973be99f 100644 --- a/nanovdb/nanovdb/CMakeLists.txt +++ b/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 ) diff --git a/nanovdb/nanovdb/GridHandle.h b/nanovdb/nanovdb/GridHandle.h index 22b7697106..53c2df9612 100644 --- a/nanovdb/nanovdb/GridHandle.h +++ b/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::value, int>::type = 0> + template::value, int>::type = 0> void* data() { return mBuffer.data(); } - template::value, int>::type = 0> + template::value, int>::type = 0> const void* data() const { return mBuffer.data(); } //@} @@ -220,7 +229,7 @@ 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::value, int>::type = 0> + template::value, int>::type = 0> const NanoGrid* grid(uint32_t n = 0) const; /// @brief Returns a host pointer to the @a n'th NanoVDB grid encoded in this GridHandle. @@ -228,7 +237,7 @@ 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::value, int>::type = 0> + template::value, int>::type = 0> NanoGrid* grid(uint32_t n = 0) {return const_cast*>(static_cast(this)->template grid(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::value, int>::type = 0> + template::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::value, int>::type = 0> + template::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::value, + static_assert(!(BufferIsDeviceOnly::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::value, + static_assert(!(BufferIsDeviceOnly::value), "GridHandle::write requires host-accessible grids: cuda::copyTo a host-readable handle first"); for (uint32_t n=0; ngridCount(); ++n) this->write(os, n); @@ -427,7 +436,7 @@ class GridHandle // --------------------------> Implementation of private methods in GridHandle <------------------------------------ template -template::value, int>::type> +template::value, int>::type> inline const GridData* GridHandle::gridData(uint32_t n) const { const void *data = this->data(); @@ -436,7 +445,7 @@ inline const GridData* GridHandle::gridData(uint32_t n) const }// const GridData* GridHandle::gridData(uint32_t n) const template -template::value, int>::type> +template::value, int>::type> inline const GridMetaData* GridHandle::gridMetaData(uint32_t n) const { const auto *data = this->data(); @@ -474,7 +483,8 @@ inline GridHandle GridHandle::copy(const OtherBufferT& ot { static_assert(!(BufferHasDeviceSingle::value || BufferHasDeviceSingle::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();// return an empty handle auto buffer = detail::createHostStorage(mBuffer.size(), other); std::memcpy(buffer.data(), mBuffer.data(), mBuffer.size());// deep copy of buffer @@ -488,19 +498,22 @@ inline GridHandle GridHandle::copy() const if constexpr (BufferHasDeviceSingle::value || BufferHasDeviceSingle::value) { static_assert(util::is_same::value && BufferHasDeviceSingle::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::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 GridHandle::copy() const template -template::value, int>::type> +template::value, int>::type> inline const NanoGrid* GridHandle::grid(uint32_t n) const { return this->template gridAt(mBuffer.data(), n); @@ -517,7 +530,7 @@ GridHandle::deviceGrid(uint32_t n) const template void GridHandle::read(std::istream& is, const BufferT& pool) { - static_assert(!BufferHasDeviceSingle::value, + static_assert(!(BufferIsDeviceOnly::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::read(std::istream& is, const BufferT& pool) template void GridHandle::read(std::istream& is, uint32_t n, const BufferT& pool) { - static_assert(!BufferHasDeviceSingle::value, + static_assert(!(BufferIsDeviceOnly::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::read(std::istream& is, uint32_t n, const BufferT& pool template void GridHandle::read(std::istream& is, const std::string &gridName, const BufferT& pool) { - static_assert(!BufferHasDeviceSingle::value, + static_assert(!(BufferIsDeviceOnly::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 class VectorT = std::vecto inline VectorT> splitGrids(const GridHandle &handle, const BufferT* other = nullptr) { - static_assert(!BufferHasDeviceSingle::value, + static_assert(!(BufferIsDeviceOnly::value), "splitGrids requires a buffer type providing create(): cuda::copyTo a HostBuffer handle first"); static_assert(!BufferHasHostSingle::value, "splitGrids requires a buffer type providing create(): copy the handle to a HostBuffer first"); @@ -634,7 +647,7 @@ template inline GridHandle mergeGrids(const GridHandle* const* handles, size_t count, const BufferT* pool = nullptr) { - static_assert(!BufferHasDeviceSingle::value, + static_assert(!(BufferIsDeviceOnly::value), "mergeGrids requires a buffer type providing create(): cuda::copyTo HostBuffer handles first"); static_assert(!BufferHasHostSingle::value, "mergeGrids requires a buffer type providing create(): copy the handles to HostBuffer first"); diff --git a/nanovdb/nanovdb/HostBuffer.h b/nanovdb/nanovdb/HostBuffer.h index 9abf01f366..cb51aa317d 100644 --- a/nanovdb/nanovdb/HostBuffer.h +++ b/nanovdb/nanovdb/HostBuffer.h @@ -123,6 +123,14 @@ template struct BufferHasHostSingle::hasHostSingle)>> { static constexpr bool value = BufferTraits::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 +struct BufferIsDeviceOnly +{ static constexpr bool value = BufferHasDeviceSingle::value && !BufferHasHostSingle::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 struct BufferHasByteElements> { 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 +struct BufferIsDefaultConstructible { static constexpr bool value = false; }; +template +struct BufferIsDefaultConstructible> +{ 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. diff --git a/nanovdb/nanovdb/NodeManager.h b/nanovdb/nanovdb/NodeManager.h index d8044a9265..2f41996257 100644 --- a/nanovdb/nanovdb/NodeManager.h +++ b/nanovdb/nanovdb/NodeManager.h @@ -60,7 +60,7 @@ class NodeManagerHandle GridType mGridType{GridType::Unknown}; BufferT mBuffer; - template::value, int>::type = 0> + template::value, int>::type = 0> const NodeManager* getMgr() const { return mGridType == toGridType() ? (const NodeManager*)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::value, int>::type = 0> + template::value, int>::type = 0> void* data() { return mBuffer.data(); } - template::value, int>::type = 0> + template::value, int>::type = 0> const void* data() const { return mBuffer.data(); } //@} diff --git a/nanovdb/nanovdb/cuda/Buffer.h b/nanovdb/nanovdb/cuda/Buffer.h index 30ae022bfa..9a07309a53 100644 --- a/nanovdb/nanovdb/cuda/Buffer.h +++ b/nanovdb/nanovdb/cuda/Buffer.h @@ -52,6 +52,18 @@ struct StreamHolder { 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 class Buffer : private detail::StreamHolder::value> { @@ -452,11 +464,15 @@ struct BufferTraits> // 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::value; + static constexpr bool hasDeviceSingle = !cuda::is_host_accessible_resource::value + || cuda::is_device_accessible_resource::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::value; }; diff --git a/nanovdb/nanovdb/cuda/DeviceBuffer.h b/nanovdb/nanovdb/cuda/DeviceBuffer.h index a39e810cff..450b5cb780 100644 --- a/nanovdb/nanovdb/cuda/DeviceBuffer.h +++ b/nanovdb/nanovdb/cuda/DeviceBuffer.h @@ -8,7 +8,7 @@ \date January 8, 2020 - \brief DeviceBuffer has one pinned host buffer and multiple device CUDA buffers + \brief DualDeviceBuffer has one pinned host buffer and multiple device CUDA buffers \note This file has no device-only kernel functions, which explains why it's a .h and not .cuh file. @@ -26,14 +26,24 @@ namespace nanovdb {// ========================================================== namespace cuda {// =================================================================== -// ----------------------------> DeviceBuffer <-------------------------------------- +// ----------------------------> DualDeviceBuffer <-------------------------------------- /// @brief Simple memory buffer using un-managed pinned host memory when compiled with NVCC. /// Obviously this class is making explicit used of CUDA so replace it with your own memory /// allocator if you are not using CUDA. /// @note While CUDA's pinned host memory allows for asynchronous memory copy between host and device /// it is significantly slower then cached (un-pinned) memory on the host. -class DeviceBuffer +/// @note This is the implementation behind the deprecated DeviceBuffer alias +/// below, renamed so the [[deprecated]] attribute reaches only code +/// that spells the public name: the GPU tools' signature defaults +/// reference this implementation, so default-using callers stay +/// warning-free until the defaults change at removal. Transitional -- +/// do not adopt this name; it is deleted together with the alias. The +/// header keeps its long-standing name and include path for the same +/// reason: renaming a header breaks existing includes outright, and +/// the old path is where external code will find the alias and its +/// migration message. +class DualDeviceBuffer { 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 @@ -55,7 +65,7 @@ class DeviceBuffer /// 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) + void freeDualDeviceBuffers(cudaStream_t stream) { int current = 0; cudaCheck(cudaGetDevice(¤t)); @@ -76,16 +86,16 @@ class DeviceBuffer public: - using PtrT = std::shared_ptr; + using PtrT = std::shared_ptr; /// @brief Default constructor of an empty buffer - DeviceBuffer() : mSize(0), mCpuData(nullptr), mGpuData(nullptr), mDeviceCount(0), mManaged(0){} + DualDeviceBuffer() : mSize(0), mCpuData(nullptr), mGpuData(nullptr), mDeviceCount(0), mManaged(0){} /// @brief Constructor with a specified device and size /// @param size byte size of buffer to be initialized /// @param device id of the device on which to initialize the buffer /// @param stream cuda stream - DeviceBuffer(uint64_t size, int device = cudaCpuDeviceId, cudaStream_t stream = 0) : DeviceBuffer() + DualDeviceBuffer(uint64_t size, int device = cudaCpuDeviceId, cudaStream_t stream = 0) : DualDeviceBuffer() { this->init(size, device, stream); } @@ -94,7 +104,7 @@ class DeviceBuffer /// @param size byte size of buffer to be initialized /// @param host If true buffer is initialized only on the host/CPU, else on the current device/GPU /// @param stream optional stream argument (defaults to stream NULL) - DeviceBuffer(uint64_t size, bool host, void* stream) : DeviceBuffer() + DualDeviceBuffer(uint64_t size, bool host, void* stream) : DualDeviceBuffer() { int device = cudaCpuDeviceId; if (!host) cudaCheck(cudaGetDevice(&device)); @@ -107,7 +117,7 @@ class DeviceBuffer /// @param gpuData device buffer, assumed to NOT be NULL; /// @note The device buffer, @c gpuData, will be associated /// with the current device ID given by cudaGetDevice - DeviceBuffer(uint64_t size, void* cpuData, void* gpuData) + DualDeviceBuffer(uint64_t size, void* cpuData, void* gpuData) : mSize(size) , mCpuData(cpuData) , mManaged(0) @@ -125,7 +135,7 @@ class DeviceBuffer /// @param size byte size of the two external buffers /// @param cpuData host buffer, assumed to NOT be NULL /// @param list list of device IDs and external device buffers, all assumed to not be NULL - DeviceBuffer(uint64_t size, void* cpuData, std::initializer_list> list) + DualDeviceBuffer(uint64_t size, void* cpuData, std::initializer_list> list) : mSize(size) , mCpuData(cpuData) , mManaged(0) @@ -141,10 +151,10 @@ class DeviceBuffer } /// @brief Disallow copy-construction - DeviceBuffer(const DeviceBuffer&) = delete; + DualDeviceBuffer(const DualDeviceBuffer&) = delete; /// @brief Move copy-constructor - DeviceBuffer(DeviceBuffer&& other) noexcept + DualDeviceBuffer(DualDeviceBuffer&& other) noexcept : mSize(other.mSize) , mCpuData(other.mCpuData) , mGpuData(other.mGpuData) @@ -161,8 +171,8 @@ class DeviceBuffer /// @param buffer host buffer from which to copy data /// @param device id of the device on which to initialize the buffer /// @param stream cuda stream - DeviceBuffer(const HostBuffer& buffer, int device = cudaCpuDeviceId, cudaStream_t stream = 0) - : DeviceBuffer(buffer.size(), device, stream) + DualDeviceBuffer(const HostBuffer& buffer, int device = cudaCpuDeviceId, cudaStream_t stream = 0) + : DualDeviceBuffer(buffer.size(), device, stream) { if (mCpuData) { cudaCheck(cudaMemcpy(mCpuData, buffer.data(), mSize, cudaMemcpyHostToHost)); @@ -174,7 +184,7 @@ 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(); }; + ~DualDeviceBuffer() { this->clear(); }; /// @brief Static factory method that return an instance of this buffer /// @param size byte size of buffer to be initialized @@ -182,51 +192,51 @@ class DeviceBuffer /// @param host If true buffer is initialized only on the host/CPU, else only on the device/GPU /// @param stream optional stream argument (defaults to stream NULL) /// @return An instance of this class using move semantics - static DeviceBuffer create(uint64_t size, const DeviceBuffer* dummy, bool host, void* stream){return DeviceBuffer(size, host, stream);} + static DualDeviceBuffer create(uint64_t size, const DualDeviceBuffer* dummy, bool host, void* stream){return DualDeviceBuffer(size, host, stream);} /// @brief Static factory method that returns an instance of this buffer /// @param size byte size of buffer to be initialized /// @param dummy this argument is currently ignored but required to match the API of the HostBuffer /// @param device id of the device on which to initialize the buffer /// @param stream cuda stream - static DeviceBuffer create(uint64_t size, const DeviceBuffer* dummy = nullptr, int device = cudaCpuDeviceId, cudaStream_t stream = 0){return DeviceBuffer(size, device, stream);} + static DualDeviceBuffer create(uint64_t size, const DualDeviceBuffer* dummy = nullptr, int device = cudaCpuDeviceId, cudaStream_t stream = 0){return DualDeviceBuffer(size, device, stream);} /// @brief Static factory method that returns an instance of this buffer that wraps externally managed memory /// @param size byte size of buffer specified by external memory /// @param cpuData pointer to externally managed host memory /// @param gpuData pointer to externally managed device memory /// @return An instance of this class using move semantics - static DeviceBuffer create(uint64_t size, void* cpuData, void* gpuData) {return DeviceBuffer(size, cpuData, gpuData);} + static DualDeviceBuffer create(uint64_t size, void* cpuData, void* gpuData) {return DualDeviceBuffer(size, cpuData, gpuData);} /// @brief Static factory method that returns an instance of this buffer that wraps externally managed host and device memory /// @param size byte size of buffer to be initialized /// @param cpuData pointer to externally managed host memory /// @param list list of device IDs and device memory pointers - static DeviceBuffer create(uint64_t size, void* cpuData, std::initializer_list> list) {return DeviceBuffer(size, cpuData, list);} + static DualDeviceBuffer create(uint64_t size, void* cpuData, std::initializer_list> list) {return DualDeviceBuffer(size, cpuData, list);} /// @brief Static factory method that returns an instance of this buffer constructed from a HostBuffer /// @param buffer host buffer from which to copy data /// @param device id of the device on which to initialize the buffer /// @param stream cuda stream - static DeviceBuffer create(const HostBuffer& buffer, int device = cudaCpuDeviceId, cudaStream_t stream = 0) {return DeviceBuffer(buffer, device, stream);} + static DualDeviceBuffer create(const HostBuffer& buffer, int device = cudaCpuDeviceId, cudaStream_t stream = 0) {return DualDeviceBuffer(buffer, device, stream);} /////////////////////////////////////////////////////////////////////// /// @{ - /// @brief Factory methods that create a shared pointer to an DeviceBuffer instance - static PtrT createPtr(uint64_t size, const DeviceBuffer* = nullptr, int device = cudaCpuDeviceId, cudaStream_t stream = 0) {return std::make_shared(size, device, stream);} - static PtrT createPtr(uint64_t size, void* cpuData, void* gpuData) {return std::make_shared(size, cpuData, gpuData);} - static PtrT createPtr(uint64_t size, void* cpuData, std::initializer_list> list) {return std::make_shared(size, cpuData, list);} - static PtrT createPtr(const HostBuffer& buffer, int device = cudaCpuDeviceId, cudaStream_t stream = 0) {return std::make_shared(buffer, device, stream);} + /// @brief Factory methods that create a shared pointer to an DualDeviceBuffer instance + static PtrT createPtr(uint64_t size, const DualDeviceBuffer* = nullptr, int device = cudaCpuDeviceId, cudaStream_t stream = 0) {return std::make_shared(size, device, stream);} + static PtrT createPtr(uint64_t size, void* cpuData, void* gpuData) {return std::make_shared(size, cpuData, gpuData);} + static PtrT createPtr(uint64_t size, void* cpuData, std::initializer_list> list) {return std::make_shared(size, cpuData, list);} + static PtrT createPtr(const HostBuffer& buffer, int device = cudaCpuDeviceId, cudaStream_t stream = 0) {return std::make_shared(buffer, device, stream);} /// @} /////////////////////////////////////////////////////////////////////// /// @brief Disallow copy assignment operation - DeviceBuffer& operator=(const DeviceBuffer&) = delete; + DualDeviceBuffer& operator=(const DualDeviceBuffer&) = delete; /// @brief Move copy assignment operation - DeviceBuffer& operator=(DeviceBuffer&& other) noexcept; + DualDeviceBuffer& operator=(DualDeviceBuffer&& other) noexcept; /////////////////////////////////////////////////////////////////////// @@ -391,16 +401,28 @@ class DeviceBuffer void clear(cudaStream_t stream = 0); void clear(void* stream){this->clear(cudaStream_t(stream));} -}; // DeviceBuffer class +}; // DualDeviceBuffer class + +/// @brief The dual-space device buffer under its long-standing public name. +/// @deprecated Grid storage is moving to the single-space cuda::Buffer: +/// build or read into a host handle and move it with +/// cuda::copyTo (see cuda/HandleStorage.h), or allocate the +/// result of a GPU tool directly in a cuda::Buffer. Transfers +/// adopt the source handle's already-validated metadata -- no +/// kernel runs -- so copyTo is callable from host-only +/// translation units directly (see the CUDA examples). The +/// dual buffer and this name are removed together after a +/// deprecation window. +using DeviceBuffer [[deprecated("grid storage is moving to cuda::Buffer: build into a host handle and use cuda::copyTo (cuda/HandleStorage.h, host-callable); see the CUDA examples")]] = DualDeviceBuffer; // --------------------------> Implementations below <------------------------------------ -inline DeviceBuffer& DeviceBuffer::operator=(DeviceBuffer&& other) noexcept +inline DualDeviceBuffer& DualDeviceBuffer::operator=(DualDeviceBuffer&& 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}); + this->freeDualDeviceBuffers(cudaStream_t{0}); } delete [] mGpuData; delete [] mEvents; @@ -419,7 +441,7 @@ inline DeviceBuffer& DeviceBuffer::operator=(DeviceBuffer&& other) noexcept return *this; } -inline void DeviceBuffer::init(uint64_t size, int device, cudaStream_t stream) +inline void DualDeviceBuffer::init(uint64_t size, int device, cudaStream_t stream) { if (size==0) return; cudaCheck(cudaGetDeviceCount(&mDeviceCount)); @@ -428,22 +450,22 @@ inline void DeviceBuffer::init(uint64_t size, int device, cudaStream_t stream) 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 - checkPtr(mCpuData, "cuda::DeviceBuffer::init: failed to allocate host buffer"); + checkPtr(mCpuData, "cuda::DualDeviceBuffer::init: failed to allocate host buffer"); } 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"); + checkPtr(mGpuData[device], "cuda::DualDeviceBuffer::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 -} // DeviceBuffer::init +} // DualDeviceBuffer::init -inline void DeviceBuffer::deviceUpload(int device, cudaStream_t stream, bool sync) +inline void DualDeviceBuffer::deviceUpload(int device, cudaStream_t stream, bool sync) { NANOVDB_ASSERT(device >= 0 && device < mDeviceCount);// should be device and not the host checkPtr(mCpuData, "uninitialized cpu source data"); if (mGpuData[device] == nullptr) { - if (mManaged==0) throw std::runtime_error("DeviceBuffer::deviceUpload called on externally managed memory that wasn\'t allocated."); + if (mManaged==0) throw std::runtime_error("DualDeviceBuffer::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! } checkPtr(mGpuData[device], "uninitialized gpu destination data"); @@ -453,21 +475,21 @@ inline void DeviceBuffer::deviceUpload(int device, cudaStream_t stream, bool syn cudaCheck(cudaMemcpyAsync(mGpuData[device], mCpuData, mSize, cudaMemcpyHostToDevice, stream)); this->recordUse(device, stream); if (sync) cudaCheck(cudaStreamSynchronize(stream)); -} // DeviceBuffer::deviceUpload +} // DualDeviceBuffer::deviceUpload -inline void DeviceBuffer::deviceUpload(cudaStream_t stream, bool sync) +inline void DualDeviceBuffer::deviceUpload(cudaStream_t stream, bool sync) { int device = 0; cudaGetDevice(&device); this->deviceUpload(device, stream, sync); -} // DeviceBuffer::deviceUpload +} // DualDeviceBuffer::deviceUpload -inline void DeviceBuffer::deviceDownload(int device, cudaStream_t stream, bool sync) +inline void DualDeviceBuffer::deviceDownload(int device, cudaStream_t stream, bool sync) { NANOVDB_ASSERT(device >= 0 && device < mDeviceCount); checkPtr(mGpuData[device], "uninitialized gpu source data");// no source data on the specified device if (mCpuData == nullptr) { - if (mManaged==0) throw std::runtime_error("DeviceBuffer::deviceDownload called on uninitialized cpu destination memory that is externally managed."); + if (mManaged==0) throw std::runtime_error("DualDeviceBuffer::deviceDownload called on uninitialized cpu destination memory that is externally managed."); 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"); @@ -475,20 +497,20 @@ inline void DeviceBuffer::deviceDownload(int device, cudaStream_t stream, bool s cudaCheck(cudaMemcpyAsync(mCpuData, mGpuData[device], mSize, cudaMemcpyDeviceToHost, stream)); this->recordUse(device, stream); if (sync) cudaCheck(cudaStreamSynchronize(stream)); -} // DeviceBuffer::deviceDownload +} // DualDeviceBuffer::deviceDownload -inline void DeviceBuffer::deviceDownload(void* stream, bool sync) +inline void DualDeviceBuffer::deviceDownload(void* stream, bool sync) { int device = 0; cudaCheck(cudaGetDevice(&device)); this->deviceDownload(device, cudaStream_t(stream), sync); -} // DeviceBuffer::deviceDownload +} // DualDeviceBuffer::deviceDownload -inline void DeviceBuffer::clear(cudaStream_t stream) +inline void DualDeviceBuffer::clear(cudaStream_t stream) { if (mManaged) {// free all the managed data buffers, ordered after every use of each cudaCheck(cudaFreeHost(mCpuData)); - this->freeDeviceBuffers(stream); + this->freeDualDeviceBuffers(stream); } delete [] mGpuData; delete [] mEvents; @@ -498,14 +520,14 @@ inline void DeviceBuffer::clear(cudaStream_t stream) mSize = 0; mDeviceCount = 0; mManaged = 0; -} // DeviceBuffer::clear +} // DualDeviceBuffer::clear }// namespace cuda -using CudaDeviceBuffer [[deprecated("Use nanovdb::cuda::DeviceBuffer instead")]] = cuda::DeviceBuffer; +using CudaDeviceBuffer [[deprecated("Use GridHandle> with cuda::copyTo instead")]] = cuda::DualDeviceBuffer; template<> -struct BufferTraits +struct BufferTraits { static constexpr bool hasDeviceDual = true; }; diff --git a/nanovdb/nanovdb/cuda/DeviceResource.h b/nanovdb/nanovdb/cuda/DeviceResource.h index b22ef44d90..d24a4c4d15 100644 --- a/nanovdb/nanovdb/cuda/DeviceResource.h +++ b/nanovdb/nanovdb/cuda/DeviceResource.h @@ -134,6 +134,18 @@ struct is_host_accessible_resource : std::false_type {}; template struct is_host_accessible_resource::type> : std::true_type {}; +/// @brief Companion detection: @c is_device_accessible_resource::value is +/// true iff @c R declares `static constexpr bool DEVICE_ACCESSIBLE = +/// true`, i.e. its allocations are also valid device addresses even +/// though they are host-accessible (e.g. ManagedResource). A handle +/// over such a resource exposes both accessor families. Purely +/// device-resident resources do not need the marker: not being +/// host-accessible already implies device residency. +template +struct is_device_accessible_resource : std::false_type {}; +template +struct is_device_accessible_resource::type> : std::true_type {}; + /// @brief Detection trait: @c is_resource::value is true iff @c R models /// the synchronous Resource concept, i.e. exposes /// allocate(size_t, size_t) and deallocate(void*, size_t, size_t). @@ -257,6 +269,7 @@ struct AsyncFromSync /// @brief The adapter is host-accessible iff the adapted resource is. static constexpr bool HOST_ACCESSIBLE = is_host_accessible_resource::value; + static constexpr bool DEVICE_ACCESSIBLE = is_device_accessible_resource::value; R resource; @@ -302,6 +315,7 @@ struct ResourceRef /// @brief A reference is host-accessible iff the referenced resource is. static constexpr bool HOST_ACCESSIBLE = is_host_accessible_resource::value; + static constexpr bool DEVICE_ACCESSIBLE = is_device_accessible_resource::value; /// @brief Constructs a ref borrowing @c resource. /// @param resource resource to allocate from; must outlive this ref diff --git a/nanovdb/nanovdb/cuda/GridHandle.cuh b/nanovdb/nanovdb/cuda/GridHandle.cuh index 1350d8c2ad..8e76fca642 100644 --- a/nanovdb/nanovdb/cuda/GridHandle.cuh +++ b/nanovdb/nanovdb/cuda/GridHandle.cuh @@ -18,6 +18,7 @@ #define NANOVDB_CUDA_GRIDHANDLE_CUH_HAS_BEEN_INCLUDED #include // for the resource-aware scratch buffers below +#include // cuda::copyTo and the storage helpers live there (host-includable) #include // required for instantiation of move c-tor of GridHandle #include // for cuda::updateChecksum #include @@ -148,39 +149,6 @@ inline ScratchT makeMetaScratch(const BufferT& buf, uint64_t count, cudaStream_t else { (void)stream; return ScratchT(buf.resource(), count, noInit); } } -/// @brief Allocates @c bytes of destination storage for a cross-space -/// transfer: single-space buffers allocate through @c proto's resource -/// (or a default-constructed resource without one), on @c stream when -/// the resource is stream-ordered; buffers providing create() go -/// through it. -template -inline DstBufferT makeTransferStorage(uint64_t bytes, cudaStream_t stream, const DstBufferT* proto) -{ - if constexpr (BufferHasDeviceSingle::value || BufferHasHostSingle::value) { - using ResourceT = typename DstBufferT::ResourceType; - if (!proto) { - // both branches of a plain conditional would instantiate the - // default-resource constructor, breaking non-default-constructible - // resources (e.g. ResourceRef) even for callers that pass a proto - if constexpr (std::is_default_constructible::value) { - if constexpr (is_async_resource::value) return DstBufferT(stream, bytes, noInit); - else return DstBufferT(bytes, noInit); - } else { - throw std::runtime_error("cuda::copyTo: a destination buffer over a non-default-constructible " - "resource requires a prototype buffer"); - } - } - if constexpr (is_async_resource::value) { - return DstBufferT(stream, proto->resource(), bytes, noInit); - } else { - (void)stream; - return DstBufferT(proto->resource(), bytes, noInit); - } - } else { - return DstBufferT::create(bytes, proto); - } -} - }// namespace detail template class VectorT = std::vector> @@ -243,79 +211,6 @@ mergeGridHandles(const VectorT> &handles, const BufferT* oth return GridHandle(std::move(buffer)); }// cuda::mergeGridHandles -/// @brief Deep-copies a grid handle into a different address space: the -/// explicit, stream-carrying transfer between single-space device -/// handles and host-readable handles (HostBuffer or a host-accessible -/// single-space buffer such as a pinned-resource cuda::Buffer). -/// @tparam DstBufferT destination buffer type (specify explicitly) -/// @param src the handle to copy; must not be dual-space (use -/// deviceUpload/deviceDownload on those) -/// @param stream stream the copy is issued on; a device destination buffer -/// with a stream-ordered resource retains it -/// @warning Passing a stream other than the source buffer's retained stream -/// makes the caller responsible for ordering: prior work on the -/// source (and the source's later destruction, which frees on its -/// own stream) must be ordered against @a stream by the caller, -/// e.g. with cudaStreamWaitEvent or a synchronization. The -/// stream-less overload below has no such requirement. -/// @param proto optional buffer whose resource (or pool, for buffers -/// providing create()) allocates the destination storage; without it -/// the destination resource is default-constructed -/// @return a handle of the destination buffer type with equal contents -/// @details The returned handle is immediately usable: a host-readable -/// destination synchronizes @c stream before returning, and a device -/// destination parses (and validates) its metadata on the -/// transferred bytes, which synchronizes internally. A pageable host -/// source or destination (HostBuffer) degrades the copy to -/// synchronous behavior; pinned single-space handles keep it -/// asynchronous. -template -inline GridHandle copyTo(const GridHandle& src, cudaStream_t stream, const DstBufferT* proto = nullptr) -{ - constexpr bool srcDev = BufferHasDeviceSingle::value; - constexpr bool dstDev = BufferHasDeviceSingle::value; - static_assert(!BufferTraits::hasDeviceDual && !BufferTraits::hasDeviceDual, - "cuda::copyTo does not support dual-space buffers: use deviceUpload/deviceDownload on the handle"); - static_assert(srcDev || dstDev, - "cuda::copyTo is for cross-space transfers involving a device buffer: use GridHandle::copy for host-to-host"); - const uint64_t bytes = src.bufferSize(); - if (bytes == 0u) { - if constexpr (std::is_default_constructible::value) { - return GridHandle(); - } else { - throw std::runtime_error("cuda::copyTo: an empty handle cannot be copied to a buffer type " - "that is not default-constructible"); - } - } - DstBufferT dst = detail::makeTransferStorage(bytes, stream, proto); - const void* srcPtr; - if constexpr (srcDev) srcPtr = src.deviceData(); - else srcPtr = src.data(); - constexpr cudaMemcpyKind kind = srcDev ? (dstDev ? cudaMemcpyDeviceToDevice : cudaMemcpyDeviceToHost) - : cudaMemcpyHostToDevice; - cudaCheck(cudaMemcpyAsync(dst.data(), srcPtr, bytes, kind, stream)); - if constexpr (dstDev) { - // A synchronous destination resource retains no stream: order the - // metadata parse (which runs on the default stream) after the copy. - if constexpr (!is_async_resource::value) - cudaCheck(cudaStreamSynchronize(stream)); - } else { - cudaCheck(cudaStreamSynchronize(stream));// the host-readable result is the postcondition - } - return GridHandle(std::move(dst));// the constructor parses (and validates) the metadata -}// cuda::copyTo - -/// @brief Convenience overload issuing the copy on the source buffer's -/// retained stream when it has one (any single-space source over a -/// stream-ordered resource), the default stream otherwise. -template -inline GridHandle copyTo(const GridHandle& src, const DstBufferT* proto = nullptr) -{ - cudaStream_t stream = 0; - if constexpr (BufferHasStream::value) stream = src.buffer().stream(); - return copyTo(src, stream, proto); -}// cuda::copyTo (retained stream) - }// namespace cuda template class VectorT = std::vector> @@ -371,11 +266,11 @@ GridHandle::GridHandle(T&& buffer) }// GridHandle(T&& buffer) for single-space device buffers // Emit the dual-buffer move constructor from every CUDA translation unit, so -// host-only translation units that construct GridHandle +// host-only translation units that construct GridHandle // (they see only the declaration) always find the symbol at link time. An // unused private function is not enough: the optimizer may drop the // complete-object constructor it instantiates. -template GridHandle::GridHandle(cuda::DeviceBuffer&&); +template GridHandle::GridHandle(cuda::DualDeviceBuffer&&); } // namespace nanovdb diff --git a/nanovdb/nanovdb/cuda/HandleStorage.h b/nanovdb/nanovdb/cuda/HandleStorage.h new file mode 100644 index 0000000000..46d93f3324 --- /dev/null +++ b/nanovdb/nanovdb/cuda/HandleStorage.h @@ -0,0 +1,246 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 + +/*! + \file nanovdb/cuda/HandleStorage.h + + \brief Allocates the device-resident storage behind a GridHandle or + NodeManagerHandle for any buffer family: through the static + create() interface for buffers that provide it (the dual-space + DeviceBuffer family), and through the buffer's memory resource + for a single-space cuda::Buffer. This is the bridge that lets + every tool entry point accept either buffer family. + + \note This header is host-includable: it calls the CUDA runtime but + launches no kernels, so a plain C++ translation unit (linked + against the CUDA runtime) can allocate storage and transfer grid + handles with cuda::copyTo. +*/ + +#ifndef NANOVDB_CUDA_HANDLESTORAGE_H_HAS_BEEN_INCLUDED +#define NANOVDB_CUDA_HANDLESTORAGE_H_HAS_BEEN_INCLUDED + +#include // for the handle cuda::copyTo transfers +#include // for the BufferTraits detectors +#include // for noInit and the resource concepts +#include // for cudaCheck + +#include // for std::move +#include // for the adopted metadata + +#include // for std::runtime_error +#include // for std::is_default_constructible + +namespace nanovdb { + +namespace cuda { + +namespace detail { + +/// @brief Allocates @c bytes of device-resident storage of buffer type +/// @c BufferT: single-space buffers allocate through @c pool's +/// resource -- on @c stream when the resource is stream-ordered -- +/// and every other buffer type goes through its static +/// create(bytes, pool, device, stream) interface. +/// @param bytes size of the allocation +/// @param proto prototype buffer or null: passed through as the pool for +/// create()-style buffers; the source of the resource for +/// single-space buffers, whose resource is default-constructed when +/// @c proto is null +/// @param device device the storage lives on; single-space buffers +/// allocate on the current device, which every call site has +/// already made current +/// @param stream stream the allocation is ordered on where supported +template +inline BufferT createDeviceStorage(uint64_t bytes, const BufferT* proto, int device, cudaStream_t stream) +{ + if constexpr (BufferHasDeviceSingle::value) { + using ResourceT = typename BufferT::ResourceType; + (void)device; + if (!proto) { + if constexpr (std::is_default_constructible::value) { + if constexpr (is_async_resource::value) return BufferT(stream, bytes, noInit); + else return BufferT(bytes, noInit); + } else { + throw std::runtime_error("createDeviceStorage: a buffer over a non-default-constructible " + "resource requires a prototype buffer to take the resource from"); + } + } + if constexpr (is_async_resource::value) { + return BufferT(stream, proto->resource(), bytes, noInit); + } else { + (void)stream; + return BufferT(proto->resource(), bytes, noInit); + } + } else { + return BufferT::create(bytes, proto, device, stream); + } +} + +/// @brief The device address of a storage buffer made by +/// createDeviceStorage: data() for a single-space buffer, whose one +/// allocation is the device allocation, and deviceData() for the +/// dual-space family. +template +inline void* deviceStorageData(BufferT& buffer) +{ + if constexpr (BufferHasDeviceSingle::value) return buffer.data(); + else return buffer.deviceData(); +} + +/// @brief Orders the host after @c stream where a handle is about to be +/// constructed from bytes still being written on it: a single-space +/// buffer over a synchronous resource retains no stream, so the +/// constructor's metadata parse (which runs on the default stream) is +/// not otherwise ordered after the producer. A no-op for dual-space +/// buffers (their constructor path predates this bridge) and for +/// stream-ordered resources (the buffer retains the stream). +template +inline void orderBeforeHandleConstruction(cudaStream_t stream) +{ + if constexpr (BufferHasDeviceSingle::value) { + if constexpr (!is_async_resource::value) + cudaCheck(cudaStreamSynchronize(stream)); + } + (void)stream; +} + +/// @brief Allocates @c bytes of destination storage for a cross-space +/// transfer: single-space buffers allocate through @c proto's resource +/// (or a default-constructed resource without one), on @c stream when +/// the resource is stream-ordered; buffers providing create() go +/// through it. +template +inline DstBufferT makeTransferStorage(uint64_t bytes, cudaStream_t stream, const DstBufferT* proto) +{ + if constexpr (BufferHasDeviceSingle::value || BufferHasHostSingle::value) { + using ResourceT = typename DstBufferT::ResourceType; + if (!proto) { + // both branches of a plain conditional would instantiate the + // default-resource constructor, breaking non-default-constructible + // resources (e.g. ResourceRef) even for callers that pass a proto + if constexpr (std::is_default_constructible::value) { + if constexpr (is_async_resource::value) return DstBufferT(stream, bytes, noInit); + else return DstBufferT(bytes, noInit); + } else { + throw std::runtime_error("cuda::copyTo: a destination buffer over a non-default-constructible " + "resource requires a prototype buffer"); + } + } + if constexpr (is_async_resource::value) { + return DstBufferT(stream, proto->resource(), bytes, noInit); + } else { + (void)stream; + return DstBufferT(proto->resource(), bytes, noInit); + } + } else { + return DstBufferT::create(bytes, proto); + } +} + + +/// @brief The one gateway for constructing a GridHandle from a buffer plus +/// metadata that is already known to be valid -- adopted from another +/// handle, whose own construction from raw bytes did the validation. +struct HandleFactory +{ + template + static GridHandle make(BufferT&& buffer, std::vector meta) + { + return GridHandle(std::move(buffer), std::move(meta)); + } + + template + static const std::vector& meta(const GridHandle& handle) + { + return handle.mMetaData; + } +}; + +}// namespace detail +/// @brief Deep-copies a grid handle into a different address space: the +/// explicit, stream-carrying transfer between single-space device +/// handles and host-readable handles (HostBuffer or a host-accessible +/// single-space buffer such as a pinned-resource cuda::Buffer). +/// @tparam DstBufferT destination buffer type (specify explicitly) +/// @param src the handle to copy; must not be dual-space (use +/// deviceUpload/deviceDownload on those) +/// @param stream stream the copy is issued on; a device destination buffer +/// with a stream-ordered resource retains it +/// @warning Passing a stream other than the source buffer's retained stream +/// makes the caller responsible for ordering: prior work on the +/// source (and the source's later destruction, which frees on its +/// own stream) must be ordered against @a stream by the caller, +/// e.g. with cudaStreamWaitEvent or a synchronization. The +/// stream-less overload below has no such requirement for a source +/// with a retained stream. A source WITHOUT one (a synchronous +/// resource, e.g. a pinned-resource buffer) is the caller's to keep +/// alive under either overload: its destruction frees host memory +/// immediately, unordered against the still-asynchronous copy, so +/// synchronize @a stream before destroying such a source. (A +/// pageable HostBuffer source is exempt: its copy degrades to +/// synchronous behavior.) +/// @param proto optional buffer whose resource (or pool, for buffers +/// providing create()) allocates the destination storage; without it +/// the destination resource is default-constructed +/// @return a handle of the destination buffer type with equal contents +/// @details A host-readable destination -- HostBuffer, pinned, or a +/// both-space managed buffer -- synchronizes @c stream before +/// returning, so its host accessors are immediately valid; a +/// device-only destination is stream-ordered, so use its +/// contents on @c stream or synchronize first. The metadata is +/// adopted from the source handle -- it was validated when that +/// handle was constructed from raw bytes -- so no kernel runs and +/// this function is callable from host-only translation units. A +/// pageable host source or destination (HostBuffer) degrades the +/// copy to synchronous behavior; pinned single-space handles keep +/// it asynchronous. +template +inline GridHandle copyTo(const GridHandle& src, cudaStream_t stream, const DstBufferT* proto = nullptr) +{ + constexpr bool srcDev = BufferHasDeviceSingle::value; + constexpr bool dstDev = BufferHasDeviceSingle::value; + static_assert(!BufferTraits::hasDeviceDual && !BufferTraits::hasDeviceDual, + "cuda::copyTo does not support dual-space buffers: use deviceUpload/deviceDownload on the handle"); + static_assert(srcDev || dstDev, + "cuda::copyTo is for cross-space transfers involving a device buffer: use GridHandle::copy for host-to-host"); + const uint64_t bytes = src.bufferSize(); + if (bytes == 0u) { + if constexpr (std::is_default_constructible::value) { + return GridHandle(); + } else { + throw std::runtime_error("cuda::copyTo: an empty handle cannot be copied to a buffer type " + "that is not default-constructible"); + } + } + DstBufferT dst = detail::makeTransferStorage(bytes, stream, proto); + const void* srcPtr; + if constexpr (srcDev) srcPtr = src.deviceData(); + else srcPtr = src.data(); + cudaCheck(cudaMemcpyAsync(dst.data(), srcPtr, bytes, cudaMemcpyDefault, stream)); + if constexpr (!dstDev || BufferHasHostSingle::value) + cudaCheck(cudaStreamSynchronize(stream)); // the host-readable result is the postcondition; covers the both-space (managed) destination, whose handle exposes host accessors immediately + // A handle-to-handle copy adopts the source's metadata, which was + // validated when that handle was constructed from raw bytes -- no kernel + // runs here, which is what keeps this header host-includable. A device + // destination is stream-ordered: use it on @a stream, or synchronize. + return detail::HandleFactory::make(std::move(dst), detail::HandleFactory::meta(src)); +}// cuda::copyTo + +/// @brief Convenience overload issuing the copy on the source buffer's +/// retained stream when it has one (any single-space source over a +/// stream-ordered resource), the default stream otherwise. +template +inline GridHandle copyTo(const GridHandle& src, const DstBufferT* proto = nullptr) +{ + cudaStream_t stream = 0; + if constexpr (BufferHasStream::value) stream = src.buffer().stream(); + return copyTo(src, stream, proto); +}// cuda::copyTo (retained stream) + + +}// namespace cuda + +}// namespace nanovdb + +#endif // NANOVDB_CUDA_HANDLESTORAGE_H_HAS_BEEN_INCLUDED diff --git a/nanovdb/nanovdb/cuda/ManagedResource.h b/nanovdb/nanovdb/cuda/ManagedResource.h new file mode 100644 index 0000000000..8e389bdacb --- /dev/null +++ b/nanovdb/nanovdb/cuda/ManagedResource.h @@ -0,0 +1,73 @@ +// Copyright Contributors to the OpenVDB Project +// SPDX-License-Identifier: Apache-2.0 +// +#ifndef NANOVDB_CUDA_MANAGEDRESOURCE_H_HAS_BEEN_INCLUDED +#define NANOVDB_CUDA_MANAGEDRESOURCE_H_HAS_BEEN_INCLUDED + +#include +#include + +#include + +namespace nanovdb { + +namespace cuda { + +/// @brief Managed (unified) memory resource. Allocations +/// (cudaMallocManaged) are accessible from both the host and the +/// device, with the driver migrating pages on demand, so a container +/// over this resource serves grids that are read on both sides -- +/// the replacement for the legacy UnifiedBuffer. A GridHandle over +/// this resource parses (and validates) its metadata through the +/// device like any device buffer -- a host-side parse could race +/// still-running producer kernels -- while the host accessors remain +/// available afterwards; ordering host reads after device writes is +/// the caller's responsibility, exactly as with UnifiedBuffer. +/// @note This resource is *synchronous*: cudaMallocManaged / cudaFree have +/// no stream-ordered form, so it models the synchronous Resource +/// concept (allocate / deallocate, no stream). The usual contract +/// applies: the caller ensures device work touching an allocation has +/// completed before it is freed. +class ManagedResource +{ +public: + // cudaMallocManaged aligns to at least 256 bytes. + static constexpr size_t DEFAULT_ALIGNMENT = 256; + + /// @brief Managed allocations are mapped into the host address space + /// (detected by nanovdb::cuda::is_host_accessible_resource). + static constexpr bool HOST_ACCESSIBLE = true; + + /// @brief Managed allocations are also valid device addresses (detected + /// by nanovdb::cuda::is_device_accessible_resource), so a handle + /// over this resource exposes the device accessors as well as the + /// host ones. + static constexpr bool DEVICE_ACCESSIBLE = true; + + /// @brief Synchronous allocation of managed memory. + /// @param bytes number of bytes to allocate + /// @param alignment requested alignment (ignored; cudaMallocManaged + /// satisfies at least DEFAULT_ALIGNMENT) + void* allocate(size_t bytes, size_t alignment) { + (void)alignment; + void* p = nullptr; + cudaCheck(cudaMallocManaged(&p, bytes)); + return p; + } + + /// @brief Synchronous deallocation. + /// @param p pointer previously returned by allocate + /// @param bytes size passed to the matching allocate (unused) + /// @param alignment alignment passed to the matching allocate (unused) + void deallocate(void* p, size_t bytes, size_t alignment) { + (void)bytes; + (void)alignment; + cudaCheck(cudaFree(p)); + } +}; + +}// namespace cuda + +}// namespace nanovdb + +#endif // NANOVDB_CUDA_MANAGEDRESOURCE_H_HAS_BEEN_INCLUDED diff --git a/nanovdb/nanovdb/cuda/NodeManager.cuh b/nanovdb/nanovdb/cuda/NodeManager.cuh index eb4c897d5f..fd3c4d27a9 100644 --- a/nanovdb/nanovdb/cuda/NodeManager.cuh +++ b/nanovdb/nanovdb/cuda/NodeManager.cuh @@ -49,7 +49,7 @@ inline BufferT makeNodeManagerStorage(cudaStream_t stream, const RefT& ref, uint /// @param buffer buffer from which to allocate the output handle /// @param stream cuda stream /// @return Handle that contains a device NodeManager -template +template inline typename util::enable_if::hasDeviceDual, NodeManagerHandle>::type createNodeManager(const NanoGrid *d_grid, const BufferT& pool = BufferT(), @@ -109,15 +109,18 @@ createNodeManager(const NanoGrid *d_grid, /// as the builders. /// /// @param d_grid device grid pointer whose nodes will be accessed sequentially -/// @param resource device-resident memory resource the handle's storage -/// allocates through; a host-accessible resource (e.g. PinnedResource) -/// is excluded, since the result would carry no device NodeManager +/// @param resource memory resource with device-valid allocations the +/// handle's storage allocates through: device-resident (e.g. +/// DeviceResource) or host-and-device-accessible (ManagedResource). +/// A host-only resource (e.g. PinnedResource) is excluded, since the +/// result would carry no device NodeManager /// @param stream cuda stream /// @return Handle over cuda::Buffer> that /// contains a device NodeManager template inline typename util::enable_if<(is_resource::value || is_async_resource::value) - && !is_host_accessible_resource::value, + && (!is_host_accessible_resource::value + || is_device_accessible_resource::value), NodeManagerHandle>>>::type createNodeManager(const NanoGrid *d_grid, ResourceT& resource, @@ -175,7 +178,7 @@ createNodeManager(const NanoGrid *d_grid, }// namespace cuda -template +template [[deprecated("Use cuda::createNodeManager instead")]] inline typename util::enable_if::hasDeviceDual, NodeManagerHandle>::type cudaCreateNodeManager(const NanoGrid *d_grid, diff --git a/nanovdb/nanovdb/examples/ex_coarsen_nanovdb_cuda/coarsen_nanovdb_cuda.cpp b/nanovdb/nanovdb/examples/ex_coarsen_nanovdb_cuda/coarsen_nanovdb_cuda.cpp index 8826a3000c..2a68cd887a 100644 --- a/nanovdb/nanovdb/examples/ex_coarsen_nanovdb_cuda/coarsen_nanovdb_cuda.cpp +++ b/nanovdb/nanovdb/examples/ex_coarsen_nanovdb_cuda/coarsen_nanovdb_cuda.cpp @@ -7,7 +7,7 @@ // the following files are from NanoVDB #include -#include +#include // host-includable: cuda::copyTo transfers grids without any kernel #include template @@ -59,7 +59,7 @@ int main(int argc, char *argv[]) // Convert to indexGrid (original, un-coarsened) cpuTimer.start("Converting openVDB input to indexGrid (original version)"); - auto handleOriginal = nanovdb::tools::openToIndexVDB( + auto handleOriginal = nanovdb::tools::openToIndexVDB( grid, 0u, // Don't copy data channel false, // No stats @@ -105,7 +105,7 @@ int main(int argc, char *argv[]) // Convert to indexGrid (coarsened) cpuTimer.start("Converting openVDB input to indexGrid (coarsened version)"); - auto handleCoarsened = nanovdb::tools::openToIndexVDB( + auto handleCoarsened = nanovdb::tools::openToIndexVDB( coarsenedGrid, 0u, // Don't copy data channel false, // No stats @@ -133,10 +133,11 @@ int main(int argc, char *argv[]) } // Copy both NanoVDB grids to GPU - handleOriginal.deviceUpload(); - handleCoarsened.deviceUpload(); - auto* deviceGridOriginal = handleOriginal.deviceGrid(); - auto* deviceGridCoarsened = handleCoarsened.deviceGrid(); + // deep-copy both grids to the device; the returned handles validate them there + auto deviceHandleOriginal = nanovdb::cuda::copyTo>(handleOriginal); + auto deviceHandleCoarsened = nanovdb::cuda::copyTo>(handleCoarsened); + auto* deviceGridOriginal = deviceHandleOriginal.deviceGrid(); + auto* deviceGridCoarsened = deviceHandleCoarsened.deviceGrid(); if (!deviceGridOriginal || !deviceGridCoarsened) OPENVDB_THROW(openvdb::RuntimeError, "Failure while uploading indexGrids to GPU"); diff --git a/nanovdb/nanovdb/examples/ex_coarsen_nanovdb_cuda/coarsen_nanovdb_cuda_kernels.cu b/nanovdb/nanovdb/examples/ex_coarsen_nanovdb_cuda/coarsen_nanovdb_cuda_kernels.cu index a54b4c86c1..b8a6827ee4 100644 --- a/nanovdb/nanovdb/examples/ex_coarsen_nanovdb_cuda/coarsen_nanovdb_cuda_kernels.cu +++ b/nanovdb/nanovdb/examples/ex_coarsen_nanovdb_cuda/coarsen_nanovdb_cuda_kernels.cu @@ -2,6 +2,7 @@ // SPDX-License-Identifier: Apache-2.0 #include +#include // for cuda::copyTo template bool bufferCheck(const T* deviceBuffer, const T* hostBuffer, size_t elem_count) { @@ -54,3 +55,4 @@ void mainCoarsenGrid( nanovdb::NanoGrid *indexGridCoarsened, uint32_t benchmark_iters ); + diff --git a/nanovdb/nanovdb/examples/ex_collide_level_set/main.cc b/nanovdb/nanovdb/examples/ex_collide_level_set/main.cc index a5028fa68b..7ce7e28eb7 100644 --- a/nanovdb/nanovdb/examples/ex_collide_level_set/main.cc +++ b/nanovdb/nanovdb/examples/ex_collide_level_set/main.cc @@ -5,13 +5,8 @@ #include #include #include -#include -#if defined(NANOVDB_USE_CUDA) -using BufferT = nanovdb::cuda::DeviceBuffer; -#else -using BufferT = nanovdb::HostBuffer; -#endif +using BufferT = nanovdb::HostBuffer; // the handle lives in host memory; the CUDA side deep-copies it to the device extern void runNanoVDB(nanovdb::GridHandle& handle, int numIterations, int numPoints, BufferT& positionBuffer, BufferT& velocityBuffer); #if defined(NANOVDB_USE_OPENVDB) diff --git a/nanovdb/nanovdb/examples/ex_collide_level_set/nanovdb.cu b/nanovdb/nanovdb/examples/ex_collide_level_set/nanovdb.cu index 312e93975a..ff1878658f 100644 --- a/nanovdb/nanovdb/examples/ex_collide_level_set/nanovdb.cu +++ b/nanovdb/nanovdb/examples/ex_collide_level_set/nanovdb.cu @@ -15,10 +15,9 @@ #include "common.h" #if defined(NANOVDB_USE_CUDA) -using BufferT = nanovdb::cuda::DeviceBuffer; -#else -using BufferT = nanovdb::HostBuffer; +#include // for cuda::copyTo, the explicit host->device grid transfer #endif +using BufferT = nanovdb::HostBuffer; using namespace nanovdb; @@ -131,17 +130,20 @@ void runNanoVDB(nanovdb::GridHandle& handle, int numIterations, int num #if defined(NANOVDB_USE_CUDA) - handle.deviceUpload(); + // deep-copy the grid and the particle state to the device + auto deviceHandle = nanovdb::cuda::copyTo>(handle); - auto* d_grid = handle.deviceGrid(); + auto* d_grid = deviceHandle.deviceGrid(); if (!d_grid) throw std::runtime_error("GridHandle does not contain a valid device grid"); - positionBuffer.deviceUpload(); - Vec3f* d_positions = reinterpret_cast(positionBuffer.deviceData()); + nanovdb::cuda::Buffer devicePositions(cudaStream_t(0), size_t(numPoints), nanovdb::cuda::noInit); + cudaMemcpy(devicePositions.data(), positionBuffer.data(), size_t(numPoints) * sizeof(Vec3f), cudaMemcpyHostToDevice); + Vec3f* d_positions = devicePositions.data(); - velocityBuffer.deviceUpload(); - Vec3f* d_velocities = reinterpret_cast(velocityBuffer.deviceData()); + nanovdb::cuda::Buffer deviceVelocities(cudaStream_t(0), size_t(numPoints), nanovdb::cuda::noInit); + cudaMemcpy(deviceVelocities.data(), velocityBuffer.data(), size_t(numPoints) * sizeof(Vec3f), cudaMemcpyHostToDevice); + Vec3f* d_velocities = deviceVelocities.data(); { float durationAvg = 0; diff --git a/nanovdb/nanovdb/examples/ex_collide_level_set/openvdb.cc b/nanovdb/nanovdb/examples/ex_collide_level_set/openvdb.cc index 3a4c456376..ce1cd0c02e 100644 --- a/nanovdb/nanovdb/examples/ex_collide_level_set/openvdb.cc +++ b/nanovdb/nanovdb/examples/ex_collide_level_set/openvdb.cc @@ -12,16 +12,11 @@ #include #include -#include #include #include "common.h" -#if defined(NANOVDB_USE_CUDA) -using BufferT = nanovdb::cuda::DeviceBuffer; -#else -using BufferT = nanovdb::HostBuffer; -#endif +using BufferT = nanovdb::HostBuffer; // the handle lives in host memory; the CUDA side deep-copies it to the device openvdb::GridBase::Ptr nanoToOpenVDB(nanovdb::GridHandle& handle); diff --git a/nanovdb/nanovdb/examples/ex_dilate_nanovdb_cuda/dilate_nanovdb_cuda.cpp b/nanovdb/nanovdb/examples/ex_dilate_nanovdb_cuda/dilate_nanovdb_cuda.cpp index 6f6000f03c..2c75ff495b 100644 --- a/nanovdb/nanovdb/examples/ex_dilate_nanovdb_cuda/dilate_nanovdb_cuda.cpp +++ b/nanovdb/nanovdb/examples/ex_dilate_nanovdb_cuda/dilate_nanovdb_cuda.cpp @@ -7,7 +7,7 @@ // the following files are from NanoVDB #include -#include +#include // host-includable: cuda::copyTo transfers grids without any kernel #include template @@ -54,7 +54,7 @@ int main(int argc, char *argv[]) // Convert to indexGrid (original, un-dilated) cpuTimer.start("Converting openVDB input to indexGrid (original version)"); - auto handleOriginal = nanovdb::tools::openToIndexVDB( + auto handleOriginal = nanovdb::tools::openToIndexVDB( grid, 0u, // Don't copy data channel false, // No stats @@ -87,7 +87,7 @@ int main(int argc, char *argv[]) // Convert to indexGrid (dilated) cpuTimer.start("Converting openVDB input to indexGrid (dilated version)"); - auto handleDilated = nanovdb::tools::openToIndexVDB( + auto handleDilated = nanovdb::tools::openToIndexVDB( grid, 0u, // Don't copy data channel false, // No stats @@ -115,10 +115,11 @@ int main(int argc, char *argv[]) } // Copy both NanoVDB grids to GPU - handleOriginal.deviceUpload(); - handleDilated.deviceUpload(); - auto* deviceGridOriginal = handleOriginal.deviceGrid(); - auto* deviceGridDilated = handleDilated.deviceGrid(); + // deep-copy both grids to the device; the returned handles validate them there + auto deviceHandleOriginal = nanovdb::cuda::copyTo>(handleOriginal); + auto deviceHandleDilated = nanovdb::cuda::copyTo>(handleDilated); + auto* deviceGridOriginal = deviceHandleOriginal.deviceGrid(); + auto* deviceGridDilated = deviceHandleDilated.deviceGrid(); if (!deviceGridOriginal || !deviceGridDilated) OPENVDB_THROW(openvdb::RuntimeError, "Failure while uploading indexGrids to GPU"); diff --git a/nanovdb/nanovdb/examples/ex_dilate_nanovdb_cuda/dilate_nanovdb_cuda_kernels.cu b/nanovdb/nanovdb/examples/ex_dilate_nanovdb_cuda/dilate_nanovdb_cuda_kernels.cu index 2dcdb5389b..f9cf3bb282 100644 --- a/nanovdb/nanovdb/examples/ex_dilate_nanovdb_cuda/dilate_nanovdb_cuda_kernels.cu +++ b/nanovdb/nanovdb/examples/ex_dilate_nanovdb_cuda/dilate_nanovdb_cuda_kernels.cu @@ -4,6 +4,7 @@ #include #include #include +#include // for cuda::copyTo template bool bufferCheck(const T* deviceBuffer, const T* hostBuffer, size_t elem_count) { @@ -50,11 +51,11 @@ void mainDilateGrid( } uint32_t dstLeafCount = nanovdb::util::cuda::DeviceGridTraits::getTreeData(dstGrid).mNodeCount[0]; - nanovdb::cuda::DeviceBuffer dstLeafMaskBuffer; + nanovdb::cuda::Buffer dstLeafMaskBuffer; nanovdb::Mask<3>* dstLeafMasks = nullptr; if (dstLeafCount) { - dstLeafMaskBuffer = nanovdb::cuda::DeviceBuffer::create( std::size_t(dstLeafCount) * sizeof(nanovdb::Mask<3>), nullptr, false ); - dstLeafMasks = static_cast*>(dstLeafMaskBuffer.deviceData()); + dstLeafMaskBuffer = nanovdb::cuda::Buffer(cudaStream_t(0), std::size_t(dstLeafCount) * sizeof(nanovdb::Mask<3>), nanovdb::cuda::noInit); + dstLeafMasks = reinterpret_cast*>(dstLeafMaskBuffer.data()); if (!dstLeafMasks) throw std::runtime_error("No GPU buffer for dstLeafMask"); } @@ -100,3 +101,4 @@ void mainDilateGrid( uint32_t nnType, uint32_t benchmark_iters ); + diff --git a/nanovdb/nanovdb/examples/ex_index_grid_cuda/index_grid_cuda.cc b/nanovdb/nanovdb/examples/ex_index_grid_cuda/index_grid_cuda.cc index 8f731b014b..5924786e03 100644 --- a/nanovdb/nanovdb/examples/ex_index_grid_cuda/index_grid_cuda.cc +++ b/nanovdb/nanovdb/examples/ex_index_grid_cuda/index_grid_cuda.cc @@ -3,7 +3,7 @@ #include #include // for nanovdb::tools::createLevelSetSphere -#include // for nanovdb::cuda::DeviceBuffer +#include // host-includable: cuda::copyTo transfers grids without any kernel extern "C" void launch_kernels(const nanovdb::NanoGrid*,// device grid const nanovdb::NanoGrid*,// host grid @@ -14,28 +14,29 @@ int main(int, char**) { using SrcGridT = nanovdb::FloatGrid; using DstBuildT = nanovdb::ValueOnIndex; - using BufferT = nanovdb::cuda::DeviceBuffer; try { // Create an NanoVDB grid of a sphere at the origin with radius 100 and voxel size 1. auto srcHandle = nanovdb::tools::createLevelSetSphere(); auto *srcGrid = srcHandle.grid(); - // Converts the FloatGrid to an IndexGrid using CUDA for memory management. - auto idxHandle = nanovdb::tools::createNanoGrid(*srcGrid, 1u, false , false);// 1 channel, no tiles or stats + // Converts the FloatGrid to an IndexGrid in host memory. + auto idxHandle = nanovdb::tools::createNanoGrid(*srcGrid, 1u, false , false); // 1 channel, no tiles or stats - cudaStream_t stream; // Create a CUDA stream to allow for asynchronous copy of pinned CUDA memory. + cudaStream_t stream; // stream that orders the transfer and the kernels below cudaStreamCreate(&stream); - - idxHandle.deviceUpload(stream, false); // Copy the NanoVDB grid to the GPU asynchronously - auto* cpuGrid = idxHandle.grid(); // get a (raw) pointer to a NanoVDB grid of value type float on the CPU - auto* gpuGrid = idxHandle.deviceGrid(); // get a (raw) pointer to a NanoVDB grid of value type float on the GPU - - if (!gpuGrid) throw std::runtime_error("GridHandle did not contain a device grid with value type float"); - if (!cpuGrid) throw std::runtime_error("GridHandle did not contain a host grid with value type float"); - - launch_kernels(cpuGrid, cpuGrid, stream); // Call a host method to print a grid value on both the CPU and GPU - - cudaStreamDestroy(stream); // Destroy the CUDA stream + { + // deep-copy the grid to the GPU (implemented in the CUDA translation unit) + auto deviceHandle = nanovdb::cuda::copyTo>(idxHandle, stream); + auto* cpuGrid = idxHandle.grid(); + auto* gpuGrid = deviceHandle.deviceGrid(); + + if (!gpuGrid) throw std::runtime_error("GridHandle did not contain a device grid with value type float"); + if (!cpuGrid) throw std::runtime_error("GridHandle did not contain a host grid with value type float"); + + launch_kernels(gpuGrid, cpuGrid, stream); // print a grid value on both the CPU and GPU + cudaStreamSynchronize(stream); // the kernels must finish before the device handle (whose buffer frees on this stream) goes away + } + cudaStreamDestroy(stream); } catch (const std::exception& e) { std::cerr << "An exception occurred: \"" << e.what() << "\"" << std::endl; diff --git a/nanovdb/nanovdb/examples/ex_index_grid_cuda/index_grid_cuda_kernel.cu b/nanovdb/nanovdb/examples/ex_index_grid_cuda/index_grid_cuda_kernel.cu index 7639ead7c9..c0647f78f7 100644 --- a/nanovdb/nanovdb/examples/ex_index_grid_cuda/index_grid_cuda_kernel.cu +++ b/nanovdb/nanovdb/examples/ex_index_grid_cuda/index_grid_cuda_kernel.cu @@ -2,7 +2,8 @@ // SPDX-License-Identifier: Apache-2.0 #include // this defined the core tree data structure of NanoVDB accessable on both the host and device -#include // required since GridHandle has device code +#include // for cuda::copyTo, the explicit host->device grid transfer + #include // for printf #include diff --git a/nanovdb/nanovdb/examples/ex_make_custom_nanovdb_cuda/make_custom_nanovdb_cuda.cc b/nanovdb/nanovdb/examples/ex_make_custom_nanovdb_cuda/make_custom_nanovdb_cuda.cc index 64f1b11c9b..321697f673 100644 --- a/nanovdb/nanovdb/examples/ex_make_custom_nanovdb_cuda/make_custom_nanovdb_cuda.cc +++ b/nanovdb/nanovdb/examples/ex_make_custom_nanovdb_cuda/make_custom_nanovdb_cuda.cc @@ -5,7 +5,7 @@ #include #include -#include +#include // host-includable: cuda::copyTo transfers grids without any kernel #include @@ -26,20 +26,23 @@ int main() printf("build::Grid: (%i,%i,%i)=%4.2f\n", 1, 2,-3, acc.getValue(nanovdb::Coord(1, 2,-3))); printf("build::Grid: (%i,%i,%i)=%4.2f\n", 1, 2, 3, acc.getValue(nanovdb::Coord(1, 2, 3))); - // convert build::grid to a nanovdb::GridHandle using a Cuda buffer - auto handle = nanovdb::tools::createNanoGrid(grid); + // convert build::grid to a nanovdb::GridHandle in host memory + auto handle = nanovdb::tools::createNanoGrid(grid); auto* cpuGrid = handle.grid(); //get a (raw) pointer to a NanoVDB grid of value type float on the CPU if (!cpuGrid) throw std::runtime_error("GridHandle does not contain a grid with value type float"); - cudaStream_t stream; // Create a CUDA stream to allow for asynchronous copy of pinned CUDA memory. + cudaStream_t stream; // stream that orders the transfer and the kernels below cudaStreamCreate(&stream); - - handle.deviceUpload(stream, false); // Copy the NanoVDB grid to the GPU asynchronously - auto* gpuGrid = handle.deviceGrid(); // get a (raw) pointer to a NanoVDB grid of value type float on the GPU - - launch_kernels(gpuGrid, cpuGrid, stream); // Call a host method to print a grid values on both the CPU and GPU - cudaStreamDestroy(stream); // Destroy the CUDA stream + { + // deep-copy the grid to the GPU (implemented in the CUDA translation unit) + auto deviceHandle = nanovdb::cuda::copyTo>(handle, stream); + auto* gpuGrid = deviceHandle.deviceGrid(); + + launch_kernels(gpuGrid, cpuGrid, stream); // print grid values on both the CPU and GPU + cudaStreamSynchronize(stream); // the kernels must finish before the device handle (whose buffer frees on this stream) goes away + } + cudaStreamDestroy(stream); } catch (const std::exception& e) { std::cerr << "An exception occurred: \"" << e.what() << "\"" << std::endl; diff --git a/nanovdb/nanovdb/examples/ex_make_custom_nanovdb_cuda/make_custom_nanovdb_cuda_kernel.cu b/nanovdb/nanovdb/examples/ex_make_custom_nanovdb_cuda/make_custom_nanovdb_cuda_kernel.cu index 70ba0e87cb..8fde3af441 100644 --- a/nanovdb/nanovdb/examples/ex_make_custom_nanovdb_cuda/make_custom_nanovdb_cuda_kernel.cu +++ b/nanovdb/nanovdb/examples/ex_make_custom_nanovdb_cuda/make_custom_nanovdb_cuda_kernel.cu @@ -24,6 +24,7 @@ __global__ void gpu_kernel(const nanovdb::NanoGrid* deviceGrid) } // This is called by the client code on the host + extern "C" void launch_kernels(const nanovdb::NanoGrid* deviceGrid, const nanovdb::NanoGrid* cpuGrid, cudaStream_t stream) diff --git a/nanovdb/nanovdb/examples/ex_make_mgpu_nanovdb/main.cu b/nanovdb/nanovdb/examples/ex_make_mgpu_nanovdb/main.cu index 7d7630976f..ad7016396f 100644 --- a/nanovdb/nanovdb/examples/ex_make_mgpu_nanovdb/main.cu +++ b/nanovdb/nanovdb/examples/ex_make_mgpu_nanovdb/main.cu @@ -192,7 +192,7 @@ void testConvolution() auto floatHandle = nanovdb::tools::createLevelSetSphere(100, nanovdb::Vec3d(0), 1, 3, nanovdb::Vec3d(0), "test"); nanovdb::FloatGrid* floatGrid = floatHandle.grid(); - using BufferT = nanovdb::cuda::DeviceBuffer; + using BufferT = nanovdb::cuda::DualDeviceBuffer; // migrates with the multi-GPU work that retires UnifiedBuffer auto indexHandle = nanovdb::tools::createNanoGrid(*floatGrid, 0u, false, false, 1); std::for_each(deviceMesh.begin(), deviceMesh.end(), [&](const nanovdb::cuda::DeviceNode& node) {// copy host buffer to all the device buffers cudaCheck(cudaSetDevice(node.id)); diff --git a/nanovdb/nanovdb/examples/ex_merge_nanovdb_cuda/merge_nanovdb_cuda.cpp b/nanovdb/nanovdb/examples/ex_merge_nanovdb_cuda/merge_nanovdb_cuda.cpp index 8b6626558b..98f393cee4 100644 --- a/nanovdb/nanovdb/examples/ex_merge_nanovdb_cuda/merge_nanovdb_cuda.cpp +++ b/nanovdb/nanovdb/examples/ex_merge_nanovdb_cuda/merge_nanovdb_cuda.cpp @@ -7,7 +7,7 @@ // the following files are from NanoVDB #include -#include +#include // host-includable: cuda::copyTo transfers grids without any kernel #include template @@ -60,7 +60,7 @@ int main(int argc, char *argv[]) // Convert to indexGrid cpuTimer.start("Converting openVDB input to indexGrid (first component)"); - auto srcHandle1 = nanovdb::tools::openToIndexVDB( + auto srcHandle1 = nanovdb::tools::openToIndexVDB( grid1, 0u, // Don't copy data channel false, // No stats @@ -71,7 +71,7 @@ int main(int argc, char *argv[]) cpuTimer.stop(); cpuTimer.start("Converting openVDB input to indexGrid (second component)"); - auto srcHandle2 = nanovdb::tools::openToIndexVDB( + auto srcHandle2 = nanovdb::tools::openToIndexVDB( grid2, 0u, // Don't copy data channel false, // No stats @@ -118,7 +118,7 @@ int main(int argc, char *argv[]) // Convert to indexGrid cpuTimer.start("Converting merged openVDB output to indexGrid"); - auto dstReferenceHandle = nanovdb::tools::openToIndexVDB( + auto dstReferenceHandle = nanovdb::tools::openToIndexVDB( mergedGrid, 0u, // Don't copy data channel false, // No stats @@ -144,13 +144,13 @@ int main(int argc, char *argv[]) std::cout << "Memory usage : " << dstReferenceGrid->gridSize() << " bytes" << std::endl; } - // Copy both NanoVDB grids to GPU - srcHandle1.deviceUpload(); - srcHandle2.deviceUpload(); - dstReferenceHandle.deviceUpload(); - auto* deviceSrcGrid1 = srcHandle1.deviceGrid(); - auto* deviceSrcGrid2 = srcHandle2.deviceGrid(); - auto* deviceDstReferenceGrid = dstReferenceHandle.deviceGrid(); + // Deep-copy all three NanoVDB grids to the GPU; the returned handles validate them there + auto deviceSrcHandle1 = nanovdb::cuda::copyTo>(srcHandle1); + auto deviceSrcHandle2 = nanovdb::cuda::copyTo>(srcHandle2); + auto deviceDstReferenceHandle = nanovdb::cuda::copyTo>(dstReferenceHandle); + auto* deviceSrcGrid1 = deviceSrcHandle1.deviceGrid(); + auto* deviceSrcGrid2 = deviceSrcHandle2.deviceGrid(); + auto* deviceDstReferenceGrid = deviceDstReferenceHandle.deviceGrid(); if (!deviceSrcGrid1 || !deviceSrcGrid2 || !deviceDstReferenceGrid) OPENVDB_THROW(openvdb::RuntimeError, "Failure while uploading indexGrids to GPU"); diff --git a/nanovdb/nanovdb/examples/ex_merge_nanovdb_cuda/merge_nanovdb_cuda_kernels.cu b/nanovdb/nanovdb/examples/ex_merge_nanovdb_cuda/merge_nanovdb_cuda_kernels.cu index d29b849b8b..136e552db0 100644 --- a/nanovdb/nanovdb/examples/ex_merge_nanovdb_cuda/merge_nanovdb_cuda_kernels.cu +++ b/nanovdb/nanovdb/examples/ex_merge_nanovdb_cuda/merge_nanovdb_cuda_kernels.cu @@ -2,6 +2,7 @@ // SPDX-License-Identifier: Apache-2.0 #include +#include // for cuda::copyTo template bool bufferCheck(const T* deviceBuffer, const T* hostBuffer, size_t elem_count) { @@ -57,3 +58,4 @@ void mainMergeGrids( nanovdb::NanoGrid *hostSrcGrid2, nanovdb::NanoGrid *hostDstReferenceGrid, uint32_t benchmark_iters); + diff --git a/nanovdb/nanovdb/examples/ex_modify_nanovdb_thrust/modify_nanovdb_thrust.cc b/nanovdb/nanovdb/examples/ex_modify_nanovdb_thrust/modify_nanovdb_thrust.cc index f2bb4ecd3e..270592c5b2 100644 --- a/nanovdb/nanovdb/examples/ex_modify_nanovdb_thrust/modify_nanovdb_thrust.cc +++ b/nanovdb/nanovdb/examples/ex_modify_nanovdb_thrust/modify_nanovdb_thrust.cc @@ -5,7 +5,7 @@ /// modified on the device. It depends on NanoVDB and CUDA thrust. #include -#include +#include // host-includable: cuda::copyTo transfers grids without any kernel extern "C" void scaleActiveVoxels(nanovdb::FloatGrid *grid_d, uint64_t leafCount, float scale); @@ -13,13 +13,14 @@ int main() { try { // Create an NanoVDB grid of a sphere at the origin with radius 100 and voxel size 1. - auto handle = nanovdb::tools::createLevelSetSphere(100.0f); + auto handle = nanovdb::tools::createLevelSetSphere(100.0f); using GridT = nanovdb::FloatGrid; - handle.deviceUpload(nullptr, false); // Copy the NanoVDB grid to the GPU asynchronously + // deep-copy the grid to the device -- callable right here, in a host-only file + auto deviceHandle = nanovdb::cuda::copyTo>(handle); - const GridT* grid = handle.grid(); // get a (raw) const pointer to a NanoVDB grid of value type float on the CPU - GridT* deviceGrid = handle.deviceGrid(); // get a (raw) pointer to a NanoVDB grid of value type float on the GPU + const GridT* grid = handle.grid(); // a (raw) const pointer to the grid on the CPU + GridT* deviceGrid = deviceHandle.deviceGrid(); // and its deep copy on the GPU if (!deviceGrid || !grid) { throw std::runtime_error("GridHandle did not contain a grid with value type float"); @@ -32,9 +33,10 @@ int main() scaleActiveVoxels(deviceGrid, grid->tree().nodeCount(0), 2.0f); - handle.deviceDownload(nullptr, true); // Copy the NanoVDB grid to the CPU synchronously + // copy the modified grid back to the host and read the result from the returned handle + auto result = nanovdb::cuda::copyTo(deviceHandle); - std::cout << "Value after scaling = " << grid->tree().getValue(nanovdb::Coord(101,0,0)) << std::endl; + std::cout << "Value after scaling = " << result.grid()->tree().getValue(nanovdb::Coord(101,0,0)) << std::endl; } catch (const std::exception& e) { std::cerr << "An exception occurred: \"" << e.what() << "\"" << std::endl; diff --git a/nanovdb/nanovdb/examples/ex_modify_nanovdb_thrust/modify_nanovdb_thrust.cu b/nanovdb/nanovdb/examples/ex_modify_nanovdb_thrust/modify_nanovdb_thrust.cu index 4c602f8a44..4df4e0944a 100644 --- a/nanovdb/nanovdb/examples/ex_modify_nanovdb_thrust/modify_nanovdb_thrust.cu +++ b/nanovdb/nanovdb/examples/ex_modify_nanovdb_thrust/modify_nanovdb_thrust.cu @@ -7,8 +7,8 @@ #include #include -#include -#include +#include // for cuda::copyTo, the explicit host<->device grid transfer + extern "C" void scaleActiveVoxels(nanovdb::FloatGrid *grid_d, uint64_t leafCount, float scale) { diff --git a/nanovdb/nanovdb/examples/ex_nodemanager_cuda/nodemanager_cuda.cc b/nanovdb/nanovdb/examples/ex_nodemanager_cuda/nodemanager_cuda.cc index ae912a2493..5b05f79958 100644 --- a/nanovdb/nanovdb/examples/ex_nodemanager_cuda/nodemanager_cuda.cc +++ b/nanovdb/nanovdb/examples/ex_nodemanager_cuda/nodemanager_cuda.cc @@ -3,53 +3,49 @@ #include // replace with your own dependencies for generating the OpenVDB grid #include // converter from OpenVDB to NanoVDB (includes NanoVDB.h and GridManager.h) -#include +#include // host-includable: cuda::copyTo transfers grids without any kernel +#include #include extern "C" void launch_kernels(const nanovdb::NodeManager*,// device NaodeManager const nanovdb::NodeManager*,// host NodeManager cudaStream_t stream); -extern "C" void cudaCreateNodeManager(const nanovdb::NanoGrid*,// device grid - nanovdb::NodeManagerHandle*);// Handle to device NodeManager +extern nanovdb::NodeManagerHandle>> +uploadNodeManager(const nanovdb::NanoGrid* d_grid, cudaStream_t stream); // constructs a NodeManager for a device grid /// @brief This examples depends on OpenVDB, NanoVDB and CUDA. int main() { using SrcGridT = openvdb::FloatGrid; - using BufferT = nanovdb::cuda::DeviceBuffer; try { - cudaStream_t stream; // Create a CUDA stream to allow for asynchronous copy of pinned CUDA memory. + cudaStream_t stream; // stream that orders the transfers and the kernels below cudaStreamCreate(&stream); + { // Create an OpenVDB grid of a sphere at the origin with radius 100 and voxel size 1. auto srcGrid = openvdb::tools::createLevelSetSphere(100.0f, openvdb::Vec3f(0.0f), 1.0f); // Converts the OpenVDB to NanoVDB and returns a GridHandle that uses CUDA for memory management. - auto gridHandle = nanovdb::tools::createNanoGrid(*srcGrid); - gridHandle.deviceUpload(stream, false); // Copy the NanoVDB grid to the GPU asynchronously - auto* grid = gridHandle.grid(); // get a (raw) pointer to a NanoVDB grid of value type float on the CPU - auto* deviceGrid = gridHandle.deviceGrid(); // get a (raw) pointer to a NanoVDB grid of value type float on the GPU + auto gridHandle = nanovdb::tools::createNanoGrid(*srcGrid); + auto deviceGridHandle = nanovdb::cuda::copyTo>(gridHandle, stream); // deep-copy the grid to the GPU + auto* grid = gridHandle.grid(); // a (raw) pointer to the grid on the CPU + auto* deviceGrid = deviceGridHandle.deviceGrid(); // and its deep copy on the GPU if (!deviceGrid || !grid) { throw std::runtime_error("GridHandle did not contain a grid with value type float"); } - auto nodeHandle = nanovdb::createNodeManager(*grid); + auto nodeHandle = nanovdb::createNodeManager(*grid); // host NodeManager over the host grid auto *nodeMgr = nodeHandle.template mgr(); -#if 0// this approach copies a NodeManager from host to device - nodeHandle.deviceUpload(deviceGrid, stream, false); - auto *deviceNodeMgr = nodeHandle.template deviceMgr(); -#else// the approach below constructs a new NodeManager directly for a device grid - nanovdb::NodeManagerHandle nodeHandle2; - cudaCreateNodeManager(deviceGrid, &nodeHandle2); + auto nodeHandle2 = uploadNodeManager(deviceGrid, stream); // device NodeManager constructed for the device grid auto *deviceNodeMgr = nodeHandle2.template deviceMgr(); -#endif if (!deviceNodeMgr || !nodeMgr) { throw std::runtime_error("NodeManagerHandle did not contain a grid with value type float"); } launch_kernels(deviceNodeMgr, nodeMgr, stream); // Call a host method to print a grid value on both the CPU and GPU - + cudaStreamSynchronize(stream); // the kernels must finish before the device handles (whose buffers free on this stream) go away + } cudaStreamDestroy(stream); // Destroy the CUDA stream } catch (const std::exception& e) { diff --git a/nanovdb/nanovdb/examples/ex_nodemanager_cuda/nodemanager_cuda_kernel.cu b/nanovdb/nanovdb/examples/ex_nodemanager_cuda/nodemanager_cuda_kernel.cu index 4fc6b251fe..4e5e889a43 100644 --- a/nanovdb/nanovdb/examples/ex_nodemanager_cuda/nodemanager_cuda_kernel.cu +++ b/nanovdb/nanovdb/examples/ex_nodemanager_cuda/nodemanager_cuda_kernel.cu @@ -30,8 +30,10 @@ extern "C" void launch_kernels(const nanovdb::NodeManager* deviceMgr, } // Simple wrapper that makes sure nanovdb::cuda::createNodeManager is initiated -extern "C" void cudaCreateNodeManager(const nanovdb::NanoGrid *d_grid, - nanovdb::NodeManagerHandle *handle) +// Constructs a NodeManager for a device grid, allocated through the default +// device memory resource (a program-lifetime singleton the handle borrows). +nanovdb::NodeManagerHandle>> +uploadNodeManager(const nanovdb::NanoGrid *d_grid, cudaStream_t stream) { - *handle = std::move(nanovdb::cuda::createNodeManager(d_grid)); + return nanovdb::cuda::createNodeManager(d_grid, nanovdb::cuda::default_resource(), stream); } diff --git a/nanovdb/nanovdb/examples/ex_openvdb_to_nanovdb_cuda/openvdb_to_nanovdb_cuda.cc b/nanovdb/nanovdb/examples/ex_openvdb_to_nanovdb_cuda/openvdb_to_nanovdb_cuda.cc index 73441474ea..2f25b12a37 100644 --- a/nanovdb/nanovdb/examples/ex_openvdb_to_nanovdb_cuda/openvdb_to_nanovdb_cuda.cc +++ b/nanovdb/nanovdb/examples/ex_openvdb_to_nanovdb_cuda/openvdb_to_nanovdb_cuda.cc @@ -3,7 +3,7 @@ #include // replace with your own dependencies for generating the OpenVDB grid #include // converter from OpenVDB to NanoVDB (includes NanoVDB.h and GridManager.h) -#include +#include // host-includable: cuda::copyTo transfers grids without any kernel extern "C" void launch_kernels(const nanovdb::NanoGrid*, const nanovdb::NanoGrid*, @@ -18,21 +18,23 @@ int main(int, char**) auto srcGrid = openvdb::tools::createLevelSetSphere(100.0f, openvdb::Vec3f(0.0f), 1.0f); // Converts the OpenVDB to NanoVDB and returns a GridHandle that uses CUDA for memory management. - auto handle = nanovdb::tools::createNanoGrid(*srcGrid); + auto handle = nanovdb::tools::createNanoGrid(*srcGrid); - cudaStream_t stream; // Create a CUDA stream to allow for asynchronous copy of pinned CUDA memory. + cudaStream_t stream; // stream that orders the transfer and the kernels below cudaStreamCreate(&stream); - - handle.deviceUpload(stream, false); // Copy the NanoVDB grid to the GPU asynchronously + { + // deep-copy the grid to the GPU (implemented in the CUDA translation unit) + auto deviceHandle = nanovdb::cuda::copyTo>(handle, stream); auto* grid = handle.grid(); // get a (raw) pointer to a NanoVDB grid of value type float on the CPU - auto* deviceGrid = handle.deviceGrid(); // get a (raw) pointer to a NanoVDB grid of value type float on the GPU + auto* deviceGrid = deviceHandle.deviceGrid(); // get a (raw) pointer to a NanoVDB grid of value type float on the GPU if (!deviceGrid || !grid) throw std::runtime_error("GridHandle did not contain a grid with value type float"); launch_kernels(deviceGrid, grid, stream); // Call a host method to print a grid value on both the CPU and GPU - + cudaStreamSynchronize(stream); // the kernels must finish before the device handle (whose buffer frees on this stream) goes away + } cudaStreamDestroy(stream); // Destroy the CUDA stream } catch (const std::exception& e) { diff --git a/nanovdb/nanovdb/examples/ex_openvdb_to_nanovdb_cuda/openvdb_to_nanovdb_cuda_kernel.cu b/nanovdb/nanovdb/examples/ex_openvdb_to_nanovdb_cuda/openvdb_to_nanovdb_cuda_kernel.cu index e0bf9a021d..b47233a685 100644 --- a/nanovdb/nanovdb/examples/ex_openvdb_to_nanovdb_cuda/openvdb_to_nanovdb_cuda_kernel.cu +++ b/nanovdb/nanovdb/examples/ex_openvdb_to_nanovdb_cuda/openvdb_to_nanovdb_cuda_kernel.cu @@ -2,7 +2,7 @@ // SPDX-License-Identifier: Apache-2.0 #include // this defined the core tree data structure of NanoVDB accessable on both the host and device -#include // required since GridHandle has device code +#include #include // for printf // This is called by the host only diff --git a/nanovdb/nanovdb/examples/ex_raytrace_fog_volume/main.cc b/nanovdb/nanovdb/examples/ex_raytrace_fog_volume/main.cc index a82708b136..a0d8bac30f 100644 --- a/nanovdb/nanovdb/examples/ex_raytrace_fog_volume/main.cc +++ b/nanovdb/nanovdb/examples/ex_raytrace_fog_volume/main.cc @@ -6,12 +6,7 @@ #include #include -#if defined(NANOVDB_USE_CUDA) -#include -using BufferT = nanovdb::cuda::DeviceBuffer; -#else -using BufferT = nanovdb::HostBuffer; -#endif +using BufferT = nanovdb::HostBuffer; // the handle lives in host memory; the CUDA side deep-copies it to the device extern void runNanoVDB(nanovdb::GridHandle& handle, int numIterations, int width, int height, BufferT& imageBuffer); #if defined(NANOVDB_USE_OPENVDB) diff --git a/nanovdb/nanovdb/examples/ex_raytrace_fog_volume/nanovdb.cu b/nanovdb/nanovdb/examples/ex_raytrace_fog_volume/nanovdb.cu index e0e7577610..a9fab82604 100644 --- a/nanovdb/nanovdb/examples/ex_raytrace_fog_volume/nanovdb.cu +++ b/nanovdb/nanovdb/examples/ex_raytrace_fog_volume/nanovdb.cu @@ -8,11 +8,9 @@ #include #if defined(NANOVDB_USE_CUDA) -#include -using BufferT = nanovdb::cuda::DeviceBuffer; -#else -using BufferT = nanovdb::HostBuffer; +#include // for cuda::copyTo, the explicit host->device grid transfer #endif +using BufferT = nanovdb::HostBuffer; #include #include #include @@ -86,14 +84,15 @@ void runNanoVDB(nanovdb::GridHandle& handle, int numIterations, int wid } #if defined(NANOVDB_USE_CUDA) - handle.deviceUpload(); + // deep-copy the grid to the device; the returned handle validates it there + auto deviceHandle = nanovdb::cuda::copyTo>(handle); - auto* d_grid = handle.deviceGrid(); + auto* d_grid = deviceHandle.deviceGrid(); if (!d_grid) throw std::runtime_error("GridHandle does not contain a valid device grid"); - imageBuffer.deviceUpload(); - float* d_outImage = reinterpret_cast(imageBuffer.deviceData()); + nanovdb::cuda::Buffer deviceImage(cudaStream_t(0), size_t(width) * height, nanovdb::cuda::noInit); + float* d_outImage = deviceImage.data(); { float durationAvg = 0; @@ -105,7 +104,7 @@ void runNanoVDB(nanovdb::GridHandle& handle, int numIterations, int wid durationAvg /= numIterations; std::cout << "Average Duration(NanoVDB-Cuda) = " << durationAvg << " ms" << std::endl; - imageBuffer.deviceDownload(); + cudaMemcpy(imageBuffer.data(), deviceImage.data(), size_t(width) * height * sizeof(float), cudaMemcpyDeviceToHost); saveImage("raytrace_fog_volume-nanovdb-cuda.pfm", width, height, (float*)imageBuffer.data()); } #endif diff --git a/nanovdb/nanovdb/examples/ex_raytrace_fog_volume/openvdb.cc b/nanovdb/nanovdb/examples/ex_raytrace_fog_volume/openvdb.cc index 16360157fd..9b458b6164 100644 --- a/nanovdb/nanovdb/examples/ex_raytrace_fog_volume/openvdb.cc +++ b/nanovdb/nanovdb/examples/ex_raytrace_fog_volume/openvdb.cc @@ -17,12 +17,7 @@ #include "common.h" -#if defined(NANOVDB_USE_CUDA) -#include -using BufferT = nanovdb::cuda::DeviceBuffer; -#else -using BufferT = nanovdb::HostBuffer; -#endif +using BufferT = nanovdb::HostBuffer; // the handle lives in host memory; the CUDA side deep-copies it to the device void runOpenVDB(nanovdb::GridHandle& handle, int numIterations, int width, int height, BufferT& imageBuffer) { diff --git a/nanovdb/nanovdb/examples/ex_raytrace_iso_surface/main.cc b/nanovdb/nanovdb/examples/ex_raytrace_iso_surface/main.cc index db5fb6952d..48c3aaf0a3 100644 --- a/nanovdb/nanovdb/examples/ex_raytrace_iso_surface/main.cc +++ b/nanovdb/nanovdb/examples/ex_raytrace_iso_surface/main.cc @@ -6,13 +6,8 @@ #include #include #include -#include -#if defined(NANOVDB_USE_CUDA) -using BufferT = nanovdb::cuda::DeviceBuffer; -#else -using BufferT = nanovdb::HostBuffer; -#endif +using BufferT = nanovdb::HostBuffer; // the handle lives in host memory; the CUDA side deep-copies it to the device extern void runNanoVDB(nanovdb::GridHandle& handle, int numIterations, int width, int height, BufferT& imageBuffer, bool usePersistentThreads); diff --git a/nanovdb/nanovdb/examples/ex_raytrace_iso_surface/nanovdb.cu b/nanovdb/nanovdb/examples/ex_raytrace_iso_surface/nanovdb.cu index 3eafd79ce8..ab15700f0b 100644 --- a/nanovdb/nanovdb/examples/ex_raytrace_iso_surface/nanovdb.cu +++ b/nanovdb/nanovdb/examples/ex_raytrace_iso_surface/nanovdb.cu @@ -8,11 +8,9 @@ #include #if defined(NANOVDB_USE_CUDA) -#include -using BufferT = nanovdb::cuda::DeviceBuffer; -#else -using BufferT = nanovdb::HostBuffer; +#include // for cuda::copyTo, the explicit host->device grid transfer #endif +using BufferT = nanovdb::HostBuffer; #include #include #include @@ -32,12 +30,13 @@ void runNanoVDB(nanovdb::GridHandle& handle, int numIterations, int wid renderOp.saveImage("raytrace_iso_surface-nanovdb-host.pfm", (float*)imageBuffer.data()); #if defined(NANOVDB_USE_CUDA) - handle.deviceUpload(); + // deep-copy the grid to the device; the returned handle validates it there + auto deviceHandle = nanovdb::cuda::copyTo>(handle); using BuildT = typename nanovdb::util::remove_pointer_t::BuildType; - auto* d_grid = handle.deviceGrid(); + auto* d_grid = deviceHandle.deviceGrid(); if (!d_grid) throw std::runtime_error("GridHandle does not contain a valid device grid"); - imageBuffer.deviceUpload(); - float* d_outImage = reinterpret_cast(imageBuffer.deviceData()); + nanovdb::cuda::Buffer deviceImage(cudaStream_t(0), size_t(width) * height, nanovdb::cuda::noInit); + float* d_outImage = deviceImage.data(); sum = 0; if (usePersistentThreads) { int* d_nextPixel = nullptr; @@ -45,12 +44,12 @@ void runNanoVDB(nanovdb::GridHandle& handle, int numIterations, int wid for (int i = 0; i < numIterations; ++i, sum += renderOp.renderImagePersistent(d_outImage, d_grid, d_nextPixel)); NANOVDB_CUDA_CHECK_ERROR(cudaFree(d_nextPixel), __FILE__, __LINE__); std::cout << "Average of " << numIterations << " renderings (NanoVDB-Cuda-Persistent) = " << (sum/numIterations) << " ms " << std::endl; - imageBuffer.deviceDownload(); + cudaMemcpy(imageBuffer.data(), deviceImage.data(), size_t(width) * height * sizeof(float), cudaMemcpyDeviceToHost); renderOp.saveImage("raytrace_iso_surface-nanovdb-cuda-persistent.pfm", (float*)imageBuffer.data()); } else { for (int i = 0; i < numIterations; ++i, sum += renderOp.renderImage(true/*useCuda*/, d_outImage, d_grid)); std::cout << "Average of " << numIterations << " renderings (NanoVDB-Cuda) = " << (sum/numIterations) << " ms " << std::endl; - imageBuffer.deviceDownload(); + cudaMemcpy(imageBuffer.data(), deviceImage.data(), size_t(width) * height * sizeof(float), cudaMemcpyDeviceToHost); renderOp.saveImage("raytrace_iso_surface-nanovdb-cuda.pfm", (float*)imageBuffer.data()); } #endif diff --git a/nanovdb/nanovdb/examples/ex_raytrace_level_set/main.cc b/nanovdb/nanovdb/examples/ex_raytrace_level_set/main.cc index 46a3c6f48f..7f57ac18a3 100644 --- a/nanovdb/nanovdb/examples/ex_raytrace_level_set/main.cc +++ b/nanovdb/nanovdb/examples/ex_raytrace_level_set/main.cc @@ -5,13 +5,8 @@ #include #include #include -#include -#if defined(NANOVDB_USE_CUDA) -using BufferT = nanovdb::cuda::DeviceBuffer; -#else -using BufferT = nanovdb::HostBuffer; -#endif +using BufferT = nanovdb::HostBuffer; // the handle lives in host memory; the CUDA side deep-copies it to the device extern void runNanoVDB(nanovdb::GridHandle& handle, int numIterations, int width, int height, BufferT& imageBuffer); #if defined(NANOVDB_USE_OPENVDB) diff --git a/nanovdb/nanovdb/examples/ex_raytrace_level_set/nanovdb.cu b/nanovdb/nanovdb/examples/ex_raytrace_level_set/nanovdb.cu index ded2c6e1d2..a1212ca205 100644 --- a/nanovdb/nanovdb/examples/ex_raytrace_level_set/nanovdb.cu +++ b/nanovdb/nanovdb/examples/ex_raytrace_level_set/nanovdb.cu @@ -10,11 +10,9 @@ #include #if defined(NANOVDB_USE_CUDA) -#include -using BufferT = nanovdb::cuda::DeviceBuffer; -#else -using BufferT = nanovdb::HostBuffer; +#include // for cuda::copyTo, the explicit host->device grid transfer #endif +using BufferT = nanovdb::HostBuffer; #include #include #include @@ -116,14 +114,15 @@ void runNanoVDB(nanovdb::GridHandle& handle, int numIterations, int wid } #if defined(NANOVDB_USE_CUDA) - handle.deviceUpload(); + // deep-copy the grid to the device; the returned handle validates it there + auto deviceHandle = nanovdb::cuda::copyTo>(handle); - auto* d_grid = handle.deviceGrid(); + auto* d_grid = deviceHandle.deviceGrid(); if (!d_grid) throw std::runtime_error("GridHandle does not contain a valid device grid"); - imageBuffer.deviceUpload(); - float* d_outImage = reinterpret_cast(imageBuffer.deviceData()); + nanovdb::cuda::Buffer deviceImage(cudaStream_t(0), size_t(width) * height, nanovdb::cuda::noInit); + float* d_outImage = deviceImage.data(); { for (int i = 0; i < NUM_WARMUP_ITERATIONS; ++i) { @@ -137,7 +136,7 @@ void runNanoVDB(nanovdb::GridHandle& handle, int numIterations, int wid } reportStats("Duration(NanoVDB-Cuda):", samples); - imageBuffer.deviceDownload(); + cudaMemcpy(imageBuffer.data(), deviceImage.data(), size_t(width) * height * sizeof(float), cudaMemcpyDeviceToHost); saveImage("raytrace_level_set-nanovdb-cuda.pfm", width, height, (float*)imageBuffer.data()); } #endif diff --git a/nanovdb/nanovdb/examples/ex_raytrace_level_set/openvdb.cc b/nanovdb/nanovdb/examples/ex_raytrace_level_set/openvdb.cc index 75ee170acc..04b0b02c30 100644 --- a/nanovdb/nanovdb/examples/ex_raytrace_level_set/openvdb.cc +++ b/nanovdb/nanovdb/examples/ex_raytrace_level_set/openvdb.cc @@ -17,14 +17,9 @@ #include "common.h" -#if defined(NANOVDB_USE_CUDA) -#include -using BufferT = nanovdb::cuda::DeviceBuffer; -#else using BufferT = nanovdb::HostBuffer; -#endif -void runOpenVDB(nanovdb::GridHandle& handle, int numIterations, int width, int height, BufferT& imageBuffer) +void runOpenVDB(nanovdb::GridHandle& handle, int numIterations, int width, int height, BufferT& imageBuffer) { using GridT = openvdb::FloatGrid; using CoordT = openvdb::Coord; diff --git a/nanovdb/nanovdb/examples/ex_read_nanovdb_sphere_accessor_cuda/read_nanovdb_sphere_accessor_cuda.cu b/nanovdb/nanovdb/examples/ex_read_nanovdb_sphere_accessor_cuda/read_nanovdb_sphere_accessor_cuda.cu index d2c7263a48..0badd508ef 100644 --- a/nanovdb/nanovdb/examples/ex_read_nanovdb_sphere_accessor_cuda/read_nanovdb_sphere_accessor_cuda.cu +++ b/nanovdb/nanovdb/examples/ex_read_nanovdb_sphere_accessor_cuda/read_nanovdb_sphere_accessor_cuda.cu @@ -3,8 +3,7 @@ //! [read_nanovdb_sphere_accessor_cuda] #include // this is required to read (and write) NanoVDB files on the host -#include // required for CUDA memory management -#include +#include // for cuda::copyTo, the explicit host<->device grid transfer extern "C" void launch_kernels(const nanovdb::NanoGrid*, const nanovdb::NanoGrid*, @@ -16,23 +15,27 @@ extern "C" void launch_kernels(const nanovdb::NanoGrid*, int main(int, char**) { try { - // returns a GridHandle using CUDA for memory management. - auto handle = nanovdb::io::readGrid("data/sphere.nvdb"); + // read the grid into host memory (HostBuffer is the default buffer type) + auto handle = nanovdb::io::readGrid("data/sphere.nvdb"); - cudaStream_t stream; // Create a CUDA stream to allow for asynchronous copy of pinned CUDA memory. + cudaStream_t stream; // stream that orders the transfer and the kernels below cudaStreamCreate(&stream); + { + // Deep-copy the grid to the GPU: the copy is ordered on the stream, and the + // returned handle validates the transferred grid on the device. + auto deviceHandle = nanovdb::cuda::copyTo>(handle, stream); - handle.deviceUpload(stream, false); // Copy the NanoVDB grid to the GPU asynchronously + auto* cpuGrid = handle.grid(); // a (raw) pointer to the grid of value type float on the CPU + auto* deviceGrid = deviceHandle.deviceGrid(); // and its deep copy on the GPU - auto* cpuGrid = handle.grid(); // get a (raw) pointer to a NanoVDB grid of value type float on the CPU - auto* deviceGrid = handle.deviceGrid(); // get a (raw) pointer to a NanoVDB grid of value type float on the GPU + if (!deviceGrid || !cpuGrid) + throw std::runtime_error("GridHandle did not contain a grid with value type float"); - if (!deviceGrid || !cpuGrid) - throw std::runtime_error("GridHandle did not contain a grid with value type float"); + launch_kernels(deviceGrid, cpuGrid, stream); // print grid values on both the CPU and GPU - launch_kernels(deviceGrid, cpuGrid, stream); // Call a host method to print a grid values on both the CPU and GPU - - cudaStreamDestroy(stream); // Destroy the CUDA stream + cudaStreamSynchronize(stream); // the kernels must finish before the device handle (whose buffer frees on this stream) goes away + } + cudaStreamDestroy(stream); // safe: nothing outlives the stream now } catch (const std::exception& e) { std::cerr << "An exception occurred: \"" << e.what() << "\"" << std::endl; diff --git a/nanovdb/nanovdb/examples/ex_refine_nanovdb_cuda/refine_nanovdb_cuda.cpp b/nanovdb/nanovdb/examples/ex_refine_nanovdb_cuda/refine_nanovdb_cuda.cpp index 19795c5dfd..73f1db2b82 100644 --- a/nanovdb/nanovdb/examples/ex_refine_nanovdb_cuda/refine_nanovdb_cuda.cpp +++ b/nanovdb/nanovdb/examples/ex_refine_nanovdb_cuda/refine_nanovdb_cuda.cpp @@ -11,7 +11,7 @@ // the following files are from NanoVDB #include -#include +#include // host-includable: cuda::copyTo transfers grids without any kernel #include template @@ -90,7 +90,7 @@ int main(int argc, char *argv[]) // Convert to indexGrid (original, un-refined) cpuTimer.start("Converting openVDB input to indexGrid (original version)"); - auto handleOriginal = nanovdb::tools::openToIndexVDB( + auto handleOriginal = nanovdb::tools::openToIndexVDB( grid, 0u, // Don't copy data channel false, // No stats @@ -154,7 +154,7 @@ int main(int argc, char *argv[]) // Convert to indexGrid (refined) cpuTimer.start("Converting openVDB input to indexGrid (refineed version)"); - auto handleRefined = nanovdb::tools::openToIndexVDB( + auto handleRefined = nanovdb::tools::openToIndexVDB( refinedGrid, 0u, // Don't copy data channel false, // No stats @@ -182,10 +182,11 @@ int main(int argc, char *argv[]) } // Copy both NanoVDB grids to GPU - handleOriginal.deviceUpload(); - handleRefined.deviceUpload(); - auto* deviceGridOriginal = handleOriginal.deviceGrid(); - auto* deviceGridRefined = handleRefined.deviceGrid(); + // deep-copy both grids to the device; the returned handles validate them there + auto deviceHandleOriginal = nanovdb::cuda::copyTo>(handleOriginal); + auto deviceHandleRefined = nanovdb::cuda::copyTo>(handleRefined); + auto* deviceGridOriginal = deviceHandleOriginal.deviceGrid(); + auto* deviceGridRefined = deviceHandleRefined.deviceGrid(); if (!deviceGridOriginal || !deviceGridRefined) OPENVDB_THROW(openvdb::RuntimeError, "Failure while uploading indexGrids to GPU"); diff --git a/nanovdb/nanovdb/examples/ex_refine_nanovdb_cuda/refine_nanovdb_cuda_kernels.cu b/nanovdb/nanovdb/examples/ex_refine_nanovdb_cuda/refine_nanovdb_cuda_kernels.cu index f590f0cf19..6cd03cf2b7 100644 --- a/nanovdb/nanovdb/examples/ex_refine_nanovdb_cuda/refine_nanovdb_cuda_kernels.cu +++ b/nanovdb/nanovdb/examples/ex_refine_nanovdb_cuda/refine_nanovdb_cuda_kernels.cu @@ -2,6 +2,7 @@ // SPDX-License-Identifier: Apache-2.0 #include +#include // for cuda::copyTo template bool bufferCheck(const T* deviceBuffer, const T* hostBuffer, size_t elem_count) { @@ -54,3 +55,4 @@ void mainRefineGrid( nanovdb::NanoGrid *indexGridRefined, uint32_t benchmark_iters ); + diff --git a/nanovdb/nanovdb/examples/ex_voxels_to_grid_cuda/ex_voxels_to_grid_cuda.cu b/nanovdb/nanovdb/examples/ex_voxels_to_grid_cuda/ex_voxels_to_grid_cuda.cu index 21d9a1f496..dc4128cf47 100644 --- a/nanovdb/nanovdb/examples/ex_voxels_to_grid_cuda/ex_voxels_to_grid_cuda.cu +++ b/nanovdb/nanovdb/examples/ex_voxels_to_grid_cuda/ex_voxels_to_grid_cuda.cu @@ -3,6 +3,7 @@ #include #include +#include // for cuda::copyTo, the explicit device->host grid transfer /// @brief Demonstrates how to create a NanoVDB grid from voxel coordinates on the GPU int main() @@ -10,20 +11,22 @@ int main() try { // Define list of voxel coordinates and copy them to the device const size_t numVoxels = 3; - nanovdb::Coord coords[numVoxels] = {nanovdb::Coord(1, 2, 3), nanovdb::Coord(-1,3,6), nanovdb::Coord(-90,100,5678)}, *d_coords = nullptr; - cudaCheck(cudaMalloc(&d_coords, numVoxels * sizeof(nanovdb::Coord))); + nanovdb::Coord coords[numVoxels] = {nanovdb::Coord(1, 2, 3), nanovdb::Coord(-1,3,6), nanovdb::Coord(-90,100,5678)}; + nanovdb::cuda::Buffer coordBuffer(cudaStream_t(0), numVoxels, nanovdb::cuda::noInit); + nanovdb::Coord *d_coords = coordBuffer.data(); cudaCheck(cudaMemcpy(d_coords, coords, numVoxels * sizeof(nanovdb::Coord), cudaMemcpyHostToDevice));// coords CPU -> GPU - // Generate a NanoVDB grid that contains the list of voxels on the device - auto handle = nanovdb::tools::cuda::voxelsToGrid(d_coords, numVoxels); + // Generate a NanoVDB grid from the voxels, stored in a single-space device buffer + auto handle = nanovdb::tools::cuda::voxelsToGrid>(d_coords, numVoxels); auto *d_grid = handle.deviceGrid(); // Define a list of values and copy them to the device - float values[numVoxels] = {1.4f, 6.7f, -5.0f}, *d_values; - cudaCheck(cudaMalloc(&d_values, numVoxels * sizeof(float))); + float values[numVoxels] = {1.4f, 6.7f, -5.0f}; + nanovdb::cuda::Buffer valueBuffer(cudaStream_t(0), numVoxels, nanovdb::cuda::noInit); + float *d_values = valueBuffer.data(); cudaCheck(cudaMemcpy(d_values, values, numVoxels * sizeof(float), cudaMemcpyHostToDevice));// values CPU -> GPU - // Launch a device kernel that sets the values of voxels define above and prints them + // Launch a device kernel that sets the values of the voxels defined above and prints them const unsigned int numThreads = 128, numBlocks = nanovdb::util::cuda::blocksPerGrid(numVoxels, numThreads); nanovdb::util::cuda::lambdaKernel<<>>(numVoxels, [=] __device__(size_t tid) { using OpT = nanovdb::SetVoxel;// defines type of random-access operation (set value) @@ -32,17 +35,14 @@ int main() printf("GPU: voxel # %zu, grid(%4i,%4i,%4i) = %5.1f\n", tid, ijk[0], ijk[1], ijk[2], d_grid->tree().getValue(ijk)); }); cudaCheckError(); - // Copy grid from GPU to CPU and print the voxel values for validation - handle.deviceDownload();// creates a copy on the CPU - auto *grid = handle.grid(); + // Deep-copy the grid to a host handle and print the voxel values for validation + auto hostHandle = nanovdb::cuda::copyTo(handle); + auto *grid = hostHandle.grid(); for (size_t i=0; itree().getValue(ijk)); } - // free arrays allocated on the device - cudaCheck(cudaFree(d_coords)); - cudaCheck(cudaFree(d_values)); } catch (const std::exception& e) { std::cerr << "An exception occurred: \"" << e.what() << "\"" << std::endl; diff --git a/nanovdb/nanovdb/python/PyGridValidator.cc b/nanovdb/nanovdb/python/PyGridValidator.cc index db9a46fc8f..64e4214066 100644 --- a/nanovdb/nanovdb/python/PyGridValidator.cc +++ b/nanovdb/nanovdb/python/PyGridValidator.cc @@ -32,7 +32,7 @@ template void defineValidateGrids(nb::module_& m) template void defineValidateGrids(nb::module_&); #ifdef NANOVDB_USE_CUDA -template void defineValidateGrids(nb::module_&); +template void defineValidateGrids(nb::module_&); #endif namespace { @@ -117,7 +117,7 @@ void defineGridValidatorModule(nb::module_& toolsModule) "whole handle."); #ifdef NANOVDB_USE_CUDA toolsModule.def("validateGrid", - &tools::validateGrid>, + &tools::validateGrid>, "handle"_a, "gridID"_a, "mode"_a = CheckMode::Default, "verbose"_a = false, nb::call_guard(), diff --git a/nanovdb/nanovdb/python/PyIO.cc b/nanovdb/nanovdb/python/PyIO.cc index 4573b93d02..61df4fe403 100644 --- a/nanovdb/nanovdb/python/PyIO.cc +++ b/nanovdb/nanovdb/python/PyIO.cc @@ -143,7 +143,7 @@ void defineHostReadWriteGrid(nb::module_& m) #ifdef NANOVDB_USE_CUDA void defineDeviceReadWriteGrid(nb::module_& m) { - using BufferT = cuda::DeviceBuffer; + using BufferT = cuda::DualDeviceBuffer; defineReadWriteGrid(m); m.def("deviceWriteGrid", diff --git a/nanovdb/nanovdb/python/PyPrimitives.cc b/nanovdb/nanovdb/python/PyPrimitives.cc index e061435af9..3ea1dc0aa7 100644 --- a/nanovdb/nanovdb/python/PyPrimitives.cc +++ b/nanovdb/nanovdb/python/PyPrimitives.cc @@ -531,7 +531,7 @@ template void definePrimitives(nb::module_& m) template void definePrimitives(nb::module_&); #ifdef NANOVDB_USE_CUDA -template void definePrimitives(nb::module_&); +template void definePrimitives(nb::module_&); #endif } // namespace pynanovdb diff --git a/nanovdb/nanovdb/python/PyTools.cc b/nanovdb/nanovdb/python/PyTools.cc index 9aa8f93452..63dbeb5415 100644 --- a/nanovdb/nanovdb/python/PyTools.cc +++ b/nanovdb/nanovdb/python/PyTools.cc @@ -53,9 +53,9 @@ void defineToolsModule(nb::module_& m) nb::module_ cudaModule = m.def_submodule("cuda"); cudaModule.doc() = "A submodule that implements CUDA-accelerated tools"; - defineValidateGrids(m); + defineValidateGrids(m); - definePrimitives(cudaModule); + definePrimitives(cudaModule); defineSignedFloodFill(cudaModule, "signedFloodFill"); defineSignedFloodFill(cudaModule, "signedFloodFill"); diff --git a/nanovdb/nanovdb/python/cuda/PyDeviceBuffer.cc b/nanovdb/nanovdb/python/cuda/PyDeviceBuffer.cc index e024d63939..6f2c9a67ef 100644 --- a/nanovdb/nanovdb/python/cuda/PyDeviceBuffer.cc +++ b/nanovdb/nanovdb/python/cuda/PyDeviceBuffer.cc @@ -13,7 +13,7 @@ namespace pynanovdb { void defineDeviceBuffer(nb::module_& m) { - nb::class_(m, "DeviceBuffer", + nb::class_(m, "DeviceBuffer", "CUDA device-side buffer used to back a DeviceGridHandle. Holds a " "host mirror and a device pointer; deviceUpload / deviceDownload on " "the handle move bytes between the two."); diff --git a/nanovdb/nanovdb/python/cuda/PyDeviceGridHandle.cu b/nanovdb/nanovdb/python/cuda/PyDeviceGridHandle.cu index 0214474420..a53c799ab7 100644 --- a/nanovdb/nanovdb/python/cuda/PyDeviceGridHandle.cu +++ b/nanovdb/nanovdb/python/cuda/PyDeviceGridHandle.cu @@ -22,7 +22,7 @@ namespace pynanovdb { // or the BuildT is not Python-visible. static nb::object pyDeviceGrid(nb::handle py_handle, uint32_t n) { - using BufferT = nanovdb::cuda::DeviceBuffer; + using BufferT = nanovdb::cuda::DualDeviceBuffer; auto& handle = nb::cast&>(py_handle); if (n >= handle.gridCount()) return nb::none(); switch (handle.gridType(n)) { @@ -58,7 +58,7 @@ static nb::object pyDeviceGrid(nb::handle py_handle, uint32_t n) void defineDeviceGridHandle(nb::module_& m) { - using BufferT = nanovdb::cuda::DeviceBuffer; + using BufferT = nanovdb::cuda::DualDeviceBuffer; defineGridHandle(m, "DeviceGridHandle") .def( "__init__", diff --git a/nanovdb/nanovdb/tools/VoxelBlockManager.h b/nanovdb/nanovdb/tools/VoxelBlockManager.h index 15760523c3..f1f730753f 100644 --- a/nanovdb/nanovdb/tools/VoxelBlockManager.h +++ b/nanovdb/nanovdb/tools/VoxelBlockManager.h @@ -181,6 +181,23 @@ class VoxelBlockManagerHandle typename util::enable_if::hasDeviceDual, const uint64_t*>::type deviceJumpMap() const { return static_cast(mJumpMap.deviceData()); } + //@{ + /// @brief For a single-space buffer the device data is the buffer itself. + /// @warning Note that the return pointer can be NULL if the VoxelBlockManagerHandle was not initialized + template + typename util::enable_if::value, uint32_t*>::type + deviceFirstLeafID() { return reinterpret_cast(mFirstLeafID.data()); } + template + typename util::enable_if::value, const uint32_t*>::type + deviceFirstLeafID() const { return reinterpret_cast(mFirstLeafID.data()); } + template + typename util::enable_if::value, uint64_t*>::type + deviceJumpMap() { return reinterpret_cast(mJumpMap.data()); } + template + typename util::enable_if::value, const uint64_t*>::type + deviceJumpMap() const { return reinterpret_cast(mJumpMap.data()); } + //@} + /// @brief Returns the number of voxel blocks in the VoxelBlockManager uint64_t blockCount() const { return mBlockCount; } diff --git a/nanovdb/nanovdb/tools/cuda/AddBlindData.cuh b/nanovdb/nanovdb/tools/cuda/AddBlindData.cuh index 13507dd5b3..21c221c2b6 100644 --- a/nanovdb/nanovdb/tools/cuda/AddBlindData.cuh +++ b/nanovdb/nanovdb/tools/cuda/AddBlindData.cuh @@ -26,6 +26,7 @@ #include #include // for std::strcpy +#include namespace nanovdb {// ================================================ @@ -45,7 +46,7 @@ namespace tools::cuda {// ============================================ /// @param pool optional pool used for allocation /// @param stream optional CUDA stream (defaults to CUDA stream 0) /// @return GridHandle with blind data appended -template +template GridHandle addBlindData(const NanoGrid *d_grid, const BlindDataT *d_blindData, @@ -61,7 +62,6 @@ addBlindData(const NanoGrid *d_grid, // Out: |-----------|----------|----------|-----------|------------| // old grid old meta new meta old data new data - static_assert(BufferTraits::hasDeviceDual, "Expected BufferT to support device allocation"); static_assert(nanovdb::cuda::is_async_resource::value, "addBlindData allocates stream-ordered scratch and requires an AsyncResource"); @@ -89,8 +89,9 @@ addBlindData(const NanoGrid *d_grid, sizeof(BlindDataT), semantics, blindClass, toGridType()}; if (!metaData.isValid()) throw std::runtime_error("cudaAddBlindData: invalid combination of blind meta data"); std::strcpy(metaData.mName, name); - auto buffer = BufferT::create(tmp[GRID] + tmp[META] + sizeof(GridBlindMetaData) + tmp[DATA] + metaData.blindDataSize(), &pool, false); - void *d_data = buffer.deviceData(); + auto buffer = nanovdb::cuda::detail::createDeviceStorage(tmp[GRID] + tmp[META] + sizeof(GridBlindMetaData) + tmp[DATA] + metaData.blindDataSize(), + &pool, util::cuda::currentDevice(), stream); + void *d_data = nanovdb::cuda::detail::deviceStorageData(buffer); // 1: |-----------|----------| // old grid old meta @@ -130,12 +131,13 @@ addBlindData(const NanoGrid *d_grid, Checksum cs(tmp[CHECKSUM]); cuda::updateChecksum(reinterpret_cast(d_data), cs.mode(), stream); + nanovdb::cuda::detail::orderBeforeHandleConstruction(stream); return GridHandle(std::move(buffer)); }// cudaAddBlindData }// namespace tools::cuda -template +template [[deprecated("Use nanovdb::cuda::addBlindData instead")]] GridHandle cudaAddBlindData(const NanoGrid *d_grid, diff --git a/nanovdb/nanovdb/tools/cuda/CoarsenGrid.cuh b/nanovdb/nanovdb/tools/cuda/CoarsenGrid.cuh index 0f4cebd15c..cca0ac00ac 100644 --- a/nanovdb/nanovdb/tools/cuda/CoarsenGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/CoarsenGrid.cuh @@ -64,7 +64,7 @@ public: /// @tparam BufferT Buffer type used for allocation of the grid handle /// @param buffer optional buffer (currently ignored) /// @return returns a handle with a grid of type NanoGrid - template + template GridHandle getHandle(const BufferT &buffer = BufferT()); @@ -196,13 +196,12 @@ void CoarsenGrid::coarsenRoot() // Package the new root topology into a RootNode plus Tile list; upload to the GPU uint64_t rootSize = RootT::memUsage(coarsenedTiles.size()); - mBuilder.mProcessedRoot = nanovdb::cuda::DeviceBuffer::create(rootSize); - auto coarsenedRootPtr = static_cast(mBuilder.mProcessedRoot.data()); + auto coarsenedRootPtr = mBuilder.allocateProcessedRoot(rootSize); coarsenedRootPtr->mTableSize = coarsenedTiles.size(); uint32_t t = 0; for (const auto& [key, tile] : coarsenedTiles) *coarsenedRootPtr->tile(t++) = tile; - mBuilder.mProcessedRoot.deviceUpload(device, mStream, false); + mBuilder.uploadProcessedRoot(mStream); }// CoarsenGrid::coarsenRoot //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- diff --git a/nanovdb/nanovdb/tools/cuda/DilateGrid.cuh b/nanovdb/nanovdb/tools/cuda/DilateGrid.cuh index 5c1bb36e5e..3244fc964d 100644 --- a/nanovdb/nanovdb/tools/cuda/DilateGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/DilateGrid.cuh @@ -68,7 +68,7 @@ public: /// @tparam BufferT Buffer type used for allocation of the grid handle /// @param buffer optional buffer (currently ignored) /// @return returns a handle with a grid of type NanoGrid - template + template GridHandle getHandle(const BufferT &buffer = BufferT()); @@ -219,13 +219,12 @@ void DilateGrid::dilateRoot() // Package the new root topology into a RootNode plus Tile list; upload to the GPU uint64_t rootSize = RootT::memUsage(dilatedTiles.size()); - mBuilder.mProcessedRoot = nanovdb::cuda::DeviceBuffer::create(rootSize); - auto dilatedRootPtr = static_cast(mBuilder.mProcessedRoot.data()); + auto dilatedRootPtr = mBuilder.allocateProcessedRoot(rootSize); dilatedRootPtr->mTableSize = dilatedTiles.size(); uint32_t t = 0; for (const auto& [key, tile] : dilatedTiles) *dilatedRootPtr->tile(t++) = tile; - mBuilder.mProcessedRoot.deviceUpload(device, mStream, false); + mBuilder.uploadProcessedRoot(mStream); }// DilateGrid::dilateRoot //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- diff --git a/nanovdb/nanovdb/tools/cuda/GridChecksum.cuh b/nanovdb/nanovdb/tools/cuda/GridChecksum.cuh index 100a405416..825c4b6fe9 100644 --- a/nanovdb/nanovdb/tools/cuda/GridChecksum.cuh +++ b/nanovdb/nanovdb/tools/cuda/GridChecksum.cuh @@ -288,7 +288,7 @@ template void crc32TailOld(const NanoGrid *d_grid, const GridData *gridData, const uint32_t *d_lut, uint32_t *d_crc, cudaStream_t stream) { static constexpr unsigned int threadsPerBlock = 128;// seems faster than the old value of 256! - auto nodeMgrHandle = nanovdb::cuda::createNodeManager(d_grid, nanovdb::cuda::DeviceBuffer(), stream); + auto nodeMgrHandle = nanovdb::cuda::createNodeManager(d_grid, nanovdb::cuda::DualDeviceBuffer(), stream); auto *d_nodeMgr = nodeMgrHandle.template deviceMgr(); NANOVDB_ASSERT(isAligned(d_nodeMgr)); const uint32_t nodeCount[3]={gridData->template nodeCount<0>(), gridData->template nodeCount<1>(), gridData->template nodeCount<2>()}; diff --git a/nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh b/nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh index 2bf1b44e0f..ea2949b9a5 100644 --- a/nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh @@ -23,6 +23,7 @@ #include #include #include +#include namespace nanovdb {// ================================================================ @@ -43,12 +44,12 @@ namespace tools::cuda {// ====================================================== /// NanoRoot::FloatType, e.g. if DstBuildT=Vec3f then NanoRoot::FloatType=float, /// in which case average and standard-deviation is undefined in the output grid. /// @return returns handle to grid that combined IndexGrid and values -template +template typename util::enable_if::is_index, GridHandle>::type indexToGrid(const NanoGrid *d_srcGrid, const typename BuildToValueMap::type *d_srcValues, const BufferT &pool = BufferT(), cudaStream_t stream = 0); -template +template typename util::enable_if::is_index, GridHandle>::type createNanoGrid(const NanoGrid *d_srcGrid, const typename BuildToValueMap::type *d_srcValues, const BufferT &pool = BufferT(), cudaStream_t stream = 0) { @@ -90,7 +91,7 @@ public: /// @param srcValues pointer to values that will be inserted into the output grid /// @param buffer optional buffer used for memory allocation /// @return A new GridHandle with the grid of type @c DstBuildT - template + template GridHandle getHandle(const typename BuildToValueMap::type *srcValues, const BufferT &buffer = BufferT()); private: @@ -379,7 +380,7 @@ GridHandle IndexToGrid::getHandle(const typename updateChecksum((GridData*)mNodeAcc.d_dstPtr, mStream); if (mVerbose) mTimer.stop(); - //cudaStreamSynchronize(mStream);// finish all device tasks in mStream + nanovdb::cuda::detail::orderBeforeHandleConstruction(mStream); return GridHandle(std::move(buffer)); }// IndexToGrid::getHandle @@ -400,8 +401,8 @@ inline BufferT IndexToGrid::getBuffer(const BufferT &pool) mNodeAcc.size = mNodeAcc.blind;// end of buffer int device = 0; cudaCheck(cudaGetDevice(&device)); - auto buffer = BufferT::create(mNodeAcc.size, &pool, device, mStream); - mNodeAcc.d_dstPtr = buffer.deviceData(); + auto buffer = nanovdb::cuda::detail::createDeviceStorage(mNodeAcc.size, &pool, device, mStream); + mNodeAcc.d_dstPtr = nanovdb::cuda::detail::deviceStorageData(buffer); 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 @@ -436,7 +437,7 @@ indexToGrid(const NanoGrid *d_srcGrid, const typename BuildToValueMap }// namespace tools::cuda ============================================================= -template +template [[deprecated("Use nanovdb::cuda::indexToGrid instead")]] typename util::enable_if::is_index, GridHandle>::type cudaIndexToGrid(const NanoGrid *d_srcGrid, const typename BuildToValueMap::type *d_srcValues, const BufferT &pool = BufferT(), cudaStream_t stream = 0) @@ -445,7 +446,7 @@ cudaIndexToGrid(const NanoGrid *d_srcGrid, const typename BuildToValu } -template +template [[deprecated("Use nanovdb::cuda::indexToGrid instead")]] typename util::enable_if::is_index, GridHandle>::type cudaCreateNanoGrid(const NanoGrid *d_srcGrid, const typename BuildToValueMap::type *d_srcValues, const BufferT &pool = BufferT(), cudaStream_t stream = 0) diff --git a/nanovdb/nanovdb/tools/cuda/MergeGrids.cuh b/nanovdb/nanovdb/tools/cuda/MergeGrids.cuh index 70149db567..6ac712be93 100644 --- a/nanovdb/nanovdb/tools/cuda/MergeGrids.cuh +++ b/nanovdb/nanovdb/tools/cuda/MergeGrids.cuh @@ -77,7 +77,7 @@ public: /// @tparam BufferT Buffer type used for allocation of the grid handle /// @param buffer optional buffer (currently ignored) /// @return returns a handle with a grid of type NanoGrid - template + template GridHandle getHandle(const BufferT &buffer = BufferT()); @@ -220,13 +220,12 @@ void MergeGrids::mergeRoot() // Package the new root topology into a RootNode plus Tile list; upload to the GPU uint64_t rootSize = RootT::memUsage(mergedTiles.size()); - mBuilder.mProcessedRoot = nanovdb::cuda::DeviceBuffer::create(rootSize); - auto mergedRootPtr = static_cast(mBuilder.mProcessedRoot.data()); + auto mergedRootPtr = mBuilder.allocateProcessedRoot(rootSize); mergedRootPtr->mTableSize = mergedTiles.size(); uint32_t t = 0; for (const auto& [key, tile] : mergedTiles) *mergedRootPtr->tile(t++) = tile; - mBuilder.mProcessedRoot.deviceUpload(device, mStream, false); + mBuilder.uploadProcessedRoot(mStream); }// MergeGrids::mergeRoot //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- diff --git a/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh b/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh index ee37375627..4962392b7b 100644 --- a/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh @@ -64,6 +64,10 @@ class MeshToGrid "MeshToGrid allocates stream-ordered scratch and requires an AsyncResource"); using PointT = nanovdb::Vec3f; + using ScratchT = nanovdb::cuda::Buffer>; + + nanovdb::cuda::ResourceRef ref() { return mBuilder.ref(); } + using TriangleIndexT = nanovdb::Vec3i; using TriangleT = Triangle; using GridT = NanoGrid; @@ -92,7 +96,10 @@ public: ResourceT& resource = nanovdb::cuda::default_resource() ) : mStream(stream), mTimer(stream), mBuilder(stream, resource), mDevicePoints(devicePoints), mPointCount(pointCount), - mDeviceTriangles(deviceTriangles), mTriangleCount(triangleCount), mMap(map), mTempDevicePool(resource) + mDeviceTriangles(deviceTriangles), mTriangleCount(triangleCount), mMap(map), mTempDevicePool(resource), + mXformedTriangles(stream, nanovdb::cuda::ResourceRef(resource), 0, nanovdb::cuda::noInit), + mBoxTrianglePairsBuffer(stream, nanovdb::cuda::ResourceRef(resource), 0, nanovdb::cuda::noInit), + mUniqueRootOriginsBuffer(stream, nanovdb::cuda::ResourceRef(resource), 0, nanovdb::cuda::noInit) {} /// @brief Toggle on and off verbose mode @@ -121,7 +128,7 @@ public: /// @tparam BufferT Buffer type used for allocation of the grid handle /// @param buffer optional buffer (currently ignored) /// @return returns a handle with a grid of type NanoGrid - template + template GridHandle getHandle(const BufferT &buffer = BufferT()); @@ -138,8 +145,8 @@ public: /// @param buffer optional allocator for the grid handle (currently ignored) /// @param sidecarBuffer optional allocator for the UDF sidecar (currently ignored) /// @return std::pair of grid handle and UDF sidecar buffer - template + template std::pair, SidecarBufferT> getHandleAndUDF(const GridBufferT& buffer = GridBufferT(), const SidecarBufferT& sidecarBuffer = SidecarBufferT()); @@ -178,15 +185,16 @@ private: const uint32_t mTriangleCount; const nanovdb::Map mMap; - nanovdb::cuda::DeviceBuffer mXformedTriangles; - nanovdb::cuda::DeviceBuffer mBoxTrianglePairsBuffer; + ScratchT mXformedTriangles; + ScratchT mBoxTrianglePairsBuffer; uint64_t mBoxTrianglePairCount{0}; - nanovdb::cuda::DeviceBuffer mUniqueRootOriginsBuffer; + ScratchT mUniqueRootOriginsBuffer; uint64_t mUniqueRootTileCount{0}; - auto deviceXformedTriangles() { return static_cast(mXformedTriangles.deviceData()); } - auto deviceBoxTrianglePairs() { return static_cast(mBoxTrianglePairsBuffer.deviceData()); } - auto deviceUniqueRootOrigins() const { return static_cast(mUniqueRootOriginsBuffer.deviceData()); } + auto deviceXformedTriangles() { return reinterpret_cast(mXformedTriangles.data()); } + auto deviceBoxTrianglePairs() { return reinterpret_cast(mBoxTrianglePairsBuffer.data()); } + auto deviceUniqueRootOrigins() { return reinterpret_cast(mUniqueRootOriginsBuffer.data()); } + auto deviceUniqueRootOrigins() const { return reinterpret_cast(mUniqueRootOriginsBuffer.data()); } nanovdb::cuda::TempPool mTempDevicePool; }; // tools::cuda::MeshToGrid @@ -282,8 +290,8 @@ GridHandle MeshToGrid::getHandle(const BufferT &buff rasterizeLeafNodes(); if (mVerbose==1) mTimer.stop(); if (mBoxTrianglePairCount) { - mXformedTriangles.clear(mStream); - mBoxTrianglePairsBuffer.clear(mStream); + mXformedTriangles.destroy(mStream); + mBoxTrianglePairsBuffer.destroy(mStream); } // Update leaf value offsets (prefix sums of per-leaf active voxel counts) @@ -307,13 +315,12 @@ GridHandle MeshToGrid::getHandle(const BufferT &buff const uint32_t leafCount = mBuilder.data()->nodeCount[0]; auto handle = GridHandle(std::move(gridBuffer)); if (leafCount) { - nanovdb::cuda::DeviceBuffer retainMaskBuffer = nanovdb::cuda::DeviceBuffer::create( - uint64_t(leafCount) * sizeof(nanovdb::Mask<3>), nullptr, device, mStream); - cudaCheck(cudaMemsetAsync(retainMaskBuffer.deviceData(), 0xFF, + ScratchT retainMaskBuffer = ScratchT(mStream, this->ref(), uint64_t(leafCount) * sizeof(nanovdb::Mask<3>), nanovdb::cuda::noInit); + cudaCheck(cudaMemsetAsync(retainMaskBuffer.data(), 0xFF, uint64_t(leafCount) * sizeof(nanovdb::Mask<3>), mStream)); tools::cuda::PruneGrid pruner( static_cast(handle.deviceData()), - static_cast*>(retainMaskBuffer.deviceData()), + reinterpret_cast*>(retainMaskBuffer.data()), mStream); handle = pruner.template getHandle(buffer); } @@ -352,8 +359,8 @@ void MeshToGrid::transformTriangles() int device = 0; cudaGetDevice(&device); - mXformedTriangles = nanovdb::cuda::DeviceBuffer::create(mTriangleCount*sizeof(TriangleT), nullptr, device, mStream); - if (mXformedTriangles.deviceData() == nullptr) throw std::runtime_error("Failed to allocate transofmed upper mask buffer on device"); + mXformedTriangles = ScratchT(mStream, this->ref(), mTriangleCount*sizeof(TriangleT), nanovdb::cuda::noInit); + if (mXformedTriangles.data() == nullptr) throw std::runtime_error("Failed to allocate transofmed upper mask buffer on device"); util::cuda::lambdaKernel<<>>( mTriangleCount, @@ -490,15 +497,15 @@ void MeshToGrid::processRootTrianglePairs() // Pass 1: Count intersecting root boxes per triangle - nanovdb::cuda::DeviceBuffer - rootBoxCounts = nanovdb::cuda::DeviceBuffer::create(mTriangleCount * sizeof(uint64_t), nullptr, device, mStream); - if (rootBoxCounts.deviceData() == nullptr) throw std::runtime_error("Failed to allocate root box counts buffer"); + ScratchT + rootBoxCounts = ScratchT(mStream, this->ref(), mTriangleCount * sizeof(uint64_t), nanovdb::cuda::noInit); + if (rootBoxCounts.data() == nullptr) throw std::runtime_error("Failed to allocate root box counts buffer"); util::cuda::lambdaKernel<<>>( mTriangleCount, topology::detail::CountRootBoxesFunctor{ deviceXformedTriangles(), - static_cast(rootBoxCounts.deviceData()), + reinterpret_cast(rootBoxCounts.data()), mBandWidth } ); @@ -506,29 +513,27 @@ void MeshToGrid::processRootTrianglePairs() // Pass 2: InclusiveSum Scan to compute offsets and total allocations - nanovdb::cuda::DeviceBuffer rootBoxOffsets = - nanovdb::cuda::DeviceBuffer::create((mTriangleCount+1)*sizeof(uint64_t), nullptr, device, mStream); - if (rootBoxOffsets.deviceData() == nullptr) throw std::runtime_error("Failed to allocate root box offsets buffer"); + ScratchT rootBoxOffsets = ScratchT(mStream, this->ref(), (mTriangleCount+1)*sizeof(uint64_t), nanovdb::cuda::noInit); + if (rootBoxOffsets.data() == nullptr) throw std::runtime_error("Failed to allocate root box offsets buffer"); - cudaCheck(cudaMemsetAsync(rootBoxOffsets.deviceData(), 0, sizeof(uint64_t), mStream)); + cudaCheck(cudaMemsetAsync(rootBoxOffsets.data(), 0, sizeof(uint64_t), mStream)); CALL_CUBS(DeviceScan::InclusiveSum, - static_cast(rootBoxCounts.deviceData()), - static_cast(rootBoxOffsets.deviceData())+1, + reinterpret_cast(rootBoxCounts.data()), + reinterpret_cast(rootBoxOffsets.data())+1, mTriangleCount); - cudaCheck(cudaMemcpyAsync(&mBoxTrianglePairCount, static_cast(rootBoxOffsets.deviceData())+mTriangleCount, sizeof(uint64_t), cudaMemcpyDeviceToHost, mStream)); + cudaCheck(cudaMemcpyAsync(&mBoxTrianglePairCount, reinterpret_cast(rootBoxOffsets.data())+mTriangleCount, sizeof(uint64_t), cudaMemcpyDeviceToHost, mStream)); cudaStreamSynchronize(mStream); // Pass 3: Re-enumerate intersections of (padded) root boxes and triangles, and scatter to allocated list - mBoxTrianglePairsBuffer = nanovdb::cuda::DeviceBuffer::create( - mBoxTrianglePairCount * sizeof(MeshToGridBoxTrianglePair), nullptr, device, mStream); - if (mBoxTrianglePairsBuffer.deviceData() == nullptr) throw std::runtime_error("Failed to allocate pairs buffer"); + mBoxTrianglePairsBuffer = ScratchT(mStream, this->ref(), mBoxTrianglePairCount * sizeof(MeshToGridBoxTrianglePair), nanovdb::cuda::noInit); + if (mBoxTrianglePairsBuffer.data() == nullptr) throw std::runtime_error("Failed to allocate pairs buffer"); util::cuda::lambdaKernel<<>>( mTriangleCount, topology::detail::ScatterRootTrianglePairsFunctor{ deviceXformedTriangles(), - static_cast(rootBoxOffsets.deviceData()), + reinterpret_cast(rootBoxOffsets.data()), deviceBoxTrianglePairs(), mBandWidth } @@ -781,9 +786,8 @@ void MeshToGrid::enumerateRootTiles() cudaGetDevice(&device); // Step 1: Encode each pair's root origin as a sortable uint64_t key - nanovdb::cuda::DeviceBuffer keysBuffer = nanovdb::cuda::DeviceBuffer::create( - mBoxTrianglePairCount * sizeof(uint64_t), nullptr, device, mStream); - auto *dKeys = static_cast(keysBuffer.deviceData()); + ScratchT keysBuffer = ScratchT(mStream, this->ref(), mBoxTrianglePairCount * sizeof(uint64_t), nanovdb::cuda::noInit); + auto *dKeys = reinterpret_cast(keysBuffer.data()); util::cuda::lambdaKernel<<>>( mBoxTrianglePairCount, @@ -792,20 +796,17 @@ void MeshToGrid::enumerateRootTiles() cudaCheckError(); // Step 2: Sort keys (SortKeys requires separate in/out buffers) - nanovdb::cuda::DeviceBuffer sortedKeysBuffer = nanovdb::cuda::DeviceBuffer::create( - mBoxTrianglePairCount * sizeof(uint64_t), nullptr, device, mStream); - auto *dSortedKeys = static_cast(sortedKeysBuffer.deviceData()); + ScratchT sortedKeysBuffer = ScratchT(mStream, this->ref(), mBoxTrianglePairCount * sizeof(uint64_t), nanovdb::cuda::noInit); + auto *dSortedKeys = reinterpret_cast(sortedKeysBuffer.data()); CALL_CUBS(DeviceRadixSort::SortKeys, dKeys, dSortedKeys, (int)mBoxTrianglePairCount, 0, 64); // Step 3: Select unique keys - nanovdb::cuda::DeviceBuffer uniqueKeysBuffer = nanovdb::cuda::DeviceBuffer::create( - mBoxTrianglePairCount * sizeof(uint64_t), nullptr, device, mStream); - auto *dUniqueKeys = static_cast(uniqueKeysBuffer.deviceData()); + ScratchT uniqueKeysBuffer = ScratchT(mStream, this->ref(), mBoxTrianglePairCount * sizeof(uint64_t), nanovdb::cuda::noInit); + auto *dUniqueKeys = reinterpret_cast(uniqueKeysBuffer.data()); - nanovdb::cuda::DeviceBuffer numSelectedBuffer = nanovdb::cuda::DeviceBuffer::create( - sizeof(int32_t), nullptr, device, mStream); - auto *dNumSelected = static_cast(numSelectedBuffer.deviceData()); + ScratchT numSelectedBuffer = ScratchT(mStream, this->ref(), sizeof(int32_t), nanovdb::cuda::noInit); + auto *dNumSelected = reinterpret_cast(numSelectedBuffer.data()); CALL_CUBS(DeviceSelect::Unique, dSortedKeys, dUniqueKeys, dNumSelected, (int)mBoxTrianglePairCount); @@ -815,8 +816,7 @@ void MeshToGrid::enumerateRootTiles() mUniqueRootTileCount = static_cast(uniqueCount); // Step 4: Decode unique keys back to Coord origins - mUniqueRootOriginsBuffer = nanovdb::cuda::DeviceBuffer::create( - mUniqueRootTileCount * sizeof(nanovdb::Coord), nullptr, device, mStream); + mUniqueRootOriginsBuffer = ScratchT(mStream, this->ref(), mUniqueRootTileCount * sizeof(nanovdb::Coord), nanovdb::cuda::noInit); auto *dOrigins = deviceUniqueRootOrigins(); util::cuda::lambdaKernel<<>>( @@ -844,8 +844,7 @@ void MeshToGrid::buildRasterizedRoot() // Only the NanoVDB tile key is set here; child pointers and values are // filled by TopologyBuilder's subsequent pipeline stages. uint64_t rootSize = RootT::memUsage(tileCount); - mBuilder.mProcessedRoot = nanovdb::cuda::DeviceBuffer::create(rootSize); - auto *rootPtr = static_cast(mBuilder.mProcessedRoot.data()); + auto *rootPtr = mBuilder.allocateProcessedRoot(rootSize); rootPtr->mTableSize = tileCount; rootPtr->mBackground = typename RootT::ValueType{}; @@ -856,8 +855,8 @@ void MeshToGrid::buildRasterizedRoot() tileCount * sizeof(nanovdb::Coord), cudaMemcpyDeviceToHost)); for (uint32_t t = 0; t < tileCount; ++t) *rootPtr->tile(t) = typename RootT::DataType::Tile{RootT::CoordToKey(hostOrigins[t])}; - mBuilder.mProcessedRoot.deviceUpload(device, mStream, false); - mUniqueRootOriginsBuffer.clear(mStream); + mBuilder.uploadProcessedRoot(mStream); + mUniqueRootOriginsBuffer.destroy(mStream); } } // MeshToGrid::buildRasterizedRoot @@ -937,21 +936,19 @@ void MeshToGrid::processLeafTrianglePairs() for (int pass = 0; pass < 3; ++pass) { // Allocate Mask<3> buffer for the CTA hit results // Size: mBoxTrianglePairCount * sizeof(nanovdb::Mask<3>) - nanovdb::cuda::DeviceBuffer maskBuffer = nanovdb::cuda::DeviceBuffer::create( - mBoxTrianglePairCount * sizeof(nanovdb::Mask<3>), nullptr, device, mStream); - if (maskBuffer.deviceData() == nullptr) { + ScratchT maskBuffer = ScratchT(mStream, this->ref(), mBoxTrianglePairCount * sizeof(nanovdb::Mask<3>), nanovdb::cuda::noInit); + if (maskBuffer.data() == nullptr) { throw std::runtime_error("Failed to allocate mask buffer for subdivision pass"); } - auto* dMasks = static_cast*>(maskBuffer.deviceData()); + auto* dMasks = reinterpret_cast*>(maskBuffer.data()); // Allocate Counts buffer for Prefix Sum // Size: mBoxTrianglePairCount * sizeof(uint64_t) - nanovdb::cuda::DeviceBuffer countsBuffer = nanovdb::cuda::DeviceBuffer::create( - mBoxTrianglePairCount * sizeof(uint64_t), nullptr, device, mStream); - if (countsBuffer.deviceData() == nullptr) { + ScratchT countsBuffer = ScratchT(mStream, this->ref(), mBoxTrianglePairCount * sizeof(uint64_t), nanovdb::cuda::noInit); + if (countsBuffer.data() == nullptr) { throw std::runtime_error("Failed to allocate counts buffer for subdivision pass"); } - auto* dCounts = static_cast(countsBuffer.deviceData()); + auto* dCounts = reinterpret_cast(countsBuffer.data()); // Evaluate & Count: 1 CTA per parent pair, 512 threads per CTA. // Uses AABB-only test for large child scales (>= mSATThreshold), full SAT below. @@ -971,11 +968,10 @@ void MeshToGrid::processLeafTrianglePairs() // Prefix Sum: element [i+1] = exclusive write offset for parent i's children, // element [0] = 0, element [mBoxTrianglePairCount] = total child pair count. - nanovdb::cuda::DeviceBuffer offsetsBuffer = nanovdb::cuda::DeviceBuffer::create( - (mBoxTrianglePairCount + 1) * sizeof(uint64_t), nullptr, device, mStream); - if (offsetsBuffer.deviceData() == nullptr) + ScratchT offsetsBuffer = ScratchT(mStream, this->ref(), (mBoxTrianglePairCount + 1) * sizeof(uint64_t), nanovdb::cuda::noInit); + if (offsetsBuffer.data() == nullptr) throw std::runtime_error("Failed to allocate offsets buffer for subdivision pass"); - auto* dOffsets = static_cast(offsetsBuffer.deviceData()); + auto* dOffsets = reinterpret_cast(offsetsBuffer.data()); cudaCheck(cudaMemsetAsync(dOffsets, 0, sizeof(uint64_t), mStream)); CALL_CUBS(DeviceScan::InclusiveSum, @@ -989,11 +985,10 @@ void MeshToGrid::processLeafTrianglePairs() cudaStreamSynchronize(mStream); // Allocate new child pair buffer - nanovdb::cuda::DeviceBuffer newPairsBuffer = nanovdb::cuda::DeviceBuffer::create( - newPairCount * sizeof(BoxTrianglePair), nullptr, device, mStream); - if (newPairsBuffer.deviceData() == nullptr) + ScratchT newPairsBuffer = ScratchT(mStream, this->ref(), newPairCount * sizeof(BoxTrianglePair), nanovdb::cuda::noInit); + if (newPairsBuffer.data() == nullptr) throw std::runtime_error("Failed to allocate child pairs buffer for subdivision pass"); - auto* dNewPairs = static_cast(newPairsBuffer.deviceData()); + auto* dNewPairs = reinterpret_cast(newPairsBuffer.data()); // Scatter surviving child pairs into the new buffer util::cuda::lambdaKernel<<>>( @@ -1053,7 +1048,7 @@ struct FinalizeSidecarFunctor template template std::pair, SidecarBufferT> -MeshToGrid::getHandleAndUDF(const GridBufferT& buffer, const SidecarBufferT&) +MeshToGrid::getHandleAndUDF(const GridBufferT& buffer, const SidecarBufferT& sidecarProto) { cudaStreamSynchronize(mStream); @@ -1130,13 +1125,12 @@ MeshToGrid::getHandleAndUDF(const GridBufferT& buffer, const const uint32_t leafCount = mBuilder.data()->nodeCount[0]; auto handle = GridHandle(std::move(gridBuffer)); if (leafCount) { - nanovdb::cuda::DeviceBuffer retainMaskBuffer = nanovdb::cuda::DeviceBuffer::create( - uint64_t(leafCount) * sizeof(nanovdb::Mask<3>), nullptr, device, mStream); - cudaCheck(cudaMemsetAsync(retainMaskBuffer.deviceData(), 0xFF, + ScratchT retainMaskBuffer = ScratchT(mStream, this->ref(), uint64_t(leafCount) * sizeof(nanovdb::Mask<3>), nanovdb::cuda::noInit); + cudaCheck(cudaMemsetAsync(retainMaskBuffer.data(), 0xFF, uint64_t(leafCount) * sizeof(nanovdb::Mask<3>), mStream)); tools::cuda::PruneGrid pruner( static_cast(handle.deviceData()), - static_cast*>(retainMaskBuffer.deviceData()), + reinterpret_cast*>(retainMaskBuffer.data()), mStream); handle = pruner.template getHandle(buffer); } @@ -1149,9 +1143,9 @@ MeshToGrid::getHandleAndUDF(const GridBufferT& buffer, const const uint64_t activeVoxelCount = util::cuda::DeviceGridTraits::getActiveVoxelCount( handle.template deviceGrid()); - auto sidecarBuffer = nanovdb::cuda::DeviceBuffer::create( - (activeVoxelCount + 1) * sizeof(float), nullptr, device, mStream); - auto *dSidecar = static_cast(sidecarBuffer.deviceData()); + auto sidecarBuffer = nanovdb::cuda::detail::createDeviceStorage( + (activeVoxelCount + 1) * sizeof(float), &sidecarProto, device, mStream); + auto *dSidecar = static_cast(nanovdb::cuda::detail::deviceStorageData(sidecarBuffer)); if (mVerbose==1) mTimer.start("Initializing UDF sidecar"); util::cuda::lambdaKernel<<>>( @@ -1169,8 +1163,8 @@ MeshToGrid::getHandleAndUDF(const GridBufferT& buffer, const handle.template deviceGrid(), dSidecar, mBandWidth * mBandWidth }); cudaCheckError(); - mXformedTriangles.clear(mStream); - mBoxTrianglePairsBuffer.clear(mStream); + mXformedTriangles.destroy(mStream); + mBoxTrianglePairsBuffer.destroy(mStream); } if (mVerbose==1) mTimer.stop(); diff --git a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh index 20011a3ba9..62e2932bf2 100644 --- a/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh @@ -28,6 +28,7 @@ #include #include #include +#include #include #include #include @@ -54,7 +55,7 @@ namespace tools::cuda {// ====================================================== /// @param stream optional CUDA stream (defaults to CUDA stream 0) /// @return Returns a handle with a grid of type NanoGrid where point information, e.g. coordinates, /// are represented as blind data defined by @c type. -template +template GridHandle pointsToGrid(const PtrT dWorldPoints, int pointCount, @@ -83,7 +84,7 @@ pointsToGrid(const PtrT dWorldPoints, /// @param stream optional CUDA stream (defaults to CUDA stream 0) /// @return Returns a handle with a grid of type NanoGrid where point information, e.g. coordinates, /// are represented as blind data defined by @c type. -template +template GridHandle pointsToGrid(const PtrT dWorldPoints, int pointCount, @@ -96,7 +97,7 @@ pointsToGrid(const PtrT dWorldPoints, //----------------------------------------------------------------------------------------------------- -template +template GridHandle pointsToGrid(std::vector> pointSet, const BufferT &buffer = BufferT(), @@ -116,7 +117,7 @@ pointsToGrid(std::vector> pointSe /// @param voxelSize Size of a voxel in world units used for the output grid /// @param buffer Instance of the device buffer used for memory allocation /// @return Returns a handle with the grid of type NanoGrid -template +template GridHandle voxelsToGrid(const PtrT dGridVoxels, size_t voxelCount, @@ -126,7 +127,7 @@ voxelsToGrid(const PtrT dGridVoxels, //------------------------------------------------------------------------------------------------------- -template +template GridHandle voxelsToGrid(std::vector> pointSet, const BufferT &buffer = BufferT(), @@ -365,7 +366,7 @@ public: /// @param pointCount number of input points or voxels /// @param buffer optional buffer (currently ignored) /// @return returns a handle with a grid of type NanoGrid - template + template GridHandle getHandle(const PtrT points, size_t pointCount, const BufferT &buffer = BufferT()); @@ -531,7 +532,7 @@ PointsToGrid::getHandle(const PtrT points, if (mVerbose==1) mTimer.stop(); if (mVerbose==1) mTimer.restart("Computation of checksum"); - tools::cuda::updateChecksum((GridData*)buffer.deviceData(), mChecksum, mStream); + tools::cuda::updateChecksum((GridData*)nanovdb::cuda::detail::deviceStorageData(buffer), mChecksum, mStream); if (mVerbose==1) mTimer.stop(); cudaStreamSynchronize(mStream); @@ -858,10 +859,10 @@ inline BufferT PointsToGrid::getBuffer(const PtrT, size_t poi int device = 0; cudaGetDevice(&device); - auto buffer = BufferT::create(mData.size, &pool, device, mStream);// only allocate buffer on the device + auto buffer = nanovdb::cuda::detail::createDeviceStorage(mData.size, &pool, device, mStream); // only allocate buffer on the device - mData.d_bufferPtr = buffer.deviceData(); - if (mData.d_bufferPtr == nullptr) throw std::runtime_error("Failed to allocate grid buffer on the device"); + mData.d_bufferPtr = nanovdb::cuda::detail::deviceStorageData(buffer); + if (mData.d_bufferPtr == nullptr) throw std::runtime_error("The grid buffer type produced no device-accessible memory"); // 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 @@ -1460,7 +1461,7 @@ voxelsToGrid(std::vector> vec, const Buffer //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template +template [[deprecated("Use cuda::pointsToGrid instead")]] GridHandle cudaPointsToGrid(const PtrT dWorldPoints, @@ -1475,7 +1476,7 @@ cudaPointsToGrid(const PtrT dWorldPoints, //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template +template [[deprecated("Use cuda::pointsToGrid instead")]] GridHandle cudaPointsToGrid(std::vector> pointSet, @@ -1487,7 +1488,7 @@ cudaPointsToGrid(std::vector> poi //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template +template [[deprecated("Use cuda::voxelsToGrid instead")]] GridHandle cudaVoxelsToGrid(const PtrT dGridVoxels, @@ -1501,7 +1502,7 @@ cudaVoxelsToGrid(const PtrT dGridVoxels, //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- -template +template [[deprecated("Use cuda::voxelsToGrid instead")]] GridHandle cudaVoxelsToGrid(std::vector> pointSet, diff --git a/nanovdb/nanovdb/tools/cuda/PruneGrid.cuh b/nanovdb/nanovdb/tools/cuda/PruneGrid.cuh index ee4f3f6e1f..7c2f0c9a8d 100644 --- a/nanovdb/nanovdb/tools/cuda/PruneGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/PruneGrid.cuh @@ -65,7 +65,7 @@ public: /// @tparam BufferT Buffer type used for allocation of the grid handle /// @param buffer optional buffer (currently ignored) /// @return returns a handle with a grid of type NanoGrid - template + template GridHandle getHandle(const BufferT &buffer = BufferT()); @@ -199,13 +199,12 @@ void PruneGrid::pruneRoot() // Package the duplicated root topology into a RootNode plus Tile list; upload to the GPU uint64_t rootSize = RootT::memUsage(prunedTiles.size()); - mBuilder.mProcessedRoot = nanovdb::cuda::DeviceBuffer::create(rootSize); - auto prunedRootPtr = static_cast(mBuilder.mProcessedRoot.data()); + auto prunedRootPtr = mBuilder.allocateProcessedRoot(rootSize); prunedRootPtr->mTableSize = prunedTiles.size(); uint32_t t = 0; for (const auto& [key, tile] : prunedTiles) *prunedRootPtr->tile(t++) = tile; - mBuilder.mProcessedRoot.deviceUpload(device, mStream, false); + mBuilder.uploadProcessedRoot(mStream); }// PruneGrid::pruneRoot //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- diff --git a/nanovdb/nanovdb/tools/cuda/RefineGrid.cuh b/nanovdb/nanovdb/tools/cuda/RefineGrid.cuh index 410173c49d..0140e10ebe 100644 --- a/nanovdb/nanovdb/tools/cuda/RefineGrid.cuh +++ b/nanovdb/nanovdb/tools/cuda/RefineGrid.cuh @@ -64,7 +64,7 @@ public: /// @tparam BufferT Buffer type used for allocation of the grid handle /// @param buffer optional buffer (currently ignored) /// @return returns a handle with a grid of type NanoGrid - template + template GridHandle getHandle(const BufferT &buffer = BufferT()); @@ -211,13 +211,12 @@ void RefineGrid::refineRoot() // Package the new root topology into a RootNode plus Tile list; upload to the GPU uint64_t rootSize = RootT::memUsage(refinedTiles.size()); - mBuilder.mProcessedRoot = nanovdb::cuda::DeviceBuffer::create(rootSize); - auto refinedRootPtr = static_cast(mBuilder.mProcessedRoot.data()); + auto refinedRootPtr = mBuilder.allocateProcessedRoot(rootSize); refinedRootPtr->mTableSize = refinedTiles.size(); uint32_t t = 0; for (const auto& [key, tile] : refinedTiles) *refinedRootPtr->tile(t++) = tile; - mBuilder.mProcessedRoot.deviceUpload(device, mStream, false); + mBuilder.uploadProcessedRoot(mStream); }// RefineGrid::refineRoot //------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- diff --git a/nanovdb/nanovdb/tools/cuda/SignedFloodFill.cuh b/nanovdb/nanovdb/tools/cuda/SignedFloodFill.cuh index 1ebda2d542..422af5e904 100644 --- a/nanovdb/nanovdb/tools/cuda/SignedFloodFill.cuh +++ b/nanovdb/nanovdb/tools/cuda/SignedFloodFill.cuh @@ -25,7 +25,7 @@ #include #include #include -#include +#include #include #include #include @@ -108,18 +108,23 @@ void processRoot(NanoTree *d_tree, cudaStream_t 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) - if (!uBuffer.data()->isRootNext()) throw std::runtime_error("ERROR: expected no padding between tree and root!"); - if ( uBuffer.data()->root().tileCount() == 0) return;// empty root node so nothing to do - uBuffer.resize(sizeof(TreeT) + uBuffer.data()->root().memUsage());// likely does nothing since we reserved 64 tiles - RootT *root = &uBuffer.data()->root(); + // First copy the tree and root and then its tiles, which is of unknown size; + // managed memory, because the scanline pass below interleaves host access + // with device reads of the same bytes + using ManagedBufT = nanovdb::cuda::Buffer; + ManagedBufT uBuffer(sizeof(TreeT) + sizeof(RootT) + 64*sizeof(TileT), nanovdb::cuda::noInit); + cudaCheck(cudaMemcpy(uBuffer.data(), d_tree, sizeof(TreeT) + sizeof(RootT), cudaMemcpyDeviceToHost)); // copy Tree and Root (minus tiles) + auto *tree = reinterpret_cast(uBuffer.data()); + if (!tree->isRootNext()) throw std::runtime_error("ERROR: expected no padding between tree and root!"); + if ( tree->root().tileCount() == 0) return; // empty root node so nothing to do + uBuffer.resize(sizeof(TreeT) + tree->root().memUsage()); // grows (with a copy) past the 64 reserved tiles + tree = reinterpret_cast(uBuffer.data()); // resize may reallocate + RootT *root = &tree->root(); cudaCheck(cudaMemcpy(root + 1, (char*)(d_tree + 1) + sizeof(RootT), root->tileCount()*sizeof(TileT), cudaMemcpyDeviceToHost));// copy tiles // Sort the child nodes of the root in lexicographic order - nanovdb::cuda::UnifiedBuffer nodeBuffer(root->tileCount()*sizeof(ChildT));// potential over-allocation - auto *first = nodeBuffer.data(), *last = first; + ManagedBufT nodeBuffer(root->tileCount()*sizeof(ChildT), nanovdb::cuda::noInit); // potential over-allocation + auto *first = reinterpret_cast(nodeBuffer.data()), *last = first; for (auto it=root->beginChild(); it; ++it) *last++ = ChildT(it.getCoord(), it.pos()); if (last - first < 2) return;// zero or one child node so nothing to do! std::sort(first, last, ChildT());// lexicographic ordering diff --git a/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh b/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh index 9159002927..4f82f1e0f5 100644 --- a/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh +++ b/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh @@ -21,6 +21,8 @@ #include #include #include +#include +#include // for the pinned host staging of the processed root namespace nanovdb { @@ -68,6 +70,7 @@ class TopologyBuilder /// Buffer rather than the dual DeviceBuffer, whose host pointer and /// per-device array they would leave unused. using ScratchT = nanovdb::cuda::Buffer>; + using HostStagingT = nanovdb::cuda::Buffer; public: @@ -83,10 +86,11 @@ public: , mVoxelOffsets(stream, resource, 0, nanovdb::cuda::noInit) , mLowerParents(stream, resource, 0, nanovdb::cuda::noInit) , mLeafParents(stream, resource, 0, nanovdb::cuda::noInit) + , mDeviceRoot(stream, resource, 0, nanovdb::cuda::noInit) + , mDeviceData(stream, resource, 0, nanovdb::cuda::noInit) , mResource(&resource) , mTempDevicePool(resource) { - mData = nanovdb::cuda::DeviceBuffer::create(sizeof(Data)); } using Data = TopologyBuilderData; @@ -108,7 +112,8 @@ public: void postProcessGridTree(cudaStream_t stream); - nanovdb::cuda::DeviceBuffer mProcessedRoot; + HostStagingT mHostRoot; // host staging for the processed root (pinned, so the upload is asynchronous) + ScratchT mDeviceRoot; // device copy, made by uploadProcessedRoot ScratchT mUpperMasks; ScratchT mLowerMasks; ScratchT mUpperOffsets; @@ -117,15 +122,47 @@ public: ScratchT mVoxelOffsets; ScratchT mLowerParents; ScratchT mLeafParents; - nanovdb::cuda::DeviceBuffer mData; + Data mHostData{}; // host side of the builder parameters + ScratchT mDeviceData; // device copy, made by uploadData CheckMode mChecksum{CheckMode::Disable}; - auto deviceProcessedRoot() { return static_cast(mProcessedRoot.deviceData()); } - auto hostProcessedRoot() { return static_cast(mProcessedRoot.data()); } + auto deviceProcessedRoot() { return reinterpret_cast(mDeviceRoot.data()); } + auto hostProcessedRoot() { return reinterpret_cast(mHostRoot.data()); } + + /// @brief Allocates (pinned) host staging for the processed root and + /// returns it for the caller to fill; any previous root is dropped. + RootT* allocateProcessedRoot(uint64_t bytes) + { + mHostRoot = HostStagingT(bytes, nanovdb::cuda::noInit); + return reinterpret_cast(mHostRoot.data()); + } + + /// @brief Copies the host-staged processed root to the device, allocating + /// through the builder's resource when the device copy is missing + /// or too small. + void uploadProcessedRoot(cudaStream_t stream) + { + if (mDeviceRoot.size() < mHostRoot.size()) + mDeviceRoot = ScratchT(stream, nanovdb::cuda::ResourceRef(*mResource), mHostRoot.size(), nanovdb::cuda::noInit); + cudaCheck(cudaMemcpyAsync(mDeviceRoot.data(), mHostRoot.data(), mHostRoot.size(), cudaMemcpyHostToDevice, stream)); + } + + /// @brief Copies the builder parameters to the device, allocating through + /// the builder's resource on first use. + void uploadData(cudaStream_t stream) + { + if (mDeviceData.empty()) + mDeviceData = ScratchT(stream, nanovdb::cuda::ResourceRef(*mResource), sizeof(Data), nanovdb::cuda::noInit); + cudaCheck(cudaMemcpyAsync(mDeviceData.data(), &mHostData, sizeof(Data), cudaMemcpyHostToDevice, stream)); + } void* deviceUpperMasks() { return mUpperMasks.data(); } void* deviceLowerMasks() { return mLowerMasks.data(); } - Data* data() { return static_cast(mData.data()); } - Data* deviceData() { return static_cast(mData.deviceData()); } + /// @brief A borrowing reference to the builder's resource, for consumers + /// allocating sibling scratch from the same instance. + nanovdb::cuda::ResourceRef ref() { return nanovdb::cuda::ResourceRef(*mResource); } + + Data* data() { return &mHostData; } + Data* deviceData() { return reinterpret_cast(mDeviceData.data()); } private: static constexpr unsigned int mNumThreads = 128;// for kernels spawned via lambdaKernel (others may specialize) @@ -252,14 +289,14 @@ BufferT TopologyBuilder::getBuffer(const BufferT &pool, cudaS int device = 0; cudaGetDevice(&device); - auto buffer = BufferT::create(data()->size, &pool, device, stream);// only allocate buffer on the device - cudaCheck(cudaMemsetAsync(buffer.deviceData(), 0, data()->size, stream)); + auto buffer = nanovdb::cuda::detail::createDeviceStorage(data()->size, &pool, device, stream); // only allocate buffer on the device + cudaCheck(cudaMemsetAsync(nanovdb::cuda::detail::deviceStorageData(buffer), 0, data()->size, stream)); - data()->d_bufferPtr = buffer.deviceData(); + data()->d_bufferPtr = nanovdb::cuda::detail::deviceStorageData(buffer); if (data()->d_bufferPtr == nullptr) throw std::runtime_error("Failed to allocate grid buffer on the device"); if (data()->nodeCount[2] != 0) // Unless the result is an empty grid data()->d_upperOffsets = reinterpret_cast(mUpperOffsets.data()); - mData.deviceUpload(device, stream, false); + this->uploadData(stream); return buffer; }// TopologyBuilder::getBuffer @@ -483,7 +520,8 @@ inline void TopologyBuilder::processLowerNodes(cudaStream_t s cudaCheckError(); } - mProcessedRoot.clear(stream); + mHostRoot.destroy(); + mDeviceRoot.destroy(stream); mUpperMasks.destroy(stream); mLowerMasks.destroy(stream); mLowerOffsets.destroy(stream); diff --git a/nanovdb/nanovdb/tools/cuda/VoxelBlockManager.cuh b/nanovdb/nanovdb/tools/cuda/VoxelBlockManager.cuh index ca31864798..cb119f1fc2 100644 --- a/nanovdb/nanovdb/tools/cuda/VoxelBlockManager.cuh +++ b/nanovdb/nanovdb/tools/cuda/VoxelBlockManager.cuh @@ -38,6 +38,7 @@ #include #include #include +#include namespace nanovdb { @@ -373,7 +374,7 @@ void buildVoxelBlockManager( /// Returns a fully-constructed VoxelBlockManagerHandle backed by device memory. /// Grid dimensions (when not supplied) are read from device memory via DeviceGridTraits. /// @tparam Log2BlockWidth Log2 of the number of active voxels per VBM block -/// @tparam BufferT Device buffer type (default: nanovdb::cuda::DeviceBuffer) +/// @tparam BufferT Device buffer type (default: nanovdb::cuda::DualDeviceBuffer) /// @param d_grid Device-side grid pointer /// @param firstOffset First active-voxel offset covered by this VBM; must satisfy /// firstOffset == 1 (mod BlockWidth). Pass 0 (default) to use 1, @@ -385,14 +386,15 @@ void buildVoxelBlockManager( /// (default) to use the minimum required capacity. /// @param stream CUDA stream (default 0) /// @return A fully constructed VoxelBlockManagerHandle backed by device memory -template +template nanovdb::tools::VoxelBlockManagerHandle buildVoxelBlockManager( NanoGrid* d_grid, uint64_t firstOffset = 0, uint64_t lastOffset = 0, uint64_t nBlocks = 0, - cudaStream_t stream = 0) + cudaStream_t stream = 0, + const BufferT* proto = nullptr) { static constexpr uint64_t BlockWidth = uint64_t(1) << Log2BlockWidth; static constexpr uint64_t JumpMapLength = BlockWidth / 64; @@ -400,15 +402,29 @@ buildVoxelBlockManager( using Traits = util::cuda::DeviceGridTraits; if (!firstOffset) firstOffset = 1; if (!lastOffset) lastOffset = Traits::getActiveVoxelCount(d_grid); - if (lastOffset < firstOffset) return nanovdb::tools::VoxelBlockManagerHandle{}; + if (lastOffset < firstOffset) {// empty grid: an empty handle, through the pool's resource when one is needed + if constexpr (BufferIsDefaultConstructible::value) { + return nanovdb::tools::VoxelBlockManagerHandle{}; + } else { + if (!proto) + throw std::runtime_error("buildVoxelBlockManager: an empty handle over a buffer type that is " + "not default-constructible requires a prototype buffer to take the resource from"); + int device = 0; + cudaCheck(cudaGetDevice(&device)); + return nanovdb::tools::VoxelBlockManagerHandle( + nanovdb::cuda::detail::createDeviceStorage(0, proto, device, stream), + nanovdb::cuda::detail::createDeviceStorage(0, proto, device, stream), + 0, firstOffset, lastOffset); // zero-size buffers allocate nothing + } + } NANOVDB_ASSERT(!((firstOffset - 1) & (BlockWidth - 1))); // firstOffset == 1 (mod BlockWidth) if (!nBlocks) nBlocks = (lastOffset - firstOffset + BlockWidth) >> Log2BlockWidth; int device = 0; cudaCheck(cudaGetDevice(&device)); - auto firstLeafIDBuf = BufferT::create(nBlocks * sizeof(uint32_t), nullptr, device, stream); - auto jumpMapBuf = BufferT::create(nBlocks * JumpMapLength * sizeof(uint64_t), nullptr, device, stream); + auto firstLeafIDBuf = nanovdb::cuda::detail::createDeviceStorage(nBlocks * sizeof(uint32_t), proto, device, stream); + auto jumpMapBuf = nanovdb::cuda::detail::createDeviceStorage(nBlocks * JumpMapLength * sizeof(uint64_t), proto, device, stream); nanovdb::tools::VoxelBlockManagerHandle handle( std::move(firstLeafIDBuf), std::move(jumpMapBuf), diff --git a/nanovdb/nanovdb/unittest/TestBuffer.cu b/nanovdb/nanovdb/unittest/TestBuffer.cu index 1afe5f4b01..fc92955b19 100644 --- a/nanovdb/nanovdb/unittest/TestBuffer.cu +++ b/nanovdb/nanovdb/unittest/TestBuffer.cu @@ -12,6 +12,9 @@ #include #include #include +#include // for the voxelsToGrid entry-point test +#include +#include // for the single-space entry-point test #include #include @@ -1026,6 +1029,20 @@ TEST(TestBuffer, GridHandleCopyToPinnedRoundTrip) EXPECT_NE(dev2.deviceGrid(1), nullptr); } +TEST(TestBuffer, GridHandleCopyToManagedSynchronizes) +{ + auto host = nanovdb::tools::createLevelSetSphere(20.0, nanovdb::Vec3d(0), 1.0, 3.0, nanovdb::Vec3d(0), "sphere"); + using ManagedBufT = nanovdb::cuda::Buffer; + // A managed destination is host-readable, so copyTo synchronizes before + // returning: the host accessors must be valid immediately, with no + // synchronization by the caller. + auto managed = nanovdb::cuda::copyTo(host, cudaStream_t(0)); + EXPECT_EQ(1u, managed.gridCount()); + ASSERT_NE(managed.grid(), nullptr); + EXPECT_EQ(0, std::memcmp(managed.data(), host.data(), host.bufferSize())); + EXPECT_NE(managed.deviceGrid(), nullptr);// the same allocation serves the device accessors +} + TEST(TestBuffer, GridHandleCopyToProtoResource) { auto host = nanovdb::tools::createLevelSetSphere(20.0, nanovdb::Vec3d(0), 1.0, 3.0, nanovdb::Vec3d(0), "sphere"); @@ -1037,7 +1054,7 @@ TEST(TestBuffer, GridHandleCopyToProtoResource) DevBufT proto(cudaStream_t(0), RefT(res), 16, nanovdb::cuda::noInit);// alloc #1: an exemplar carrying the borrowed resource auto dev = nanovdb::cuda::copyTo(host, cudaStream_t(0), &proto); ASSERT_EQ(cudaSuccess, cudaStreamSynchronize(0)); - EXPECT_EQ(3, counters.allocs);// #2: the grid storage, #3: the metadata scratch, all through the proto's resource + EXPECT_EQ(2, counters.allocs); // #2: the grid storage; the metadata is adopted from the source handle, so no scratch EXPECT_NE(dev.deviceGrid(), nullptr); nanovdb::GridHandle empty; @@ -1229,4 +1246,104 @@ TEST(TestBuffer, SingleSpaceNodeManager) EXPECT_EQ(counters.allocs, counters.deallocs); } +TEST(TestBuffer, SingleSpaceToolEntryPoints) +{ + // The builders allocate their result handle through createDeviceStorage, + // so a single-space buffer type works wherever a dual-space one does. + // voxelsToGrid stands in for the whole PointsToGrid family; the pool + // buffer supplies the resource for the handle storage. + nanovdb::Coord coords[2] = {nanovdb::Coord(1,2,3), nanovdb::Coord(10,20,8)}, *d_coords = nullptr; + ASSERT_EQ(cudaSuccess, cudaMalloc(&d_coords, 2*sizeof(nanovdb::Coord))); + ASSERT_EQ(cudaSuccess, cudaMemcpy(d_coords, coords, 2*sizeof(nanovdb::Coord), cudaMemcpyHostToDevice)); + + Counters counters; + CountingResource res{&counters}; + using RefT = nanovdb::cuda::ResourceRef; + using BufT = nanovdb::cuda::Buffer; + { + BufT pool(cudaStream_t(0), RefT(res), 16, nanovdb::cuda::noInit); // alloc #1: exemplar carrying the borrowed resource + auto handle = nanovdb::tools::cuda::voxelsToGrid(d_coords, 2, 1.0, pool); + ASSERT_EQ(cudaSuccess, cudaStreamSynchronize(0)); + EXPECT_EQ(1u, handle.gridCount()); + EXPECT_NE(handle.deviceGrid(), nullptr); + EXPECT_EQ(3, counters.allocs); // #2: the grid storage, #3: the handle's metadata scratch + } + ASSERT_EQ(cudaSuccess, cudaStreamSynchronize(0)); + EXPECT_EQ(counters.allocs, counters.deallocs); + ASSERT_EQ(cudaSuccess, cudaFree(d_coords)); +} + +static_assert(nanovdb::BufferHasDeviceSingle>::value, + "managed storage is device-accessible"); +static_assert(nanovdb::BufferHasHostSingle>::value, + "managed storage is host-accessible"); + +TEST(TestBuffer, ManagedBufferBothSpaces) +{ + // A managed-resource buffer serves grids read on both sides: the handle + // parses metadata on the host and exposes BOTH accessor families over + // the same allocation. + auto host = nanovdb::tools::createLevelSetSphere(20.0, nanovdb::Vec3d(0), 1.0, 3.0, nanovdb::Vec3d(0), "sphere"); + using BufT = nanovdb::cuda::Buffer; + BufT buf(host.bufferSize(), nanovdb::cuda::noInit); + std::memcpy(buf.data(), host.data(), host.bufferSize()); // managed memory is host-writable + + nanovdb::GridHandle handle(std::move(buf)); + EXPECT_EQ(1u, handle.gridCount()); + ASSERT_NE(handle.grid(), nullptr); // host accessor + ASSERT_NE(handle.deviceGrid(), nullptr); // device accessor, same bytes + EXPECT_EQ((const void*)handle.grid(), (const void*)handle.deviceGrid()); + EXPECT_EQ(std::string("sphere"), handle.grid()->gridName()); + + // built on the device through a tool entry point, read back on the host + nanovdb::Coord coords[2] = {nanovdb::Coord(1,2,3), nanovdb::Coord(10,20,8)}, *d_coords = nullptr; + ASSERT_EQ(cudaSuccess, cudaMalloc(&d_coords, 2*sizeof(nanovdb::Coord))); + ASSERT_EQ(cudaSuccess, cudaMemcpy(d_coords, coords, 2*sizeof(nanovdb::Coord), cudaMemcpyHostToDevice)); + auto built = nanovdb::tools::cuda::voxelsToGrid(d_coords, 2); + ASSERT_EQ(cudaSuccess, cudaDeviceSynchronize()); // device writes must land before host reads + ASSERT_NE(built.grid(), nullptr); + EXPECT_TRUE(built.grid()->tree().isActive(nanovdb::Coord(1,2,3))); + + nanovdb::cuda::ManagedResource managed; + auto mgr = nanovdb::cuda::createNodeManager(built.deviceGrid(), managed); + ASSERT_EQ(cudaSuccess, cudaStreamSynchronize(0)); + EXPECT_NE(mgr.deviceMgr(), nullptr); // a managed NodeManager serves the device... + EXPECT_NE(mgr.mgr(), nullptr); // ...and the host + ASSERT_EQ(cudaSuccess, cudaFree(d_coords)); +} + +TEST(TestBuffer, SingleSpaceVoxelBlockManager) +{ + // The VoxelBlockManager entry point allocates through createDeviceStorage + // and its handle maps the device accessors onto single-space buffers. + nanovdb::Coord coords[2] = {nanovdb::Coord(1,2,3), nanovdb::Coord(10,20,8)}, *d_coords = nullptr; + ASSERT_EQ(cudaSuccess, cudaMalloc(&d_coords, 2*sizeof(nanovdb::Coord))); + ASSERT_EQ(cudaSuccess, cudaMemcpy(d_coords, coords, 2*sizeof(nanovdb::Coord), cudaMemcpyHostToDevice)); + auto gridHandle = nanovdb::tools::cuda::voxelsToGrid(d_coords, 2); + auto* d_grid = gridHandle.deviceGrid(); + ASSERT_NE(d_grid, nullptr); + + using BufT = nanovdb::cuda::Buffer; + auto vbm = nanovdb::tools::cuda::buildVoxelBlockManager<6, BufT>(d_grid); + ASSERT_EQ(cudaSuccess, cudaStreamSynchronize(0)); + EXPECT_GT(vbm.blockCount(), 0u); + EXPECT_NE(vbm.deviceFirstLeafID(), nullptr); + EXPECT_NE(vbm.deviceJumpMap(), nullptr); + + Counters counters; + CountingResource res{&counters}; + using RefT = nanovdb::cuda::ResourceRef; + using RefBufT = nanovdb::cuda::Buffer; + { + RefBufT proto(cudaStream_t(0), RefT(res), 16, nanovdb::cuda::noInit); // exemplar carrying the borrowed resource + auto vbm2 = nanovdb::tools::cuda::buildVoxelBlockManager<6, RefBufT>(d_grid, 0, 0, 0, cudaStream_t(0), &proto); + ASSERT_EQ(cudaSuccess, cudaStreamSynchronize(0)); + EXPECT_EQ(3, counters.allocs); // proto + firstLeafID + jumpMap, all through the resource + EXPECT_NE(vbm2.deviceFirstLeafID(), nullptr); + } + ASSERT_EQ(cudaSuccess, cudaStreamSynchronize(0)); + EXPECT_EQ(counters.allocs, counters.deallocs); + ASSERT_EQ(cudaSuccess, cudaFree(d_coords)); +} + } // unnamed namespace diff --git a/nanovdb/nanovdb/unittest/TestMemoryResource.cu b/nanovdb/nanovdb/unittest/TestMemoryResource.cu index b6f6c9eb6f..82589773c4 100644 --- a/nanovdb/nanovdb/unittest/TestMemoryResource.cu +++ b/nanovdb/nanovdb/unittest/TestMemoryResource.cu @@ -27,6 +27,21 @@ #include +// These tests deliberately keep exercising the deprecated dual-space +// DeviceBuffer surface until its removal; the deprecation warnings are +// suppressed for this translation unit only. New code must use +// cuda::Buffer and cuda::copyTo instead. +#if defined(__GNUC__) || defined(__clang__) +#pragma GCC diagnostic ignored "-Wdeprecated-declarations" +#endif +#if defined(_MSC_VER) +#pragma warning(disable : 4996) +#endif +#if defined(__CUDACC__) +#pragma nv_diag_suppress 20199 +#endif + + namespace { //====================================================================== diff --git a/nanovdb/nanovdb/unittest/TestNanoVDB.cu b/nanovdb/nanovdb/unittest/TestNanoVDB.cu index 37a447e815..1923d9b360 100644 --- a/nanovdb/nanovdb/unittest/TestNanoVDB.cu +++ b/nanovdb/nanovdb/unittest/TestNanoVDB.cu @@ -40,6 +40,21 @@ #include // for std::setw, std::setfill #include // for std::thread +// These tests deliberately keep exercising the deprecated dual-space +// DeviceBuffer surface until its removal; the deprecation warnings are +// suppressed for this translation unit only. New code must use +// cuda::Buffer and cuda::copyTo instead. +#if defined(__GNUC__) || defined(__clang__) +#pragma GCC diagnostic ignored "-Wdeprecated-declarations" +#endif +#if defined(_MSC_VER) +#pragma warning(disable : 4996) +#endif +#if defined(__CUDACC__) +#pragma nv_diag_suppress 20199 +#endif + + namespace nanovdb {// this namespace is required by gtest namespace test { diff --git a/pendingchanges/nanovdb.txt b/pendingchanges/nanovdb.txt index 5fbc59f018..6574287838 100644 --- a/pendingchanges/nanovdb.txt +++ b/pendingchanges/nanovdb.txt @@ -4,6 +4,8 @@ NanoVDB: - Added new _hostdev_ function named nanovdb::math::isoCrossing, which intersects a ray against a user-defined iso-surface. - Added doc/nanovdb/TEACHME, an interactive tutorial for the NanoVDB user API: lesson documents written to be loaded by an LLM coding agent, which then teaches the API interactively to a developer new to NanoVDB. Ships with a cheat sheet, a GPU ray-marching capstone, and a CI harness that compiles every code block in the lesson against the NanoVDB headers so the tutorial cannot drift from the API. - Added nanovdb::cuda::Buffer and nanovdb::cuda::BufferView (CUDA): a typed, resource-aware, stream-ordered container that allocates from an injectable memory resource and frees on its retained stream, and a non-owning view over externally managed memory that a GridHandle can wrap without copying. Member names follow cuda::buffer (destroy, set_stream, swap). Also added the synchronous resource concept nanovdb::cuda::is_resource alongside is_async_resource. + - Added nanovdb::cuda::ManagedResource (CUDA), a synchronous cudaMallocManaged-backed memory resource whose allocations are host- and device-accessible: a GridHandle or NodeManagerHandle over cuda::Buffer parses its metadata on the host and exposes both the host and the device accessors over the same allocation -- the migration target for UnifiedBuffer users. Detected via the new nanovdb::cuda::is_device_accessible_resource trait. + - The GPU tools' entry points now allocate their result handle through either buffer family: a single-space cuda::Buffer type works wherever a dual-space buffer does, allocating through the pool buffer's memory resource (new bridge header nanovdb/cuda/HandleStorage.h). - Added nanovdb::cuda::copyTo (CUDA), the explicit, stream-carrying transfer between grid handles in different address spaces: single-space device handles copy to host-readable handles (HostBuffer or a pinned-resource cuda::Buffer) and back, and device-to-device across different resources. An optional prototype buffer routes the destination allocation through a caller-supplied resource. Dual-space handles keep using deviceUpload/deviceDownload. Improvements: @@ -15,6 +17,9 @@ NanoVDB: - GridHandle::reset, NodeManagerHandle::reset and tools::VoxelBlockManager::reset now release storage through the buffer's destroy() when it provides one, and cuda::Buffer::clear/cuda::BufferView::clear are deprecated in favor of destroy() (BufferView::destroy detaches the non-owning view). The legacy HostBuffer/DeviceBuffer clear() methods are unaffected. - The bug-fix to the nanovdb::ReadAccessor (see below) improves random-access performance in some use-cases (especially on the CPU). + Deprecations: + - nanovdb::cuda::DeviceBuffer is deprecated: grid storage is moving to the single-space nanovdb::cuda::Buffer, with explicit transfers via nanovdb::cuda::copyTo. Code that uses the GPU tools' default buffer type is unaffected until the removal (the defaults now name the transitional DualDeviceBuffer implementation); only code naming DeviceBuffer sees the warning, whose message carries the migration recipe. Note that constructing a device grid handle validates the grid with a kernel, so transfers belong in CUDA translation units; host-only source files hold the returned handle and delegate the copy to a small helper in a CUDA file (the CUDA examples demonstrate this). The tools' internal scratch, the examples and the NanoVDB tests are fully migrated; UnifiedBuffer follows in the multi-GPU work, in the same release. + Fixes: - nanovdb::cuda::createNodeManager now synchronizes the stream before reading back the device-computed NodeManager size: the host previously consumed the destination of an asynchronous copy without ordering, which happened to work only because device-to-pageable copies degrade to synchronous behavior. - tools::cuda::addBlindData and tools::cuda::indexToGrid now normalize the output grid's index and count fields (mGridIndex=0, mGridCount=1): both tools copy the source grid's header, so a source grid taken from a multi-grid buffer used to leave a stale index/count in the single-grid output, corrupting the metadata parse of the returned handle.