diff --git a/nanovdb/nanovdb/cuda/DeviceResource.h b/nanovdb/nanovdb/cuda/DeviceResource.h
index a18ee789b6..bef41f96d9 100644
--- a/nanovdb/nanovdb/cuda/DeviceResource.h
+++ b/nanovdb/nanovdb/cuda/DeviceResource.h
@@ -145,6 +145,106 @@ struct is_resource<R, std::void_t<
     decltype(std::declval<R&>().deallocate(std::declval<void*>(), size_t{0}, size_t{0}))>>
     : std::true_type {};
 
+/// @brief CRTP base supplying the synchronous half of the resource concept in
+///        terms of the stream-ordered half, so a custom stream-ordered resource
+///        only has to write allocate_async and deallocate_async.
+/// @tparam Derived the resource deriving from this base
+/// @details A stream-ordered resource must also model the synchronous concept
+///          (is_async_resource implies is_resource), which means writing four
+///          methods where two would do. The synchronous pair is not a bare
+///          delegate: memory from allocate must be usable immediately on any
+///          stream, so the null-stream allocation has to be synchronized before
+///          it is returned. Omitting that synchronization yields memory that
+///          satisfies the concept but is not actually synchronous -- a race
+///          rather than a compile error -- so it lives here rather than being
+///          rewritten per resource.
+/// @code
+/// struct MyResource : nanovdb::cuda::SyncFromAsync<MyResource> {
+///     static constexpr size_t DEFAULT_ALIGNMENT = 256;
+///     void* allocate_async(size_t bytes, size_t alignment, cudaStream_t stream);
+///     void  deallocate_async(void* p, size_t bytes, size_t alignment, cudaStream_t stream);
+/// };
+/// @endcode
+template <class Derived>
+struct SyncFromAsync
+{
+    /// @brief Allocates @c bytes usable on any stream when this returns.
+    /// @param bytes number of bytes to allocate
+    /// @param alignment requested alignment
+    /// @note Every call synchronizes the null stream; on hot paths prefer the
+    ///       stream-ordered pair.
+    void* allocate(size_t bytes, size_t alignment)
+    {
+        void* p = static_cast<Derived&>(*this).allocate_async(bytes, alignment, cudaStream_t{0});
+        cudaCheck(cudaStreamSynchronize(cudaStream_t{0}));
+        return p;
+    }
+
+    /// @brief Frees @c p on the null stream.
+    /// @param p pointer previously returned by allocate
+    /// @param bytes size passed to the matching allocate
+    /// @param alignment alignment passed to the matching allocate
+    /// @note No synchronization here: the synchronous concept's contract is
+    ///       that the memory is already quiescent when deallocate is called.
+    void deallocate(void* p, size_t bytes, size_t alignment)
+    {
+        static_cast<Derived&>(*this).deallocate_async(p, bytes, alignment, cudaStream_t{0});
+    }
+};
+
+/// @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
+/// @details Types that hold their resource by value -- cuda::Buffer, matching
+///          cuda::buffer -- select their ownership semantics by what is placed
+///          in that slot: a concrete resource is owned as a copy, while a
+///          ResourceRef borrows. This is the same division cuda::mr draws
+///          between any_resource (owning) and resource_ref (borrowing), and the
+///          same shape as std::pmr::polymorphic_allocator over memory_resource*.
+///          Use it when a resource is stateful or long-lived and a container
+///          must allocate through *that* instance rather than a copy of it.
+/// @warning The referenced resource must outlive every use of this ref and of
+///          all copies of it, including any container holding one.
+template <class R>
+struct ResourceRef
+{
+    static_assert(is_async_resource<R>::value || is_resource<R>::value,
+                  "ResourceRef requires R to model the AsyncResource or the Resource concept");
+
+    static constexpr size_t DEFAULT_ALIGNMENT = R::DEFAULT_ALIGNMENT;
+
+    /// @brief Constructs a ref borrowing @c resource.
+    /// @param resource resource to allocate from; must outlive this ref
+    ResourceRef(R& resource) : mResource(&resource) {}
+
+    /// @{
+    /// @brief Stream-ordered pair, present only when @c R models AsyncResource,
+    ///        so a ref over a synchronous resource does not misreport its tier.
+    template<class S = R, std::enable_if_t<is_async_resource<S>::value, int> = 0>
+    void* allocate_async(size_t bytes, size_t alignment, cudaStream_t stream)
+    {
+        return mResource->allocate_async(bytes, alignment, stream);
+    }
+    template<class S = R, std::enable_if_t<is_async_resource<S>::value, int> = 0>
+    void deallocate_async(void* p, size_t bytes, size_t alignment, cudaStream_t stream)
+    {
+        mResource->deallocate_async(p, bytes, alignment, stream);
+    }
+    /// @}
+
+    /// @brief Synchronous pair, forwarding to the referenced resource.
+    void* allocate(size_t bytes, size_t alignment) { return mResource->allocate(bytes, alignment); }
+    void  deallocate(void* p, size_t bytes, size_t alignment) { mResource->deallocate(p, bytes, alignment); }
+
+    /// @brief Two refs compare equal iff they reference the same resource, i.e.
+    ///        memory allocated through one may be deallocated through the other.
+    friend bool operator==(ResourceRef lhs, ResourceRef rhs) { return lhs.mResource == rhs.mResource; }
+    friend bool operator!=(ResourceRef lhs, ResourceRef rhs) { return lhs.mResource != rhs.mResource; }
+
+private:
+    R* mResource;
+};// ResourceRef<R>
+
 }
 
 } // namespace nanovdb::cuda
