[harrism]

Adds CUDA merge and split support for GridHandle<cuda::Buffer<std::byte, R>> and deprecates the dual-space overloads. This fills the gap reported by fvdb in #2232 and prepares for removing the dual-space buffers in 14.0, with a deprecation window starting in 13.2.

The new overloads take (handles, stream, proto), matching cuda::copyTo. They allocate the output through the prototype's resource and wait on each source buffer's retained stream. Grid indices, counts and head checksums are rewritten on the device, one thread per header and one kernel per 256 grids, followed by a single synchronization, so callers can release the sources on return. (Batched renumbering by @swahtz: tools::updateGridCount is now __hostdev__, so device and host renumbering share one function; for 300 grids this is 2 launches where per-grid kernels, readbacks and checksum passes took about 1,200.) The output reuses the source metadata at the new offsets without parsing the grids again on the device. MallocResource is supported on devices without memory pools.

The deprecated overloads name their replacements in the warning. Their implementations move to detail::, following #2301, so existing in-tree callers remain warning-free.

This also fixes two previously untested API problems:

• The point-set overloads of tools::cuda::pointsToGrid and voxelsToGrid did not compile for either buffer type. They now build each set directly and merge with the appropriate implementation, without narrowing size_t to int.
• The documented pointsToGrid(points, n, voxelSize, type, buffer, stream) overload had no definition. It is now implemented.

An audit of the public entry points in tools/cuda/ and cuda/ found no other missing definitions. A compile-and-link probe covers the previously uninstantiated entry points, using both buffer types where applicable.

Validation on one RTX PRO 6000 Blackwell GPU with nvcc 12.6:

• New tests cover merge/split across streams, immediate source release, metadata, full checksums, byte-for-byte split results, balanced resource allocations, and completion on return for synchronous resources.
• A 300-grid merge and split crosses the 256-grid batch boundary, with grids that have no checksum left empty.
• Both buffer types are tested through the point-set APIs, along with the previously undefined pointsToGrid overload.
• All single-GPU CUDA, buffer, memory-resource, and CPU ctest targets pass.
• Examples, tools, and the Python module build with -Werror; all 38 Python tests pass.
• An external caller receives deprecation warnings for both dual-space overloads.

[swahtz]

Before this PR appeared, I'd prototyped single-space merge/split for openvdb/fvdb-core#770. It took a different approach to renumbering the headers. I'm not asking you to switch, but the trade-off might be useful here.

Renumber on the device with no readback. Today tools::updateGridCount is a host-only function, but its only callee, crc32Head(const GridData*), is already __hostdev__. If updateGridCount is also marked __hostdev__, one kernel thread can rewrite a header's index and count and recompute its head checksum in place. The grid pointers can travel by value in the kernel parameters, because 256 pointers take 2 KB and the guaranteed limit is 4 KB. So one launch covers 256 grids without any scratch allocation:

struct GridPtrBatch { static constexpr uint32_t Size = 256; GridData* ptr[Size]; };

static __global__ void updateGridCountBatch(GridPtrBatch batch, uint32_t n, uint32_t firstIndex,
                                            uint32_t indexStep, uint32_t gridCount)
{
    const uint32_t i = threadIdx.x;
    if (i < n) tools::updateGridCount(batch.ptr[i], firstIndex + i * indexStep, gridCount);
}

With indexStep = 1, a merge numbers its chain 0..N-1. With indexStep = 0, a split gives every output index 0.

For N grids the work compares as follows:

	this PR	device-side renumber
Renumber launches	N, one thread each	⌈N/256⌉
Dirty-flag scratch	1 allocation	none
Readback + cudaStreamSynchronize	1	0
Head-checksum refresh	per changed grid: CRC LUT allocation, LUT kernel, crc32Head kernel	done inside the renumber kernel
Merge copies	1 per grid	1 per source handle, since a handle's grids are contiguous

The cost of dropping the sync falls on lifetimes. Without it, the result is stream-ordered on stream, and callers can't release sources on return unless their frees are ordered after stream. That ordering is automatic for sources that retain stream, and it's the caller's job otherwise. A source over a synchronous resource has to stay alive until stream has consumed it. This PR's guarantee that sources can be released on return is simpler, so the change isn't a strict improvement. In exchange, with no sync, merge and split stay legal under stream capture when the resource is stream-ordered. That matters for batched callers like fvdb.

There's also a middle path. You could keep the sync, along with its release-on-return guarantee, and still do the renumbering and head checksums in the batched kernel. That removes the N launches, the per-grid LUT allocations, and the dirty-flag readback, and the contract stays as it is.

The event-based ordering after each source's retained stream in this PR is better than what my prototype did. My version left that ordering to the caller.

[swahtz]

For the paper trail: this addresses item 1 of the 2026-09-17 comment on #2232. #2371 covers item 2, the public adopt API.

[swahtz]

A few optional nits, none blocking:

• Event work in mergeGridHandles and splitGridHandles. The event is created and destroyed on every call, even when BufferT retains no stream and orderAfterRetainedStream does nothing. The merge also records and waits once per source handle, even when every source shares one stream. Guarding with if constexpr (BufferHasStream<BufferT>::value) and skipping streams already waited on would avoid both. This code likely changes anyway with the cross-device fix.
• The grid-count overflow check runs after the stream-ordering loop. Moving it into the metadata loop, or at least above the ordering, means a rejected merge doesn't insert any cross-stream waits first.
• The run-hint comment // make -j 4 testNanoVDB && ... now sits above voxelSetsToDeviceGrids. It belongs to CudaIndexGridToGrid_basic, so the new test should go above the comment.

[harrism]

Thanks @swahtz, all three taken in afa9724:

• Event work: a small detail::RetainedStreamOrdering replaces the per-call event. It creates the event only when a wait is actually needed (never for a buffer type without a retained stream, or when every source already retains the target stream), and waits once per distinct retained stream, so a merge of handles sharing one stream inserts a single wait.
• Grid-count check: now runs before the ordering, so a rejected merge inserts no waits.
• Run-hint comment: back above CudaIndexGridToGrid_basic.

Full single-GPU suite passes.

[swahtz]

Looks good, thanks for addressing my notes