NanoVDB: run the builders on a synchronous memory resource (CUDA)

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<R> 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<class R>
+struct AsyncFromSync
+{
+    static_assert(is_resource<R>::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<R>
+
 /// @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

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<nanovdb::cuda::MallocResource>::value,
+              "MallocResource must model the synchronous Resource concept");
+static_assert(!nanovdb::cuda::is_async_resource<nanovdb::cuda::MallocResource>::value,
+              "MallocResource must not claim the AsyncResource concept");
+static_assert(nanovdb::cuda::is_async_resource<
+                  nanovdb::cuda::AsyncFromSync<nanovdb::cuda::MallocResource>>::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<SyncCountingResource>;
+    using VgpuT = nanovdb::cuda::AsyncFromSync<RefT>;
+    SyncCountingResource base;
+    VgpuT res{RefT(base)};
+
+    const std::vector<nanovdb::Coord> 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<nanovdb::ValueOnIndex, VgpuT> converter(nanovdb::Map(1.0), cudaStream_t{0}, res);
+        auto handle = converter.getHandle(d_voxels, voxels.size());
+        auto* grid = handle.deviceGrid<nanovdb::ValueOnIndex>();
+        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;

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).