diff --git a/nanovdb/nanovdb/cuda/TempPool.h b/nanovdb/nanovdb/cuda/TempPool.h
index 63a97569d2..65317e5f3b 100644
--- a/nanovdb/nanovdb/cuda/TempPool.h
+++ b/nanovdb/nanovdb/cuda/TempPool.h
@@ -10,6 +10,7 @@
 #ifndef NANOVDB_CUDA_TEMPPOOL_H_HAS_BEEN_INCLUDED
 #define NANOVDB_CUDA_TEMPPOOL_H_HAS_BEEN_INCLUDED
 
+#include <nanovdb/cuda/Buffer.h>
 #include <nanovdb/cuda/DeviceResource.h>
 
 #include <cstddef>
@@ -21,31 +22,33 @@ namespace cuda {
 
 template <class Resource>
 class TempPool {
+    static_assert(is_async_resource<Resource>::value,
+                  "TempPool allocates stream-ordered scratch and requires an AsyncResource");
+    // The buffer borrows the pool's resource through a ResourceRef rather than
+    // copying it, preserving the pool's contract that all traffic reaches the
+    // caller's resource instance (which may be stateful).
+    using BufferT = Buffer<std::byte, ResourceRef<Resource>>;
 public:
 
     /// @brief Default c-tor of an empty memory pool that uses the default
     ///        instance of @c Resource for all allocations.
-    TempPool() : mResource(&default_resource<Resource>()), mData(nullptr), mSize(0), mRequestedSize(0), mStream(nullptr) {}
+    TempPool() : TempPool(default_resource<Resource>()) {}
 
     /// @brief C-tor of an empty memory pool that routes all allocations through
     ///        the supplied @c Resource instance.
     /// @param resource resource instance to allocate from; must outlive this pool.
-    explicit TempPool(Resource& resource) : mResource(&resource), mData(nullptr), mSize(0), mRequestedSize(0), mStream(nullptr) {}
-
-    /// @brief Destructor. Frees the managed memory on the stream of the most
-    ///        recent reallocate(), so the stream-ordered free is ordered after
-    ///        the work that used the memory (rather than on the null stream).
-    ~TempPool() {
-        mRequestedSize = 0;
-        mResource->deallocate_async(mData, mSize, Resource::DEFAULT_ALIGNMENT, mStream);
-        mData = nullptr;
-        mSize = 0;
+    explicit TempPool(Resource& resource)
+        : mResource(&resource)
+        , mBuffer(cudaStream_t{0}, ResourceRef<Resource>(resource), 0, noInit)
+    {
     }
 
     /// @brief Returns a non-const void pointer to the data managed by this instance.
-    void* data() {return mData;}
+    void* data() {return mBuffer.data();}
 
     /// @brief Returns a non-const reference to the actual size of the data managed by this instance.
+    /// @note Returned by reference because cub's two-pass API takes the storage
+    ///       size as a size_t&, so this cannot forward Buffer::size() by value.
     size_t& size() {return mSize;}
 
     /// @brief Returns a non-const reference to the requested size of the data managed by this instance.
@@ -54,25 +57,27 @@ class TempPool {
 
     /// @brief Returns the stream that the managed memory was last (re)allocated on,
     ///        i.e. the stream this pool will free on at destruction.
-    cudaStream_t stream() const {return mStream;}
+    cudaStream_t stream() const {return mBuffer.stream();}
 
     /// @brief Re-allocation of the data managed by this instance. Only has affect if the pool in empty or
     ///        the requested memory is larger than the existing size.
     /// @param stream cuda stream used for asynchronous de-allocation and allocation.
+    /// @note Scratch is discarded, never resized: preserving a prefix of
+    ///       temporary storage would be a wasted copy.
     void reallocate(cudaStream_t stream) {
-        if (!mData || mRequestedSize > mSize) {
-            mResource->deallocate_async(mData, mSize, Resource::DEFAULT_ALIGNMENT, stream);
-            mData = mResource->allocate_async(mRequestedSize, Resource::DEFAULT_ALIGNMENT, stream);
-            mSize = mRequestedSize;
+        if (mBuffer.empty() || mRequestedSize > mSize) {
+            mBuffer.destroy(stream);// free the outgrown block on this stream
+            mBuffer = BufferT(stream, ResourceRef<Resource>(*mResource), mRequestedSize, noInit);
+            mSize = mBuffer.size();
+        } else {
+            mBuffer.set_stream(stream);// retained so the d-tor frees on the most-recently-used stream
         }
-        mStream = stream;// retained so the destructor frees on the most-recently-used stream
     }
 private:
-    Resource    *mResource;
-    void        *mData;
-    size_t       mSize;
-    size_t       mRequestedSize;
-    cudaStream_t mStream;
+    Resource *mResource;// non-owning; must outlive this pool and its buffer
+    BufferT   mBuffer;
+    size_t    mSize{0};
+    size_t    mRequestedSize{0};
 };// TempPool<Resource> class
 
 using TempDevicePool = TempPool<DeviceResource>;
diff --git a/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh b/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh
index 81709e39ec..3905c1a107 100644
--- a/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh
+++ b/nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh
@@ -49,9 +49,20 @@ struct Triangle {
     __hostdev__       nanovdb::Vec3f& operator[](int i)       { return v[i]; }
 };
 
-template <typename BuildT>
+/// @brief Pairing of a leaf-node origin with a triangle id. Independent of the
+///        resource the converter allocates from, so it lives outside MeshToGrid
+///        and stays one type across every ResourceT instantiation.
+struct alignas(16) MeshToGridBoxTrianglePair { // sizeof = 16B
+    nanovdb::Coord origin; // 12B
+    uint32_t triangleID;   // 4B
+};
+
+template <typename BuildT, typename ResourceT = nanovdb::cuda::DeviceResource>
 class MeshToGrid
 {
+    static_assert(nanovdb::cuda::is_async_resource<ResourceT>::value,
+                  "MeshToGrid allocates stream-ordered scratch and requires an AsyncResource");
+
     using PointT = nanovdb::Vec3f;
     using TriangleIndexT = nanovdb::Vec3i;
     using TriangleT = Triangle;
@@ -63,10 +74,7 @@ class MeshToGrid
     using LeafT  = NanoLeaf<BuildT>;
 
 public:
-    struct alignas(16) BoxTrianglePair { // sizeof(BoxTrianglePair) = 16B
-        nanovdb::Coord origin; // 12B
-        uint32_t triangleID;   // 4B
-    };
+    using BoxTrianglePair = MeshToGridBoxTrianglePair;
 
     /// @brief Constructor
     /// @param devicePoints Vertex list for input triangle surface
@@ -80,10 +88,11 @@ public:
         const nanovdb::Vec3i *deviceTriangles,
         const uint32_t triangleCount,
         const nanovdb::Map map = nanovdb::Map(),
-        cudaStream_t stream = 0
+        cudaStream_t stream = 0,
+        ResourceT& resource = nanovdb::cuda::default_resource<ResourceT>()
     )
-        : mStream(stream), mTimer(stream), mBuilder(stream), mDevicePoints(devicePoints), mPointCount(pointCount),
-         mDeviceTriangles(deviceTriangles), mTriangleCount(triangleCount), mMap(map)
+        : mStream(stream), mTimer(stream), mBuilder(stream, resource), mDevicePoints(devicePoints), mPointCount(pointCount),
+         mDeviceTriangles(deviceTriangles), mTriangleCount(triangleCount), mMap(map), mTempDevicePool(resource)
     {}
 
     /// @brief Toggle on and off verbose mode
@@ -156,7 +165,7 @@ private:
     static constexpr unsigned int mNumThreads = 128;// for kernels spawned via lambdaKernel (others may specialize)
     static unsigned int numBlocks(unsigned int n) {return (n + mNumThreads - 1) / mNumThreads;}
 
-    TopologyBuilder<BuildT>      mBuilder;
+    TopologyBuilder<BuildT, ResourceT> mBuilder;
     cudaStream_t                 mStream{0};
     std::string                  mGridName;
     util::cuda::Timer            mTimer;
@@ -179,8 +188,8 @@ private:
     auto deviceBoxTrianglePairs()  { return static_cast<BoxTrianglePair*>(mBoxTrianglePairsBuffer.deviceData()); }
     auto deviceUniqueRootOrigins() const { return static_cast<nanovdb::Coord*>(mUniqueRootOriginsBuffer.deviceData()); }
 
-    nanovdb::cuda::TempDevicePool mTempDevicePool;
-}; // tools::cuda::MeshToGrid<BuildT>
+    nanovdb::cuda::TempPool<ResourceT> mTempDevicePool;
+}; // tools::cuda::MeshToGrid<BuildT, ResourceT>
 
 //-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------
 
@@ -201,10 +210,9 @@ private:
 
 //-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------
 
-template<typename BuildT>
+template<typename BuildT, typename ResourceT>
 template<typename BufferT>
-GridHandle<BufferT>
-MeshToGrid<BuildT>::getHandle(const BufferT &buffer)
+GridHandle<BufferT> MeshToGrid<BuildT, ResourceT>::getHandle(const BufferT &buffer)
 {
     cudaStreamSynchronize(mStream);
 
@@ -311,7 +319,7 @@ MeshToGrid<BuildT>::getHandle(const BufferT &buffer)
     }
     if (mVerbose==1) mTimer.stop();
     return handle;
-} // MeshToGrid<BuildT>::getHandle
+} // MeshToGrid<BuildT, ResourceT>::getHandle
 
 //-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------
 
