NanoVDB: migrate off the dual-space buffers and deprecate the DeviceBuffer name (CUDA) - #2301
NanoVDB: migrate off the dual-space buffers and deprecate the DeviceBuffer name (CUDA)#2301harrism wants to merge 15 commits into
Conversation
…age (CUDA) The GPU tools build their result handle through the duck-typed static BufferT::create(size, pool, device, stream), so a single-space buffer type (cuda::Buffer over a memory resource) did not compile at any entry point. The new cuda/HandleStorage.h provides the bridge: createDeviceStorage allocates through create() for buffers that provide it and through the pool buffer's resource otherwise, and deviceStorageData returns the device address either way (data() for single-space, deviceData() for dual). The PointsToGrid family, TopologyBuilder (serving the topology ops and MeshToGrid), IndexToGrid, addBlindData and VoxelBlockManager entry points now allocate through the bridge; addBlindData also drops its dual-space-only static_assert and orders its allocation on the stream it already runs on. The no-argument GridHandle::copy() now names its default-construction requirement in a static_assert instead of failing inside the pool construction, backed by a BufferIsDefaultConstructible detector beside the other trait companions. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
ManagedResource is a synchronous memory resource over cudaMallocManaged, the migration target for UnifiedBuffer users: its allocations are valid on the host and the device, so a handle over cuda::Buffer<std::byte, ManagedResource> parses its metadata on the host and exposes both accessor families over the same allocation. The classification comes from a new DEVICE_ACCESSIBLE resource marker (detected by is_device_accessible_resource and forwarded by ResourceRef and AsyncFromSync, like HOST_ACCESSIBLE): a buffer whose resource is host-accessible AND device-accessible sets both hasHostSingle and hasDeviceSingle, and the handle gates now distinguish device-only single-space buffers (device metadata parse, host accessors are compile-time errors) from both-space ones (host parse, everything available). GridHandle's write/read and NodeManagerHandle follow the same rule. Also adds the four buffer/resource headers missing from NANOVDB_INCLUDE_CUDA_FILES, so installed trees ship cuda/Buffer.h and cuda/PinnedResource.h alongside the new ManagedResource.h and HandleStorage.h. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
…t and device buffers (CUDA) The builder's two remaining dual-space members carried a host side and a device side inside one DeviceBuffer: mProcessedRoot (the consumer tools build the new root topology on the host, then upload) and mData (the builder parameters, written on the host and uploaded once). They are now explicit: pinned host staging (so the uploads stay asynchronous, matching the DeviceBuffer behavior) plus a device buffer allocated through the builder's injected resource, with allocateProcessedRoot / uploadProcessedRoot / uploadData carrying the mechanics once instead of per tool. The six consumers (the five topology ops and MeshToGrid) each lose their DeviceBuffer::create + deviceUpload pair, and the early release in the empty-grid path frees both sides through the new members. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
…s (CUDA) MeshToGrid predated the builder scratch conversions, so its three device members and twelve local scratch buffers still allocated through the dual-space DeviceBuffer, paying for host mirrors nothing read and bypassing the injected resource. They are now cuda::Buffer<std::byte, ResourceRef<ResourceT>> like the other builders, allocated through the resource the class already takes (TopologyBuilder gains a public ref() accessor for consumers sharing its instance), with early releases through destroy() on the retained stream. The UDF sidecar buffer was hard-coded to DeviceBuffer, so the SidecarBufferT template parameter never compiled with any other type; it now allocates through the entry-point bridge, honoring the caller's sidecar prototype buffer. Two ride-alongs from review: PointsToGrid's storage null-check message now names the reachable failure (a buffer type producing no device memory) instead of allocation failure, and IndexToGrid's commented-out synchronization line is deleted. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
…UDA) The root-tile scanline pass interleaves host access with device reads of the same bytes, which is what UnifiedBuffer provided; it now uses cuda::Buffer over ManagedResource. One behavioral difference: the old buffer reserved virtual address space and grew in place, so pointers survived resize; the new buffer reallocates on growth past the 64 reserved tiles, and the tree pointer is re-derived after the resize. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
…nsfers (CUDA) The examples taught the dual-space workflow: build or read into a DeviceBuffer handle, deviceUpload, use both sides, deviceDownload. All 15 CUDA examples now hold the grid in a host handle and move it with cuda::copyTo -- the migration story external code will follow. Device-side image and particle scratch becomes cuda::Buffer with explicit copies for results, and voxels_to_grid demonstrates the full single-space flow: the builder allocates its result directly in a cuda::Buffer, kernels use the device handle, and copyTo brings the grid back for validation. Streams are scoped to outlive the device handles whose buffers free on them. Three pre-existing example defects fixed along the way: index_grid passed its host grid where its kernel launcher expects the device grid, so the example's GPU output was never produced on the GPU; nodemanager used the deprecated cudaCreateNodeManager shim through an output parameter, now the resource-taking overload; and a dead #if 0 alternative is deleted. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
The class is renamed to DualDeviceBuffer -- a transitional implementation name -- and the public DeviceBuffer spelling becomes a deprecated alias whose message carries the migration recipe. The GPU tools' buffer-type defaults name the implementation directly, so code that relies on the defaults compiles without any warning until the defaults change at removal; only code that names DeviceBuffer is warned, and it is the code that has an action available. The legacy CudaDeviceBuffer alias retargets to the implementation so a use fires one warning, not two. The python bindings keep their surface by naming the implementation internally, and the NanoVDB tests keep exercising the dual-space paths -- which still ship and need the coverage -- under translation-unit-local suppression with a comment stating the policy. cuda::Buffer's stream contract now documents the four-line event recipe for handing buffer contents to another stream, which replaces the event tracking DeviceBuffer carried. Also migrates the last device-only DeviceBuffer scratch in the dilate example onto cuda::Buffer, and fixes a pre-existing signature mismatch in the raytrace_level_set OpenVDB comparison path that only an OpenVDB-enabled build would have hit. UnifiedBuffer intentionally keeps its name for now: its replacement story (ManagedResource plus the DistributedPointsToGrid re-plumb) lands in the multi-GPU follow-up in the same release, and the alias treatment applies there identically. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
The topology example drivers are plain .cpp translation units, and including cuda/GridHandle.cuh there fails on every host compiler; they route transfers through helpers in their _kernels.cu companions and hold the returned handles through host-safe headers. Also fixes a pre-existing signature mismatch in raytrace_level_set's OpenVDB comparison path, owns the voxel example's input arrays as cuda::Buffer, and moves the multi-GPU example onto the implementation name it migrates away from with UnifiedBuffer. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
… host-callable (CUDA) A handle-to-handle copy re-validated bytes it had just copied from an already-validated handle: copyTo ended by constructing the destination handle, whose single-space constructor runs the chain-validation kernel. The metadata was validated once, when the source handle was constructed from raw bytes -- so copyTo now adopts it, through a HandleFactory that is the one gateway to building a handle from trusted metadata. With no kernel in the transfer path, copyTo (and the storage helpers) move to the host-includable cuda/HandleStorage.h: a plain C++ translation unit linked against the CUDA runtime can transfer grid handles directly. That deletes the transfer-helper pattern from every example -- the eight host-main programs call copyTo themselves and their .cu companions lose their uploadGrid shims -- and removes the metadata scratch allocation and two synchronizing round trips from every device-destination transfer. The validation kernel keeps guarding the real trust boundary: constructing a handle from raw bytes. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
From the pre-publication review: trailing comments this work added get a space before the slashes (closing-brace labels keep the surrounding files' form); a comment narrating the language rather than the library is deleted; the DualDeviceBuffer class doc states why the type is renamed while its header is not; and the openvdb-to-nanovdb example's host-side print function, deleted one line too eagerly by the helper-removal pass, is restored. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
0b6521d to
06d0d74
Compare
The \file tag named DualDeviceBuffer.h, a file that does not exist: the header deliberately keeps its long-standing name and include path, as its own class note explains. A \file tag that does not match the actual file name detaches doxygen's file page from the header. Signed-off-by: Mark Harris <mharris@nvidia.com>
| /// 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<typename U = BufferT, typename util::disable_if<BufferHasDeviceSingle<U>::value, int>::type = 0> | ||
| template<typename U = BufferT, typename util::disable_if<BufferHasDeviceSingle<U>::value && !BufferHasHostSingle<U>::value, int>::type = 0> |
There was a problem hiding this comment.
💅 polish: We use this BufferHasDeviceSingle<U>::value && !BufferHasHostSingle<U>::value expression 16 times here and 3 times in NodeManager.cuh, we could push this predicate into something named to DRY this up/improve robustness.
There was a problem hiding this comment.
Good call. Named it BufferIsDeviceOnly beside the other detectors in HostBuffer.h and replaced all 19 spellings across the handle headers. Fixed in bfc450d.
| #include <nanovdb/cuda/Buffer.h> | ||
| #include <nanovdb/GridHandle.h> | ||
| #include <nanovdb/cuda/UnifiedBuffer.h> | ||
| #include <nanovdb/cuda/Buffer.h> |
There was a problem hiding this comment.
⌨️ typo: Buffer.h is already included 2 lines above
| /// build or read into a host handle and move it with | ||
| /// cuda::copyTo (see cuda/GridHandle.cuh), or allocate the | ||
| /// result of a GPU tool directly in a cuda::Buffer. Transfers | ||
| /// construct the device handle with a validation kernel, so |
There was a problem hiding this comment.
❓ question: Isn't this description inaccurate? Isn't copyTo kernel-free and host-callable?
There was a problem hiding this comment.
Correct on both counts: the metadata-adoption commit made copyTo kernel-free and host-callable, and this doc predated it. Rewrote the @deprecated block and the alias message; they now point at cuda/HandleStorage.h, where copyTo lives. Fixed in bfc450d (a sibling message in GridHandle::copy caught in e8224fe).
| else srcPtr = src.data(); | ||
| constexpr cudaMemcpyKind kind = srcDev ? (dstDev ? cudaMemcpyDeviceToDevice : cudaMemcpyDeviceToHost) | ||
| : cudaMemcpyHostToDevice; | ||
| cudaCheck(cudaMemcpyAsync(dst.data(), srcPtr, bytes, kind, stream)); |
There was a problem hiding this comment.
🚩 issue: (From Claude) cuda::copyTo returns with the H2D cudaMemcpyAsync still in flight for a stream-less (sync-resource) source such as a pinned-resource Buffer, leaving a use-after-free window when the source handle is destroyed promptly.
There was a problem hiding this comment.
Real, but resolved in the contract rather than with a hidden synchronization: cross-stream ordering stays the caller's, and for a source with a retained stream it is automatic (destruction frees on the same stream as the copy). The gap was the documentation, which claimed the retained-stream overload has "no such requirement" without qualifying stream-less sources: a synchronous-resource (pinned) source frees immediately on destruction, unordered against the still-asynchronous copy. The @warning now states that responsibility explicitly — synchronize the stream before destroying such a source (pageable sources are exempt; their copies degrade to synchronous). Fixed in ec26d38.
| constexpr cudaMemcpyKind kind = srcDev ? (dstDev ? cudaMemcpyDeviceToDevice : cudaMemcpyDeviceToHost) | ||
| : cudaMemcpyHostToDevice; | ||
| cudaCheck(cudaMemcpyAsync(dst.data(), srcPtr, bytes, kind, stream)); | ||
| if constexpr (!dstDev) |
There was a problem hiding this comment.
🚩 issue: (From Claude) cuda::copyTo skips the documented pre-return synchronization for a ManagedResource destination: it is host-readable (hasHostSingle==true) yet classified dstDev==true, so host reads through the returned handle race the pending async copy.
There was a problem hiding this comment.
Confirmed, and this one is a code fix: the documented postcondition is "a host-readable destination synchronizes before returning", and a ManagedResource destination is host-readable (its handle exposes grid() immediately) but was classified by hasDeviceSingle alone. The sync condition is now !dstDev || BufferHasHostSingle<DstBufferT> — pinned and managed destinations synchronize, device-only destinations stay stream-ordered. Added TestBuffer.GridHandleCopyToManagedSynchronizes. Fixed in ec26d38.
| template <typename OtherBufferT> | ||
| inline GridHandle<OtherBufferT> GridHandle<BufferT>::copy(const OtherBufferT& other) const | ||
| { | ||
| static_assert(!(BufferHasDeviceSingle<BufferT>::value || BufferHasDeviceSingle<OtherBufferT>::value), |
There was a problem hiding this comment.
🚩 issue: (From Claude) GridHandle::copy() has no compiling route to a host buffer for ManagedResource handles, and its two error messages point at each other in a circle.
There was a problem hiding this comment.
Half-confirmed: there is a compiling route — cuda::copyTo<HostBuffer>(managedHandle) works (managed handles satisfy the device-single side of its static_assert) — so the messages terminate rather than circle. But copy(pool)'s message did bounce you through copy()'s error to find that out; both messages now name cuda::copyTo and its header directly. Fixed in bfc450d and e8224fe.
| /// @param stream cuda stream | ||
| /// @return Handle that contains a device NodeManager | ||
| template <typename BuildT, typename BufferT = DeviceBuffer> | ||
| template <typename BuildT, typename BufferT = DualDeviceBuffer> |
There was a problem hiding this comment.
❓ question: Why does NodeManager have to use DualDeviceBuffer, is it not possible to use the single-space cuda::Buffer?
There was a problem hiding this comment.
It doesn't have to be — the single-space route is the createNodeManager overload taking a memory resource (directly below this one): it returns a handle over cuda::Buffer<std::byte, ResourceRef<ResourceT>>, and GridStats uses it internally. The DualDeviceBuffer default on this overload is the legacy pool-buffer entry point, kept warning-free until the defaults flip at the removal step — the same policy as the tools' signature defaults. No change here; happy to add a cross-reference in the doc comment if that would help.
…UDA) cuda::copyTo documents that a host-readable destination synchronizes before returning, but the sync condition tested hasDeviceSingle alone, so a ManagedResource destination - host-readable and device-accessible from the same allocation - returned with the copy still in flight and host reads through the handle raced it. The condition now keys on host readability, so pinned and managed destinations synchronize and device-only destinations stay stream-ordered. The copy direction is now cudaMemcpyDefault: every pointer involved is a UVA address, so the runtime infers the direction, as the distributed builder's cross-device copies already rely on. The stream-safety warning now states the source-lifetime contract for sources without a retained stream: a synchronous-resource (pinned) source frees immediately on destruction, unordered against the still- asynchronous copy, so the caller synchronizes the stream before destroying it. Signed-off-by: Mark Harris <mharris@nvidia.com>
BufferIsDeviceOnly names the hasDeviceSingle-and-not-hasHostSingle predicate that gates the handles' host accessors, replacing its 19 hand-spelled repetitions across the handle headers. The DeviceBuffer deprecation text predated the metadata-adopting copyTo: it claimed transfers run a validation kernel and belong in CUDA translation units. copyTo is kernel-free and host-callable, and lives in cuda/HandleStorage.h - the doc, the alias message, and GridHandle::copy(pool)'s error message now say so (the latter pointed at copy(), whose error then pointed at copyTo; it now names the destination directly). Also drops a duplicated include. Signed-off-by: Mark Harris <mharris@nvidia.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
Signed-off-by: Mark Harris <mharris@nvidia.com>
2 participants
What this PR is, and why
This is the migrate-and-deprecate step of the memory-management work under #2232. Grid storage is moving from the dual-space
DeviceBuffer(one object holding a host copy and a device copy, moved withdeviceUpload/deviceDownload) to the single-spacecuda::Bufferwith explicit transfers (cuda::copyTo). This PR migrates everything inside NanoVDB and then deprecates the old name — designed so that nobody is broken and nobody gets a warning they cannot act on:cuda::DeviceBufferbecomes a deprecated alias for the renamed implementation class. The GPU tools' default buffer types name the implementation directly, so code that relies on defaults — most user code — compiles without any warning until the defaults change at removal, one or more releases from now. Only code that spellsDeviceBuffersees the warning, and the message carries the migration recipe.cuda::copyToadopts the source handle's metadata — validated once, when that handle was constructed from raw bytes — so transfers run no kernel and are callable from host-only translation units (cuda/HandleStorage.h). The validation kernel keeps guarding the real trust boundary: constructing a handle from raw bytes.DeviceBufferhandled with internal events, is the documented four-line CUDA-event pattern incuda/Buffer.h— matchingcuda::buffer(CCCL) and RMM.UnifiedBufferintentionally keeps its name here: its replacement (ManagedResource, added in this PR, plus the multi-GPU tool that defaults to it) completes in the multi-GPU follow-up in the same release, with the identical alias treatment.Commits, in review order
create()thatcuda::Bufferdoesn't have; a small bridge (cuda/HandleStorage.h) dispatches tocreate()or the buffer's memory resource. Plus the friendlycopy()diagnostic from the NanoVDB: cuda::copyTo cross-space transfers and pinned-buffer GridHandles (CUDA) #2292 review.ManagedResource— cudaMallocManaged-backed; a handle over it exposes both the host and device accessors over one allocation (the UnifiedBuffer migration target). Also restores four buffer headers missing from the install list.cuda::Buffer(it predated the earlier scratch conversions); fixes its sidecar being hard-coded toDeviceBuffer, which meant theSidecarBufferTparameter never compiled with any other type.ManagedResource(the out-of-bounds read documented in NanoVDB: cuda::SignedFloodFill's root-tile pass reads out of bounds on the device (confirmed with compute-sanitizer) #2300 is pre-existing and unchanged here; its fix belongs to that issue).#if 0)..cpp; transfers route through their.cucompanions (plus an old signature mismatch fixed).copyToadopts metadata and becomes host-callable — deletes the transfer-helper pattern the previous commit needed, removes a scratch allocation and two synchronizing round trips per device transfer (net −59 lines).Testing
Full CPU/CUDA/buffer/memory-resource/mgpu suites green at every step (155+63+46+19+11 on RTX 6000 Ada); example programs run-verified where buildable; python module builds; a two-translation-unit link check confirms the rename preserves the cross-TU constructor instantiation; a host-only probe TU calling
copyTocompiles under g++ and clang.🤖 Generated with Claude Code