Repository navigation
NanoVDB: resource forms of the GPU tools, and build() on the builder classes (CUDA) - #2374
Conversation
The tools allocate their output through a prototype buffer, of which they use only the type and resource(), and their scratch through a separate ResourceT template parameter. Add, beside the unchanged prototype forms, overloads of voxelsToGrid, both single-set pointsToGrid forms, indexToGrid and addBlindData that take (..., stream, resource): one resource, held by value (pass a ResourceRef to borrow), allocates both the scratch and the output, which is a GridHandle<cuda::Buffer<std::byte, R>>. Passing a stream selects the resource form, so calls without one resolve to the prototype forms exactly as before. addBlindData drew its scratch from default_resource<ResourceT>() only, so a resource without a default constructor, such as a ResourceRef, could not drive it. Its body moves to a detail function that takes the scratch resource; the prototype form passes the default resource as before. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
The builder classes take a resource for their scratch but allocate their output through a prototype passed to getHandle. Add build() beside getHandle(buffer) on PointsToGrid, DilateGrid, MergeGrids, PruneGrid, RefineGrid, CoarsenGrid, MeshToGrid and IndexToGrid, and buildWithUDF() beside MeshToGrid::getHandleAndUDF(): the output allocates through the builder's own resource, on its stream, as a handle over cuda::Buffer<std::byte, ResourceRef<R>> that borrows it, so one pool serves the scratch and the result. A getHandle(points, n) overload would be ambiguous with the existing defaulted-buffer one, hence the new name. DistributedPointsToGrid is left out: its output must be managed memory every device can reach, which its per-device scratch resources are not. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
The prototype forms stayed viable when a call passed a stream where the prototype goes, with BufferT deduced as cudaStream_t, so the resource form won only through partial ordering, and nvcc's front end found indexToGrid<float, ValueOnIndex>(grid, values, stream) ambiguous. Their return types now leave them out of overload resolution for a cudaStream_t BufferT, and a literal 0 or nullptr in that position fails with a message to write cudaStream_t(0) instead of failing deep inside. Constrain the resource forms on is_async_resource, matching createNodeManager's resource overload. Validate their arguments: a negative point count, a non-positive voxel size or density target, and a null grid throw std::invalid_argument. addBlindData copied the blind-data name into its 256-byte field with an unchecked strcpy, overflowing the host stack for longer names, and both forms dereferenced a null grid on the device; a null or too-long name, a null grid, or null data with a nonzero count now throws. Its public documentation moves back onto the prototype form, and the "buffer (currently ignored)" parameter notes on the builders now say what the prototype does: its resource allocates the output. Headers naming std::byte include <cstddef>. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
addBlindData checked for a null grid only after its first kernel had dereferenced it, so a null grid faulted the device instead of throwing. Check the pointer and name arguments before any launch. The prototype forms now validate the same arguments as the resource forms, through shared helpers, rather than accepting a negative count or a non-positive voxel size that the resource form rejects; indexToGrid's null-grid check no longer exists only as a debug assertion. Define the fixed-voxel-size pointsToGrid prototype form, which was declared but never defined, and test its bytes against the resource form. A synchronous resource passed to a resource form now reaches a static_assert naming cuda::AsyncFromSync, rather than "no matching function": the forms are constrained on being a resource at all, which still keeps an implicitly constructible buffer type out. Comments no longer call the prototype forms unchanged, and the headers touched here include what they use. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
swahtz
left a comment
There was a problem hiding this comment.
The resource forms look like the right direction. Taking a stream to select them, and the error messages for a literal 0 and for a synchronous resource, should make them easy to use correctly. I'm requesting changes for the createNanoGrid gap and the validation that build() skips. The other comments are smaller.
One merge-order note, outside the diff. The multi-set pointsToGrid at PointsToGrid.cuh:1548 still calls pointsToGrid<BuildT, PtrT, BufferT, ResourceT>(...) with four template arguments. The single-set templates take three, so it still fails to compile ("no instance of overloaded function pointsToGrid matches the argument list" with nvcc), even with the missing definition now in place. #2370 fixes the multi-set overloads. If this PR lands first, that form stays broken until #2370 lands.
…orms (CUDA) IndexToGrid throws on a null source grid and PointsToGrid on a non-positive voxel size or maximum point density, in their constructors, so getHandle() and build() reject what the free functions reject, before anything is launched. The free functions keep only the checks the constructors cannot make. createNanoGrid gets the same prototype-form constraint as indexToGrid and a resource form that forwards to indexToGrid's, so a call with a stream selects the resource form under either name. indexToGrid uses the shared PrototypeFormResult alias. The resource form of addBlindData asserts an AsyncResource like the others, and every resource form defaults to a copy of default_resource<ResourceT>(). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
The IndexToGrid null-grid check and the PointsToGrid density check ran in the constructor bodies, after the member initializers had already allocated scratch and recorded on the stream. Both now run in the first initializer, before anything is allocated. The Map constructor now rejects a singular or non-finite map, the voxel-size check rejects infinity, createNanoGrid's prototype form names itself in its stream-argument error, and the PointsToGrid checks all name the class. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
|
Thanks @swahtz. Every inline thread has a reply; the fixes are in 0393979 and 0dc2ddd. On the merge-order note: agreed. #2370 fixes the multi-set overloads and lands ahead of this one in the planned order (#2329 → #2370 → #2371 → #2372 → #2373 → #2374). This PR rebases onto it then, and the duplicate voxelSize |
swahtz
left a comment
There was a problem hiding this comment.
Great, thanks for addressing my notes, lgtm
…eFoundation#2370) into nanovdb-resource-tool-api Both sides defined pointsToGrid(points, count, voxelSize, type, buffer, stream); keep this branch's definition, which adds the prototype-form constraint and the argument checks. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
…esource-tool-api IndexToGrid's constructor keeps this branch's null-source check in its first initializer and takes master's borrowInternal for its scratch. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
…esource-tool-api Both branches added builder tests to TestBuffer.cu. Keep both sections, with AcademySoftwareFoundation#2372's first, and drop this branch's expectSameBytes in favour of AcademySoftwareFoundation#2372's expectSameBytes and hostImage, which it duplicated. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>
What this PR is, and why
Part of the 13.2 work for #2232. The GPU tools allocate their output through a prototype buffer. They use only its type and
resource(), and they take scratch memory through a separateResourceT. This PR adds resource forms beside the prototype forms, which keep their call shapes and are not deprecated.voxelsToGrid, bothpointsToGridforms,indexToGridandaddBlindDatatake(…, stream, resource).ResourceRefto borrow), allocates both the scratch and the output. The output is aGridHandle<cuda::Buffer<std::byte, R>>.cudaStream_tselects this form. When the stream sits where the prototype goes, the prototype forms leave overload resolution, so every compiler resolves the call the same way. (nvcc found one shape ambiguous when this relied on partial ordering.)0there fails with a message to writecudaStream_t(0), and a synchronous resource gets a message namingcuda::AsyncFromSync.build()besidegetHandle(buffer)on PointsToGrid, Dilate, Merge, Prune, Refine, Coarsen, MeshToGrid and IndexToGrid. MeshToGrid also getsbuildWithUDF()besidegetHandleAndUDF().getHandle(points, n)would have been ambiguous with the existing defaulted-buffer overload, hence the new name.Bugs fixed along the way
addBlindDataname overflow. It copied the blind-data name into its 256-byte field with an uncheckedstrcpy, so a longer name overflowed the host stack.ResourceRefcouldn't drive it. The resource form takes the resource explicitly.IndexToGridandPointsToGridconstructors reject a null source grid, a singular or non-finite map, a non-positive or non-finite voxel size, and a non-positive density target, before allocating anything, so the prototype forms, the resource forms,getHandle()andbuild()all reject them. The free functions also reject a negative point count.pointsToGrid(points, n, voxelSize, type, buffer, stream)was declared but never defined, and is now defined. NanoVDB: merge and split single-space device grid handles, and deprecate the dual overloads (CUDA) #2370 defines it too; whichever lands second keeps one copy.Testing
build()produces the same bytes as its prototype counterpart, every allocation goes through the given resource and balances, and the argument checks throw.-Werror.🤖 Generated with Claude Code