@@ -336,8 +344,8 @@ struct TransformTrianglesFunctor
 
 } // namespace topology::detail
 
-template<typename BuildT>
-void MeshToGrid<BuildT>::transformTriangles()
+template<typename BuildT, typename ResourceT>
+void MeshToGrid<BuildT, ResourceT>::transformTriangles()
 {
     if (mTriangleCount == 0) return;
 
@@ -356,7 +364,7 @@ void MeshToGrid<BuildT>::transformTriangles()
 
     cudaCheckError();
 
-} // MeshToGrid<BuildT>::transformTriangles
+} // MeshToGrid<BuildT, ResourceT>::transformTriangles
 
 //-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------
 
@@ -417,7 +425,7 @@ struct CountRootBoxesFunctor
 template <typename BuildT>
 struct ScatterRootTrianglePairsFunctor
 {
-    using PairT = typename MeshToGrid<BuildT>::BoxTrianglePair;
+    using PairT = MeshToGridBoxTrianglePair;
 
     const Triangle* dXformedTriangles;
     const uint64_t* dOffsets;
@@ -472,8 +480,8 @@ struct ScatterRootTrianglePairsFunctor
 
 } // namespace topology::detail
 
-template<typename BuildT>
-void MeshToGrid<BuildT>::processRootTrianglePairs()
+template<typename BuildT, typename ResourceT>
+void MeshToGrid<BuildT, ResourceT>::processRootTrianglePairs()
 {
     if (mTriangleCount == 0) { mBoxTrianglePairCount = 0; return; }
 
@@ -513,7 +521,7 @@ void MeshToGrid<BuildT>::processRootTrianglePairs()
     // Pass 3: Re-enumerate intersections of (padded) root boxes and triangles, and scatter to allocated list
 
     mBoxTrianglePairsBuffer = nanovdb::cuda::DeviceBuffer::create(
-        mBoxTrianglePairCount * sizeof(typename MeshToGrid<BuildT>::BoxTrianglePair), nullptr, device, mStream);
+        mBoxTrianglePairCount * sizeof(MeshToGridBoxTrianglePair), nullptr, device, mStream);
     if (mBoxTrianglePairsBuffer.deviceData() == nullptr) throw std::runtime_error("Failed to allocate pairs buffer");
 
     util::cuda::lambdaKernel<<<numBlocks(mTriangleCount), mNumThreads, 0, mStream>>>(
@@ -526,7 +534,7 @@ void MeshToGrid<BuildT>::processRootTrianglePairs()
         }
     );
 
