[copilot mirror] B2: SyncFromAsync, ResourceRef, MeshToGrid + TempPool - #2
[copilot mirror] B2: SyncFromAsync, ResourceRef, MeshToGrid + TempPool#2harrism wants to merge 5 commits into
Conversation
A stream-ordered resource must also model the synchronous concept, which means writing four methods where two would do. The synchronous pair is not a bare delegate -- memory from allocate must be usable on any stream when it returns, so the null-stream allocation has to be synchronized first -- and omitting that yields memory which satisfies the concept but is not actually synchronous. Put it in one place rather than leaving each author to rediscover it. The two resources in TestMemoryResource are the first users, and were already wrong in exactly that way: they provide only the async pair, so they never modelled is_async_resource. TempPool duck-typed and never checked, so nothing caught it. MeshToGrid was the last builder allocating from a hard-wired DeviceResource, through TempDevicePool. Give it a ResourceT parameter and thread it into both its TopologyBuilder and its pool. As with Data in the builder, BoxTrianglePair is hoisted out of the class: it does not depend on the resource, and leaving it nested would give every ResourceT its own incompatible type for the device functors to name. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
…uffer Buffer holds its resource by value, matching cuda::buffer -- whose model this completes: in CCCL the ownership semantics are selected by what is placed in the by-value slot, an owning any_resource or a borrowing resource_ref. We adopted the slot without the borrowing type, so a container like TempPool, whose contract is a non-owning pointer to a possibly stateful resource, had no way to hold a Buffer without copying that resource and stranding its state. ResourceRef is the missing piece: a non-owning reference that is itself a resource, so copying the ref shares the underlying instance. Its async methods exist only when R models AsyncResource, so a ref over a synchronous resource does not misreport its tier, and two refs compare equal exactly when they reference the same resource. TempPool now keeps its bytes in a Buffer<std::byte, ResourceRef<R>>: same resource contract, same stream retention, same discard-on-growth reallocation, but the block is freed by ownership rather than by hand. The TempPool unit tests, which assert traffic against the caller's own resource instance, pass unchanged. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
There was a problem hiding this comment.
Pull request overview
Warning
Copilot couldn't run its full agentic review because it didn't start before the timeout. Make sure your repository has a runner available, or add a copilot-code-review.yml file specifying one with the runs-on attribute. See the docs for more details.
This PR mirrors upstream B2 changes to make NanoVDB CUDA tooling allocate scratch memory via injectable memory resources, adding helpers to reduce boilerplate for stream-ordered resources and avoid accidental resource copying.
Changes:
- Added
nanovdb::cuda::SyncFromAsync(CRTP) to derive synchronous allocate/deallocate from async methods, andnanovdb::cuda::ResourceRefto borrow resources without copying. - Updated CUDA tests to cover the new resource utilities (
SyncFromAsync,ResourceRef). - Templatized
tools::cuda::MeshToGridto accept an injectableResourceTand switched scratch allocation to the newTempPool<ResourceT>.
Reviewed changes
Copilot reviewed 6 out of 6 changed files in this pull request and generated 5 comments.
Show a summary per file
| File | Description |
|---|---|
| pendingchanges/nanovdb.txt | Documents the new injectable-resource and helper types. |
| nanovdb/nanovdb/unittest/TestMemoryResource.cu | Updates test resources to use SyncFromAsync. |
| nanovdb/nanovdb/unittest/TestBuffer.cu | Adds unit tests for SyncFromAsync and ResourceRef. |
| nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh | Adds ResourceT template + routes scratch through TempPool<ResourceT>. |
| nanovdb/nanovdb/cuda/TempPool.h | Reimplements TempPool on top of Buffer + ResourceRef. |
| nanovdb/nanovdb/cuda/DeviceResource.h | Introduces SyncFromAsync and ResourceRef. |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
| /// @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;} |
There was a problem hiding this comment.
Intentional, and pre-existing API: size() returned size_t& before this change for the same reason — cub's two-pass API writes the required size through that reference on the query pass and only reads it on the execute pass, so it never desynchronizes in the CALL_CUBS pattern. Renaming the accessor is a public-API question outside this conversion; the doc comment now states why it is a reference.
| 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 | ||
| } |
There was a problem hiding this comment.
Deliberate: the pre-conversion pool freed the old block before allocating the new one, and preserving that order keeps peak device memory unchanged at growth — allocate-then-swap would briefly hold old+new. The throwing path is currently unreachable: allocation failure exits via cudaCheck rather than throwing (see AcademySoftwareFoundation#2265 for making that throw; if it lands, this ordering question gets revisited there), and checkedBytes cannot overflow for byte buffers.
| template <typename BuildT, typename ResourceT = nanovdb::cuda::DeviceResource> | ||
| class MeshToGrid |
| 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; | ||
| } |
There was a problem hiding this comment.
Doc note added in f619056: every allocate call synchronizes the null stream; hot paths should prefer the stream-ordered pair. The suggested micro-optimizations are skipped: on the current error model a failed allocation exits rather than returning null (so sync-on-success is the same thing), and a zero-byte fast path optimizes a degenerate case nobody has measured — leaving the semantics uniform is worth more.
| explicit TempPool(Resource& resource) | ||
| : mResource(&resource) | ||
| , mBuffer(cudaStream_t{0}, ResourceRef<Resource>(resource), 0, noInit) | ||
| { | ||
| } |
There was a problem hiding this comment.
These tests exist and gated this change: TestMemoryResource.TempPool_FreesOnRetainedStream (StreamRecordingResource; asserts the free lands on the most-recently-used stream) and TempPool_NoLeakAcrossGrowth (CountingResource; asserts every allocation reaches and is freed through the caller's own instance). They predate this PR, pass unchanged against it, and are the reason an earlier by-value version of this conversion was caught and rewritten. They are outside this diff, which is why they do not appear here.
Signed-off-by: Mark Harris <mharris@nvidia.com>
The scratch buffers held their resource by value, so each of the eight carried its own copy -- fine for the stateless default, wrong for a stateful resource, whose accounting would be split across copies while the caller's instance saw nothing. Borrow through ResourceRef instead, the same reconciliation TempPool uses. Assert the stream-ordered requirement directly in TopologyBuilder and MeshToGrid so a synchronous-only resource fails with a diagnostic that names the builder, not just the pool inside it. Note SyncFromAsync's synchronize cost on its allocate. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
e3c8930 to
f619056
Compare
There was a problem hiding this comment.
Pull request overview
Copilot reviewed 7 out of 7 changed files in this pull request and generated 1 comment.
Suppressed comments (1)
nanovdb/nanovdb/unittest/TestBuffer.cu:302
MixinResource{{}, &c}relies on aggregate initialization, butMixinResourcehas a base class (SyncFromAsync) and the project minimum is C++17. Create a localMixinResource, setcounters, and pass that instance toBufferinstead.
{
nanovdb::cuda::Buffer<float, MixinResource> buf(0, MixinResource{{}, &c}, 64, nanovdb::cuda::noInit);
EXPECT_EQ(c.allocs, 1);
| Counters c; | ||
| MixinResource r{{}, &c}; | ||
| // the inherited synchronous pair routes through the derived async methods |
There was a problem hiding this comment.
Not a defect: C++17 extended aggregates to include public base classes (P0017R1), so MixinResource is an aggregate and {{}, &c} is well-formed C++17 — it initializes the empty SyncFromAsync base and then counters. Verified compiling with gcc and nvcc at -std=c++17 and passing at runtime; MSVC's C++17 mode supports the same.
The byte scratch is reinterpreted as word-sized types, which is valid for every resource whose DEFAULT_ALIGNMENT is at least word alignment -- all CUDA allocation paths give 256 -- but nothing said so. Assert it, so a custom resource with a weaker guarantee fails at compile time instead of misaligning on the device. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
|
Copilot review cycles converged (final cycle: no new findings). All accepted changes are on the shared branch and reflected in the upstream PR; threads here document the dismissals. Branch retained. |
Copilot-review mirror of AcademySoftwareFoundation#2269, based on the B1 branch so only B2's payload shows.