diff --git a/nanovdb/nanovdb/cuda/DeviceResource.h b/nanovdb/nanovdb/cuda/DeviceResource.h index bef41f96d9..21da9b0ca2 100644 --- a/nanovdb/nanovdb/cuda/DeviceResource.h +++ b/nanovdb/nanovdb/cuda/DeviceResource.h @@ -192,6 +192,79 @@ struct SyncFromAsync } }; +/// @brief Synchronous device memory resource backed by cudaMalloc/cudaFree. +/// Models only the Resource concept: it never touches stream-ordered +/// allocation, so it works on devices without memory-pool support +/// (cudaDevAttrMemoryPoolsSupported == 0), where DeviceResource's +/// cudaMallocAsync path fails by design. Pair with AsyncFromSync to +/// drive the stream-ordered builders on such a device. +class MallocResource +{ +public: + // cudaMalloc aligns memory to 256 bytes by default + static constexpr size_t DEFAULT_ALIGNMENT = 256; + + /// @brief Allocates @c bytes with cudaMalloc; valid on every stream when + /// this returns. A zero request returns nullptr. + void* allocate(size_t bytes, size_t) + { + if (bytes == 0) return nullptr; + void* p = nullptr; + cudaCheck(cudaMalloc(&p, bytes)); + return p; + } + + /// @brief Frees @c p with cudaFree; the caller guarantees that device work + /// touching the memory has completed. + void deallocate(void* p, size_t, size_t) { cudaCheck(cudaFree(p)); } +};// MallocResource + +/// @brief Wrapper presenting a synchronous resource as a stream-ordered one, +/// so it can drive components that require the AsyncResource concept +/// (TempPool and the GPU builders). +/// @tparam R the wrapped synchronous resource, held by value; wrap a +/// ResourceRef to borrow a stateful instance instead. +/// @details The mirror of SyncFromAsync, and the analog of cuda::mr's +/// synchronous_resource_adapter. allocate_async forwards to +/// R::allocate, whose memory is immediately valid on every stream -- +/// a stronger guarantee than stream-ordering requires. +/// deallocate_async synchronizes @c stream before R::deallocate, +/// establishing the quiescence the synchronous contract demands. +/// @warning Every deallocation synchronizes its stream, so expect +/// serialization relative to a genuinely stream-ordered resource. +/// That is the unavoidable cost of a synchronous backend under a +/// stream-ordered algorithm; this wrapper exists so the cost is +/// explicit and chosen by the caller -- e.g. on a device without +/// memory-pool support -- rather than silently substituted. +template +struct AsyncFromSync +{ + static_assert(is_resource::value, + "AsyncFromSync requires R to model the synchronous Resource concept"); + + static constexpr size_t DEFAULT_ALIGNMENT = R::DEFAULT_ALIGNMENT; + + R resource; + + /// @brief Allocates through the synchronous resource; the result is valid + /// on every stream, hence trivially valid on @c stream. + void* allocate_async(size_t bytes, size_t alignment, cudaStream_t) { return resource.allocate(bytes, alignment); } + + /// @brief Synchronizes @c stream, then frees through the synchronous + /// resource -- the synchronize makes the quiescence contract hold. + /// Null is a no-op and skips the synchronize. + void deallocate_async(void* p, size_t bytes, size_t alignment, cudaStream_t stream) + { + if (p == nullptr) return; + cudaCheck(cudaStreamSynchronize(stream)); + resource.deallocate(p, bytes, alignment); + } + + /// @brief Synchronous pair, forwarding to the wrapped resource. + void* allocate(size_t bytes, size_t alignment) { return resource.allocate(bytes, alignment); } + void deallocate(void* p, size_t bytes, size_t alignment) { resource.deallocate(p, bytes, alignment); } +};// AsyncFromSync + /// @brief Non-owning reference to a memory resource that is itself a resource: /// copying the ref shares the underlying resource rather than copying it. /// @tparam R the referenced resource type diff --git a/nanovdb/nanovdb/unittest/TestMemoryResource.cu b/nanovdb/nanovdb/unittest/TestMemoryResource.cu index c9fac7f460..d50d87bf22 100644 --- a/nanovdb/nanovdb/unittest/TestMemoryResource.cu +++ b/nanovdb/nanovdb/unittest/TestMemoryResource.cu @@ -141,6 +141,57 @@ TEST(TestMemoryResource, PinnedResource_DefaultResourceRoundTrip) // retained stream (the stream of the most recent reallocate), not the null stream. //====================================================================== +// Synchronous-only and stateful: the shape of a vGPU or arena backend. +struct SyncCountingResource +{ + static constexpr size_t DEFAULT_ALIGNMENT = nanovdb::cuda::MallocResource::DEFAULT_ALIGNMENT; + int allocs = 0, deallocs = 0; + void* allocate(size_t bytes, size_t alignment) { + void* p = nanovdb::cuda::MallocResource{}.allocate(bytes, alignment); + if (p) ++allocs; + return p; + } + void deallocate(void* p, size_t bytes, size_t alignment) { + if (p) ++deallocs; + nanovdb::cuda::MallocResource{}.deallocate(p, bytes, alignment); + } +}; + +static_assert(nanovdb::cuda::is_resource::value, + "MallocResource must model the synchronous Resource concept"); +static_assert(!nanovdb::cuda::is_async_resource::value, + "MallocResource must not claim the AsyncResource concept"); +static_assert(nanovdb::cuda::is_async_resource< + nanovdb::cuda::AsyncFromSync>::value, + "AsyncFromSync must lift a synchronous resource to AsyncResource"); + +TEST(TestMemoryResource, PointsToGrid_RunsOnSynchronousResource) +{ + // The pool-less-device path: every scratch allocation routes through + // cudaMalloc/cudaFree via AsyncFromSync, never touching cudaMallocAsync. + // The grid handle's output buffer is the exception -- getHandle allocates + // it through BufferT::create, not through the injected resource. + using RefT = nanovdb::cuda::ResourceRef; + using VgpuT = nanovdb::cuda::AsyncFromSync; + SyncCountingResource base; + VgpuT res{RefT(base)}; + + const std::vector voxels = {{0,0,0},{1,2,3},{8,8,8},{100,100,100},{-50,20,7}}; + nanovdb::Coord* d_voxels = nullptr; + ASSERT_EQ(cudaMalloc(&d_voxels, voxels.size()*sizeof(nanovdb::Coord)), cudaSuccess); + ASSERT_EQ(cudaMemcpy(d_voxels, voxels.data(), voxels.size()*sizeof(nanovdb::Coord), cudaMemcpyHostToDevice), cudaSuccess); + { + nanovdb::tools::cuda::PointsToGrid converter(nanovdb::Map(1.0), cudaStream_t{0}, res); + auto handle = converter.getHandle(d_voxels, voxels.size()); + auto* grid = handle.deviceGrid(); + EXPECT_NE(grid, nullptr); + } + ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess); + ASSERT_EQ(cudaFree(d_voxels), cudaSuccess); + EXPECT_GT(base.allocs, 0); // scratch really routed through the sync resource + EXPECT_EQ(base.allocs, base.deallocs); // and every allocation was freed through it +} + TEST(TestMemoryResource, TempPool_FreesOnRetainedStream) { cudaStream_t s = nullptr; diff --git a/pendingchanges/nanovdb.txt b/pendingchanges/nanovdb.txt index cb3a42d13c..f36be8529e 100644 --- a/pendingchanges/nanovdb.txt +++ b/pendingchanges/nanovdb.txt @@ -5,7 +5,7 @@ NanoVDB: - 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. Improvements: - - The GPU builders now allocate all scratch through an injectable memory resource: tools::cuda::TopologyBuilder and tools::cuda::MeshToGrid gained a ResourceT template parameter (defaulted, so existing code is unaffected), joining PointsToGrid. Added nanovdb::cuda::SyncFromAsync, a CRTP base that derives the synchronous half of the resource concept from the stream-ordered half, and nanovdb::cuda::ResourceRef, a non-owning reference to a resource that is itself a resource, for containers that hold their resource by value. PointsToGrid's device scratch and intermediate arrays are now owned by cuda::Buffer as well, replacing all of its hand-paired allocate/free calls. + - The GPU builders now allocate all scratch through an injectable memory resource: tools::cuda::TopologyBuilder and tools::cuda::MeshToGrid gained a ResourceT template parameter (defaulted, so existing code is unaffected), joining PointsToGrid. Added nanovdb::cuda::SyncFromAsync, a CRTP base that derives the synchronous half of the resource concept from the stream-ordered half, and nanovdb::cuda::ResourceRef, a non-owning reference to a resource that is itself a resource, for containers that hold their resource by value. PointsToGrid's device scratch and intermediate arrays are now owned by cuda::Buffer as well, replacing all of its hand-paired allocate/free calls. Added nanovdb::cuda::MallocResource, a synchronous cudaMalloc-backed resource that works on devices without memory-pool support, and nanovdb::cuda::AsyncFromSync, which presents any synchronous resource as a stream-ordered one by synchronizing before each deallocation, so the builders can run with an injected synchronous resource. - Added the CMake option NANOVDB_CUDA_WERROR (default ON, also implied by OPENVDB_CXX_STRICT) which passes --Werror=all-warnings to NVCC so that every device-side diagnostic becomes an error when building the NanoVDB tests, tools and examples. Disable it if a newer CUDA toolkit introduces diagnostics that block your build. - The bug-fix to the nanovdb::ReadAccessor (see below) improves random-access performance in some use-cases (especially on the CPU).