-} // MeshToGrid<BuildT>::processRootTrianglePairs
+} // MeshToGrid<BuildT, ResourceT>::processRootTrianglePairs
 
 //-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------
 
@@ -546,7 +554,7 @@ namespace topology::detail {
 template <typename BuildT>
 struct ScatterChildPairsFunctor
 {
-    using PairT = typename MeshToGrid<BuildT>::BoxTrianglePair;
+    using PairT = MeshToGridBoxTrianglePair;
 
     const PairT*            dParents;
     const nanovdb::Mask<3>* dMasks;
@@ -658,7 +666,7 @@ __device__ inline bool testTriangleAABB(
 
 template <typename BuildT, bool OnlyUseAABB>
 __global__ void evaluateAndCountSubBoxesKernel(
-    const typename MeshToGrid<BuildT>::BoxTrianglePair* dParents,
+    const MeshToGridBoxTrianglePair* dParents,
     const Triangle* dXformedTriangles,
     nanovdb::Mask<3>* dMasks,
     uint64_t* dCounts,
@@ -747,7 +755,7 @@ __device__ inline nanovdb::Coord keyToCoord(uint64_t key)
 template <typename BuildT>
 struct EncodeRootOriginsFunctor
 {
-    const typename MeshToGrid<BuildT>::BoxTrianglePair* dPairs;
+    const MeshToGridBoxTrianglePair* dPairs;
     uint64_t*                                           dKeys;
 
     __device__ void operator()(size_t i) const { dKeys[i] = coordToKey(dPairs[i].origin); }
@@ -764,8 +772,8 @@ struct DecodeRootOriginsFunctor
 
 } // namespace topology::detail
 
-template<typename BuildT>
-void MeshToGrid<BuildT>::enumerateRootTiles()
+template<typename BuildT, typename ResourceT>
+void MeshToGrid<BuildT, ResourceT>::enumerateRootTiles()
 {
     if (mBoxTrianglePairCount == 0) return;
 
@@ -817,12 +825,12 @@ void MeshToGrid<BuildT>::enumerateRootTiles()
     );
     cudaCheckError();
 
-} // MeshToGrid<BuildT>::enumerateRootTiles
+} // MeshToGrid<BuildT, ResourceT>::enumerateRootTiles
 
 //-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------
 
-template<typename BuildT>
-void MeshToGrid<BuildT>::buildRasterizedRoot()
+template<typename BuildT, typename ResourceT>
+void MeshToGrid<BuildT, ResourceT>::buildRasterizedRoot()
 {
     int device = 0;
     cudaGetDevice(&device);
@@ -851,12 +859,12 @@ void MeshToGrid<BuildT>::buildRasterizedRoot()
         mBuilder.mProcessedRoot.deviceUpload(device, mStream, false);
         mUniqueRootOriginsBuffer.clear(mStream);
     }
-} // MeshToGrid<BuildT>::buildRasterizedRoot
+} // MeshToGrid<BuildT, ResourceT>::buildRasterizedRoot
 
 //-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------
 
