Skip to content

NanoVDB: resource forms of the GPU tools, and build() on the builder classes (CUDA) - #2374

Merged
swahtz merged 9 commits into
AcademySoftwareFoundation:masterfrom
harrism:nanovdb-resource-tool-api
Oct 20, 2026
Merged

swahtz merged 9 commits into
AcademySoftwareFoundation:masterfrom
harrism:nanovdb-resource-tool-api

Conversation

@harrism

@harrism harrism commented Oct 19, 2026 •

Copy link
Copy Markdown
Contributor

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 separate ResourceT. This PR adds resource forms beside the prototype forms, which keep their call shapes and are not deprecated.

  • Free functions: voxelsToGrid, both pointsToGrid forms, indexToGrid and addBlindData take (…, stream, resource).
    • One resource, held by value (pass a ResourceRef to borrow), allocates both the scratch and the output. The output is a GridHandle<cuda::Buffer<std::byte, R>>.
    • Passing a cudaStream_t selects 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.)
    • A literal 0 there fails with a message to write cudaStream_t(0), and a synchronous resource gets a message naming cuda::AsyncFromSync.
  • Builder classes: build() beside getHandle(buffer) on PointsToGrid, Dilate, Merge, Prune, Refine, Coarsen, MeshToGrid and IndexToGrid. MeshToGrid also gets buildWithUDF() beside getHandleAndUDF().
    • The output allocates through the builder's own resource, so one pool serves the scratch and the result, as a buffer that borrows that resource.
    • getHandle(points, n) would have been ambiguous with the existing defaulted-buffer overload, hence the new name.
    • DistributedPointsToGrid is left out: its output must be managed memory that every GPU can reach.

Bugs fixed along the way

  • addBlindData name overflow. It copied the blind-data name into its 256-byte field with an unchecked strcpy, so a longer name overflowed the host stack.
    • Both forms now validate before launching: a null or too-long name, a null grid, or null data with a nonzero count throws.
    • Its scratch came only from the default resource, so a ResourceRef couldn't drive it. The resource form takes the resource explicitly.
  • Argument checks in the builders. The IndexToGrid and PointsToGrid constructors 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() and build() all reject them. The free functions also reject a negative point count.
  • Missing definition. 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

  • New tests: each resource form and each build() produces the same bytes as its prototype counterpart, every allocation goes through the given resource and balances, and the argument checks throw.
  • Platform: one GPU (RTX PRO 6000 Blackwell), nvcc 12.6, -Werror.
  • Results: all single-GPU suites pass. Examples, tools and the Python module build warning-free. Merged together with the resource-equality PR, everything still builds and passes.

🤖 Generated with Claude Code

harrism and others added 4 commits October 16, 2026 05:23
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>
@harrism
harrism marked this pull request as ready for review October 20, 2026 02:53
@harrism
harrism requested a review from kmuseth as a code owner October 20, 2026 02:53

@swahtz swahtz left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

Comment thread nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh Outdated
Comment thread nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh Outdated
Comment thread nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh Outdated
Comment thread nanovdb/nanovdb/tools/cuda/AddBlindData.cuh
Comment thread nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh Outdated
harrism and others added 2 commits October 20, 2026 04:23
…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>
@harrism

harrism commented Oct 20, 2026

Copy link
Copy Markdown
Contributor Author

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 pointsToGrid definition resolves in #2370's favour.

@swahtz swahtz left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Great, thanks for addressing my notes, lgtm

harrism and others added 3 commits October 20, 2026 09:50
…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>
@swahtz
swahtz merged commit 7d9ae18 into AcademySoftwareFoundation:master Oct 20, 2026
16 checks passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants