NanoVDB: support single-space device buffers in GridHandle (CUDA) - #2288
NanoVDB: support single-space device buffers in GridHandle (CUDA)#2288harrism wants to merge 1 commit into
Conversation
e270bf0 to
2380d21
Compare
2380d21 to
e705df9
Compare
GridHandle can now own a nanovdb::cuda::Buffer<std::byte, R>: a new hasDeviceSingle buffer trait selects a constructor (implemented in cuda/GridHandle.cuh) that parses the grid metadata through the device, with every operation -- the validated header walk, metadata scratch allocation through the buffer's resource, the copy kernel and the readback -- ordered on the buffer's retained stream. deviceData()/deviceGrid() work on such handles. The host accessors (data, grid, gridData, gridMetaData) are SFINAE-removed for them, since the handle owns no host-readable bytes; the read/write I/O members and splitGrids/mergeGrids stay addressable (the python bindings take write's address via overload_cast) and fail instead with an explanatory static_assert when instantiated for a single-space handle. The device parse validates the whole grid chain (per-header bounds checks against the allocation) before launching the metadata kernel, so truncated buffers and forged mGridCount/mGridSize headers are rejected with an exception instead of an out-of-bounds device read; the pre-existing dual-space parse had the same hole and now shares the validation helper. Its metadata scratch uses MallocResource so that long-standing path keeps working on devices without memory-pool support, and the dirty-flag scratch in splitGridHandles and mergeGridHandles is a resource-aware buffer as well -- together replacing the file's six raw cudaMalloc/mallocAsync sites, one of which was missing its cudaCheck. copy() dispatches on the traits: host buffers keep the memcpy path, single-space buffers deep-copy device-to-device through their own resource on the retained stream and reuse the host-resident metadata instead of re-parsing. Buffer gains ElementType/ResourceType typedefs, a resource() accessor, and a no-argument copy() for stream-ordered resources that orders the copy on the retained stream. Host accessibility is a property of the resource, not the element type: the new is_host_accessible_resource trait detects a HOST_ACCESSIBLE marker (declared by PinnedResource and forwarded by ResourceRef and AsyncFromSync). A cuda::Buffer over a host-accessible resource is rejected at handle scope with a named error for now -- GridHandle's host paths require an allocation interface cuda::Buffer does not yet provide -- and non-byte element types fail loudly in the single-space constructor. Both hasDeviceSingle and hasHostSingle are detected rather than required, so pre-existing BufferTraits specializations in and out of tree compile unchanged. Part of AcademySoftwareFoundation#2232 (step 3). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
e705df9 to
c927ae7
Compare
swahtz
left a comment
There was a problem hiding this comment.
Reviewed with a VBM follow-up as part of #2232 in mind. The design lands the right conventions — the detected traits, the named compile errors instead of half-working host paths, and never requiring default-constructible buffers are all things the next handle types can build on directly. Two issues inline: a correctness regression where the hardened chain validation now throws on addBlindData output from multi-grid sources, and the O(N) stream syncs the validation adds to the long-standing GridHandle constructor (with two possible shapes for fixing it). The rest are non-blocking polish.
| /// the device-side metadata walk that follows can never read out of | ||
| /// bounds from a truncated buffer or a forged header. | ||
| /// @return the validated grid count | ||
| inline uint32_t validGridChainCount(const GridData *d_head, uint64_t bytes, cudaStream_t stream) |
There was a problem hiding this comment.
🚩 issue: I traced the consumers of the device-parse path, and I believe this new validation breaks tools::cuda::addBlindData for grids that came out of a multi-grid handle. addBlindData copies the source grid's GridData header verbatim (including mGridIndex/mGridCount) into its new single-grid buffer and never normalizes them before constructing GridHandle<BufferT> (AddBlindData.cuh:131, a device-only buffer, so it takes this path). With a source grid from handle.deviceGrid<T>(n) of a multi-grid handle: for n > 0 the first iteration throws "inconsistent grid index/count" (mGridIndex != 0), and for n == 0 with mGridCount > 1 the walk runs past the single-grid buffer and throws "grid chain exceeds the device buffer". The old parse accepted these buffers (albeit with OOB metadata reads — which is exactly the hole this validation fixes).
I think the fix belongs in addBlindData rather than in weakening the validation: normalize the header to index 0 / count 1 before constructing the handle, the way splitGridHandles does with detail::updateGridCount<<<1,1>>>. Might also be worth a quick audit for any other tool that constructs a device-parse handle from a copied header.
| mMetaData.resize(tmp.mGridCount); | ||
| cudaCheck(cudaMemcpy(mMetaData.data(), d_metaData,tmp.mGridCount*sizeof(GridHandleMetaData), cudaMemcpyDeviceToHost)); | ||
| cudaCheck(cudaFree(d_metaData)); | ||
| const uint32_t count = cuda::detail::validGridChainCount(d_data, mBuffer.size(), cudaStream_t(0)); |
There was a problem hiding this comment.
🚩 issue: This long-standing constructor previously did a single header read; it now does one cudaMemcpyAsync + full cudaStreamSynchronize per grid header inside validGridChainCount. cuda::mergeGridHandles ends by constructing exactly this handle, so merging hundreds of grids now issues hundreds of serialized 672-byte D2H copies, each draining the stream, where the old code did one memcpy plus one kernel.
Two shapes that would fix it, either of which I'd be happy with: (a) since the validation walk already copies every full GridData header to the host, accumulate mMetaData directly in the walk — that deletes the scratch allocation, the cpyGridHandleMeta<<<1,1>>> launch, the second D2H copy, and the final sync in both constructors, so the validation replaces the old parse machinery instead of adding to it (the round trips stay O(N), but the whole second phase disappears); or (b) move the validation into a single device-side kernel with one status readback, restoring O(1) round trips for the multi-grid case.
| /// or if the template parameter does not match the specified grid. | ||
| template<typename ValueT, typename U = BufferT> | ||
| typename util::enable_if<BufferHasDeviceSingle<U>::value, const NanoGrid<ValueT>*>::type | ||
| deviceGrid(uint32_t n=0) const { |
There was a problem hiding this comment.
💅 polish: This body is a token-for-token copy of grid() and the dual-space deviceGrid() — the same null/index/gridType checks and util::PtrAdd, differing only in the base pointer (mBuffer.data() vs mBuffer.deviceData()). A private ungated helper taking the base pointer (e.g. gridAt<ValueT>(const void* base, uint32_t n)) would serve all three, so a future change to the lookup contract doesn't need three edits — this inline copy sits far from the two out-of-line definitions and would be the one that gets missed.
| @@ -371,14 +462,40 @@ template<typename BufferT> | |||
| template <typename OtherBufferT> | |||
| inline GridHandle<OtherBufferT> GridHandle<BufferT>::copy(const OtherBufferT& other) const | |||
There was a problem hiding this comment.
💡 suggestion: Both copy() overloads carry an identical if constexpr guard and multi-line static_assert, and this overload's single-space branch just discards other and forwards to copy<OtherBufferT>() — which re-fires the same assert, so today every misuse emits the diagnostic twice. I believe the guard can live solely in the no-arg copy() with identical diagnostics for every (BufferT, OtherBufferT) combination. This matters for part 2: when cross-space transfers land, the "not supported yet" condition and message get relaxed, and with two copies of the guard the two spellings can silently drift to accepting different type combinations.
| /// BufferTraits specializations (in or out of tree) that only define | ||
| /// hasDeviceDual keep compiling unchanged. | ||
| template<typename BufferT, typename = void> | ||
| struct BufferHasDeviceSingle { static constexpr bool value = false; }; |
There was a problem hiding this comment.
💡 suggestion: Could we move BufferHasDeviceSingle/BufferHasHostSingle next to the BufferTraits primary in HostBuffer.h (or a small shared traits header)? They aren't GridHandle-specific — they're companions to the trait protocol itself — and there's a concrete second consumer lined up: VoxelBlockManagerHandle needs exactly the same dispatch to map deviceFirstLeafID()/deviceJumpMap() onto data() for single-space buffers (the VBM follow-up to this PR under #2232), and it shouldn't have to include GridHandle.h, or duplicate the detectors, to get it.
| /// (stream-ordered resources), and the (stream, resource, count, noInit) | ||
| /// constructor shape. | ||
| template<typename T, typename R> | ||
| struct BufferTraits<cuda::Buffer<T, R>> |
There was a problem hiding this comment.
💡 suggestion: This @note is effectively the definition of a "single-space buffer concept" — it lists exactly the members a buffer must provide for hasDeviceSingle to be honored. Since other handle types will consume the same contract (VoxelBlockManagerHandle is next, per the #2232 plan), could we phrase it as such — name the concept, and state that any consumer of hasDeviceSingle may rely on exactly this interface? That lets the follow-up PRs cite it instead of reverse-engineering the GridHandle constructor.
| public: | ||
| /// @brief Element and resource types, for generic code that rebinds one | ||
| /// or constructs sibling buffers over the same resource. | ||
| using ElementType = T; |
There was a problem hiding this comment.
💅 polish: Tiny one: a rebind alias here (template<class U> using rebind = Buffer<U, R>;) would let generic code construct sibling buffers without spelling the full type — makeMetaScratch spells cuda::Buffer<GridHandleMetaData, ResourceT> manually today, and the VBM follow-up will spell two more.
What this PR is, and why
Grids on the GPU currently require a dual-space buffer: a host allocation mirrored by a device allocation, even in pipelines where the host copy is never read — doubling memory for device-resident workflows and pinning allocation to NanoVDB's built-in allocator instead of the application's pool. This PR (step 3 of #2232, part 1 of 3) adds a device-only alternative:
GridHandle<cuda::Buffer<std::byte, R>>, a grid owned by a single device allocation made through any injected memory resource (an RMM-style pool, or downstream wrappers like fvdb's Torch caching allocator — the consumers the #2232 seam work exists for). Host mirror: gone. Host-facing APIs on such a handle: compile-time errors by design.Dual-space handles are untouched by this PR — today they're still the only way to move grids between host and device. That changes in part 2, which will add explicit cross-space transfers (per-direction
copy()between host-space and device-space handles, currently a named compile error). Once those cover the dual buffers' use cases, the dual buffers themselves will be deprecated and then removed: the end state is one buffer template class, with on type per memory space and explicit transfers between them.What
hasDeviceSinglebuffer trait selects a constructor incuda/GridHandle.cuhthat parses grid metadata through the device — validated header walk, metadata scratch through the buffer's own resource, kernel, and readback all ordered on the buffer's retained stream.deviceData()/deviceGrid<T>()are available;data(),grid<T>(),gridData(),gridMetaData()are SFINAE-removed; theread/writeI/O members andsplitGrids/mergeGridsstay addressable (the python bindings takewrite's address viaoverload_cast) and instead fail with explanatorystatic_asserts when instantiated for a single-space handle. The python module (PyGridHandle.hbindings included) is compiled as part of the local gate.copy()dispatches on the traits: host buffers keep the memcpy path; a single-space handle deep-copies device-to-device through its own resource on the retained stream and reuses the host-resident metadata instead of re-parsing.cuda::BuffergainsElementType/ResourceType,resource(), and a no-argcopy()for stream-ordered resources (=copy(stream())).mGridCount/mGridSizeheaders throw instead of reading out of bounds — the pre-existing dual-space parse had the same hole and now uses the same helper. Its scratch runs onMallocResource, so the long-standingGridHandle<DeviceBuffer>device parse keeps working on devices without memory-pool support. Six raw allocation sites incuda/GridHandle.cuh(one an uncheckedcudaMalloc) are replaced with resource-aware buffers.is_host_accessible_resourcetrait detects aHOST_ACCESSIBLEmarker (onPinnedResource, forwarded throughResourceRef/AsyncFromSync). A pinned-resourcecuda::Bufferhandle is a named compile error for now — GridHandle's host paths need an allocation interfacecuda::Bufferdoesn't yet provide (thecreatemapping, scheduled for the step-3 completion PR) — rather than a half-working host path.BufferHasDeviceSingle/BufferHasHostSingledefault to false when aBufferTraitsspecialization omits them, so pre-existing specializations in tree (PoolBufferexamples,DeviceBuffer,UnifiedBuffer) and out of tree compile unchanged. Every existing call site of the re-gated members was traced; none changes overload resolution.Tests
Seven new tests + compile-time classification asserts in
TestBuffer.cu: device meta parse and typed/wrong-typedeviceGrid; exact allocation accounting through a counting resource (grid bytes + meta scratch, freed viareset()); deep D2D copy with byte comparison; empty-handle and empty-copy; invalid-grid throw with leak check; synchronous-resource (MallocResource) construction and copy; multi-grid parse with a stream-recording resource proving every allocation lands on the retained (non-blocking) stream; forged-mGridCountrejection with an unpoisoned CUDA context.Verification
--Werror=all-warnings; full build (tests, tools, examples) atCMAKE_CUDA_ARCHITECTURES=80; 7/7 ctests on a GPU runner, plusg++ -fsyntax-onlyonGridHandle.hproving the header stays CUDA-free.🤖 Generated with Claude Code