-template<typename BuildT>
-void MeshToGrid<BuildT>::rasterizeInternalNodes()
+template<typename BuildT, typename ResourceT>
+void MeshToGrid<BuildT, ResourceT>::rasterizeInternalNodes()
 {
     if (mBoxTrianglePairCount == 0) return;
 
@@ -871,12 +879,12 @@ void MeshToGrid<BuildT>::rasterizeInternalNodes()
     );
     cudaCheckError();
 
-} // MeshToGrid<BuildT>::rasterizeInternalNodes
+} // MeshToGrid<BuildT, ResourceT>::rasterizeInternalNodes
 
 //-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------
 
-template<typename BuildT>
-void MeshToGrid<BuildT>::processGridTreeRoot()
+template<typename BuildT, typename ResourceT>
+void MeshToGrid<BuildT, ResourceT>::processGridTreeRoot()
 {
     // Initialize grid/tree/root metadata from scratch using the provided map.
     // InitGridTreeRootFunctor sets all GridData fields explicitly (magic, version,
@@ -896,12 +904,12 @@ void MeshToGrid<BuildT>::processGridTreeRoot()
         cudaCheck(cudaMemcpyAsync(dst, mGridName.c_str(), nameSize, cudaMemcpyHostToDevice, mStream));
     }
 
-} // MeshToGrid<BuildT>::processGridTreeRoot
+} // MeshToGrid<BuildT, ResourceT>::processGridTreeRoot
 
 //-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------
 
-template<typename BuildT>
-void MeshToGrid<BuildT>::rasterizeLeafNodes()
+template<typename BuildT, typename ResourceT>
+void MeshToGrid<BuildT, ResourceT>::rasterizeLeafNodes()
 {
     if (mBoxTrianglePairCount == 0) return;
 
@@ -912,12 +920,12 @@ void MeshToGrid<BuildT>::rasterizeLeafNodes()
                       &mBuilder.data()->getGrid(), mBandWidth * mBandWidth });
     cudaCheckError();
 
-} // MeshToGrid<BuildT>::rasterizeLeafNodes
+} // MeshToGrid<BuildT, ResourceT>::rasterizeLeafNodes
 
 //-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------
 
-template<typename BuildT>
-void MeshToGrid<BuildT>::processLeafTrianglePairs()
+template<typename BuildT, typename ResourceT>
+void MeshToGrid<BuildT, ResourceT>::processLeafTrianglePairs()
 {
     if (mBoxTrianglePairCount == 0) return;
 
@@ -1003,7 +1011,7 @@ void MeshToGrid<BuildT>::processLeafTrianglePairs()
         scale /= 8;
     }
 
-} // MeshToGrid<BuildT>::processLeafTrianglePairs
+} // MeshToGrid<BuildT, ResourceT>::processLeafTrianglePairs
 
 //-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------
 
@@ -1042,10 +1050,10 @@ struct FinalizeSidecarFunctor
 
 //-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------
 
-template<typename BuildT>
+template<typename BuildT, typename ResourceT>
 template<typename GridBufferT, typename SidecarBufferT>
 std::pair<GridHandle<GridBufferT>, SidecarBufferT>
-MeshToGrid<BuildT>::getHandleAndUDF(const GridBufferT& buffer, const SidecarBufferT&)
+MeshToGrid<BuildT, ResourceT>::getHandleAndUDF(const GridBufferT& buffer, const SidecarBufferT&)
 {
     cudaStreamSynchronize(mStream);
 
@@ -1176,7 +1184,7 @@ MeshToGrid<BuildT>::getHandleAndUDF(const GridBufferT& buffer, const SidecarBuff
     cudaStreamSynchronize(mStream);
 
     return { std::move(handle), std::move(sidecarBuffer) };
-} // MeshToGrid<BuildT>::getHandleAndUDF
+} // MeshToGrid<BuildT, ResourceT>::getHandleAndUDF
 
 //-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------
 
diff --git a/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh b/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh
index 13eb4714f4..9159002927 100644
--- a/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh
+++ b/nanovdb/nanovdb/tools/cuda/TopologyBuilder.cuh
@@ -56,11 +56,18 @@ class TopologyBuilder
     using LowerT = NanoLower<BuildT>;
     using LeafT  = NanoLeaf<BuildT>;
 
-    /// @brief Device-only scratch storage, allocated from the injected
-    ///        resource. These buffers are never read on the host, so they use
-    ///        the single-space Buffer rather than the dual DeviceBuffer, whose
-    ///        host pointer and per-device array they would leave unused.
-    using ScratchT = nanovdb::cuda::Buffer<std::byte, ResourceT>;
+    static_assert(nanovdb::cuda::is_async_resource<ResourceT>::value,
+                  "TopologyBuilder allocates stream-ordered scratch and requires an AsyncResource");
+    static_assert(ResourceT::DEFAULT_ALIGNMENT >= alignof(uint64_t),
+                  "TopologyBuilder reinterprets byte scratch as word-sized types and requires word-aligned allocations");
+
+    /// @brief Device-only scratch storage, borrowing the injected resource
+    ///        through a ResourceRef so all traffic reaches the caller's
+    ///        instance (which may be stateful) rather than a copy. These
+    ///        buffers are never read on the host, so they use the single-space
+    ///        Buffer rather than the dual DeviceBuffer, whose host pointer and
+    ///        per-device array they would leave unused.
+    using ScratchT = nanovdb::cuda::Buffer<std::byte, nanovdb::cuda::ResourceRef<ResourceT>>;
 
 public:
 
@@ -68,7 +75,15 @@ public:
     /// @param resource resource instance all device scratch is allocated from;
     ///        must outlive this builder
     TopologyBuilder(cudaStream_t stream, ResourceT& resource = nanovdb::cuda::default_resource<ResourceT>())
-        : mResource(&resource)
+        : mUpperMasks(stream, resource, 0, nanovdb::cuda::noInit)
+        , mLowerMasks(stream, resource, 0, nanovdb::cuda::noInit)
+        , mUpperOffsets(stream, resource, 0, nanovdb::cuda::noInit)
+        , mLowerOffsets(stream, resource, 0, nanovdb::cuda::noInit)
+        , mLeafOffsets(stream, resource, 0, nanovdb::cuda::noInit)
+        , mVoxelOffsets(stream, resource, 0, nanovdb::cuda::noInit)
+        , mLowerParents(stream, resource, 0, nanovdb::cuda::noInit)
+        , mLeafParents(stream, resource, 0, nanovdb::cuda::noInit)
+        , mResource(&resource)
         , mTempDevicePool(resource)
     {
         mData = nanovdb::cuda::DeviceBuffer::create(sizeof(Data));
diff --git a/nanovdb/nanovdb/unittest/TestBuffer.cu b/nanovdb/nanovdb/unittest/TestBuffer.cu
index 1ba042193f..2978539b39 100644
--- a/nanovdb/nanovdb/unittest/TestBuffer.cu
+++ b/nanovdb/nanovdb/unittest/TestBuffer.cu
@@ -302,6 +302,102 @@ TEST(TestBuffer, ClearFreesAndEmpties)
     ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess);
 }
 
+// A custom stream-ordered resource written the short way: two methods plus the
+// mixin, rather than four.
+struct MixinResource : nanovdb::cuda::SyncFromAsync<MixinResource>
+{
+    static constexpr size_t DEFAULT_ALIGNMENT = nanovdb::cuda::DeviceResource::DEFAULT_ALIGNMENT;
+    Counters* counters = nullptr;
+    void* allocate_async(size_t bytes, size_t alignment, cudaStream_t stream) {
+        void* p = nanovdb::cuda::DeviceResource{}.allocate_async(bytes, alignment, stream);
+        if (p) { ++counters->allocs; counters->allocBytes = bytes; }
+        return p;
+    }
+    void deallocate_async(void* p, size_t bytes, size_t alignment, cudaStream_t stream) {
+        if (p) { ++counters->deallocs; counters->deallocBytes = bytes; }
+        nanovdb::cuda::DeviceResource{}.deallocate_async(p, bytes, alignment, stream);
+    }
+};
+
+// The mixin supplies the synchronous half, so both concepts are satisfied.
+static_assert(nanovdb::cuda::is_async_resource<MixinResource>::value,
+              "SyncFromAsync user must still model AsyncResource");
+static_assert(nanovdb::cuda::is_resource<MixinResource>::value,
+              "SyncFromAsync must supply the synchronous half of the concept");
+
+TEST(TestBuffer, SyncFromAsyncSuppliesTheSynchronousPair)
+{
+    Counters c;
+    MixinResource r{{}, &c};
+    // the inherited synchronous pair routes through the derived async methods
+    void* p = r.allocate(1024, MixinResource::DEFAULT_ALIGNMENT);
+    ASSERT_NE(p, nullptr);
+    EXPECT_EQ(c.allocs, 1);
+    r.deallocate(p, 1024, MixinResource::DEFAULT_ALIGNMENT);
+    EXPECT_EQ(c.deallocs, 1);
+    ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess);
+}
+
+TEST(TestBuffer, BufferWorksOverAMixinResource)
+{
+    Counters c;
+    {
+        nanovdb::cuda::Buffer<float, MixinResource> buf(0, MixinResource{{}, &c}, 64, nanovdb::cuda::noInit);
+        EXPECT_EQ(c.allocs, 1);
+        EXPECT_NE(buf.data(), nullptr);
+    }
+    EXPECT_EQ(c.deallocs, 1);
+    ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess);
+}
+
+// State held inline, not behind a pointer: copying such a resource strands its
+// accounting, which is exactly what ResourceRef exists to avoid.
+struct StatefulInlineResource : nanovdb::cuda::SyncFromAsync<StatefulInlineResource>
+{
+    static constexpr size_t DEFAULT_ALIGNMENT = nanovdb::cuda::DeviceResource::DEFAULT_ALIGNMENT;
+    int allocs = 0, deallocs = 0;
+    void* allocate_async(size_t bytes, size_t alignment, cudaStream_t stream) {
+        ++allocs; return nanovdb::cuda::DeviceResource{}.allocate_async(bytes, alignment, stream);
+    }
+    void deallocate_async(void* p, size_t bytes, size_t alignment, cudaStream_t stream) {
+        ++deallocs; nanovdb::cuda::DeviceResource{}.deallocate_async(p, bytes, alignment, stream);
+    }
+};
+
+// A ref over an async resource models both tiers; over a synchronous-only
+// resource it models only the synchronous one -- it must not misreport.
+static_assert(nanovdb::cuda::is_async_resource<nanovdb::cuda::ResourceRef<nanovdb::cuda::DeviceResource>>::value,
+              "ref over an async resource must model AsyncResource");
+static_assert(nanovdb::cuda::is_resource<nanovdb::cuda::ResourceRef<nanovdb::cuda::DeviceResource>>::value,
+              "ref over an async resource must model Resource");
+static_assert(!nanovdb::cuda::is_async_resource<nanovdb::cuda::ResourceRef<nanovdb::cuda::PinnedResource>>::value,
+              "ref over a synchronous resource must not claim AsyncResource");
+static_assert(nanovdb::cuda::is_resource<nanovdb::cuda::ResourceRef<nanovdb::cuda::PinnedResource>>::value,
+              "ref over a synchronous resource must model Resource");
+
+TEST(TestBuffer, ResourceRefSharesTheUnderlyingResource)
+{
+    StatefulInlineResource res;   // the original; a by-value copy would strand these counters
+    {
+        nanovdb::cuda::Buffer<int, nanovdb::cuda::ResourceRef<StatefulInlineResource>> buf(
+            cudaStream_t{0}, nanovdb::cuda::ResourceRef<StatefulInlineResource>(res), 1024, nanovdb::cuda::noInit);
+        EXPECT_NE(buf.data(), nullptr);
+        EXPECT_EQ(res.allocs, 1);   // traffic reaches the original, not a copy
+        EXPECT_EQ(res.deallocs, 0);
+    }
+    ASSERT_EQ(cudaStreamSynchronize(0), cudaSuccess);
+    EXPECT_EQ(res.allocs, 1);
+    EXPECT_EQ(res.deallocs, 1);
+}
+
+TEST(TestBuffer, ResourceRefEqualityIsIdentity)
+{
+    StatefulInlineResource a, b;
+    nanovdb::cuda::ResourceRef<StatefulInlineResource> ra(a), raAgain(a), rb(b);
+    EXPECT_TRUE(ra == raAgain);   // same underlying resource
+    EXPECT_TRUE(ra != rb);        // different underlying resources
+}
+
 TEST(TestBuffer, DestroyIsTheSpellingClearDelegatesTo)
 {
     Counters c;
diff --git a/nanovdb/nanovdb/unittest/TestMemoryResource.cu b/nanovdb/nanovdb/unittest/TestMemoryResource.cu
index 136e409323..ed5591055b 100644
--- a/nanovdb/nanovdb/unittest/TestMemoryResource.cu
+++ b/nanovdb/nanovdb/unittest/TestMemoryResource.cu
@@ -25,7 +25,7 @@ namespace {
 
 /// @brief Resource that counts (non-null) allocations and deallocations so
 ///        leaks can be asserted. Delegates the actual work to DeviceResource.
-struct CountingResource
+struct CountingResource : nanovdb::cuda::SyncFromAsync<CountingResource>
 {
     static constexpr size_t DEFAULT_ALIGNMENT = nanovdb::cuda::DeviceResource::DEFAULT_ALIGNMENT;
     int allocs = 0;
@@ -43,7 +43,7 @@ struct CountingResource
 
 /// @brief Resource that records the stream of every allocation/deallocation,
 ///        to verify stream-ordered teardown. Delegates work to DeviceResource.
-struct StreamRecordingResource
+struct StreamRecordingResource : nanovdb::cuda::SyncFromAsync<StreamRecordingResource>
 {
     static constexpr size_t DEFAULT_ALIGNMENT = nanovdb::cuda::DeviceResource::DEFAULT_ALIGNMENT;
     std::vector<cudaStream_t> allocStreams;
diff --git a/pendingchanges/nanovdb.txt b/pendingchanges/nanovdb.txt
index f342f5c47a..bf0e59fce6 100644
--- a/pendingchanges/nanovdb.txt
+++ b/pendingchanges/nanovdb.txt
@@ -5,6 +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.
   - The bug-fix to the nanovdb::ReadAccessor (see below) improves random-access performance in some use-cases (especially on the CPU).
 
   Fixes: