NanoVDB: inject a resource into TopologyBuilder scratch, and align cuda::Buffer with cuda::buffer (CUDA)

[harrism]

Follow-up to #2231 and #2251, part of #2232 (NanoVDB injectable CUDA memory resources). This is B1 of step 2's retrofit; see the roadmap for the B1/B2/B3 split.

Two related changes: the TopologyBuilder scratch retrofit, and the cuda::Buffer naming alignment that the retrofit surfaced.

1. TopologyBuilder scratch from an injected resource

tools::cuda::TopologyBuilder gains a ResourceT template parameter, and its device-only scratch moves from the dual cuda::DeviceBuffer to the single-space cuda::Buffer added in #2251. This is an injectability and ownership change, not a performance one — see below.

An audit of host (.data()) vs device (.deviceData()) use across the builder's ten buffers:

Counts are call sites of .data() (host pointer) and .deviceData() (device pointer) on each member:

member	host uses	device uses	converted
mUpperMasks	0	3	yes
mLowerMasks	0	3	yes
mUpperOffsets	0	5	yes
mLowerOffsets	0	5	yes
mLeafOffsets	0	4	yes
mVoxelOffsets	0	6	yes
mLowerParents	0	2	yes
mLeafParents	0	2	yes
mProcessedRoot	1	1	no — dual-space
mData	1	1	no — dual-space

The eight with zero host uses are converted (plus three function-local buffers, eleven DeviceBuffer::create calls in total). mProcessedRoot and mData are genuinely dual-space and are left alone; they are step 3's problem, when DeviceBuffer is re-implemented as a composition of two cuda::Buffers.

This is not a performance change

An earlier revision of this description claimed the conversion removes a synchronous cudaMallocHost per buffer. That was wrong, and I have measured it. DeviceBuffer::init allocates host memory or device memory, never both:

if (device == cudaCpuDeviceId) {
    cudaCheck(cudaMallocHost((void**)&mCpuData, size));            // host only
} else {
    cudaCheck(util::cuda::mallocAsync(mGpuData+device, size, stream)); // device only
}

TopologyBuilder always passed a real device id from cudaGetDevice(), so every one of these buffers took the else branch. There was no host allocation on this path to remove.

Dilate benchmark, 50 iterations per run, three runs each, RTX 6000 Ada / CUDA 12.6:

points	baseline (ms/dilate)	this PR (ms/dilate)
1,000	4.304 / 4.113 / 4.101	4.322 / 4.091 / 4.095
20,000	4.233 / 4.243 / 4.276	4.226 / 4.232 / 4.210
200,000	5.178 / 5.189 / 5.253	5.207 / 5.215 / 5.211

Within run-to-run noise at every size, as expected once the above is understood.

What it is worth

• Injectability — the eleven allocation sites now route through ResourceT instead of a hard-wired DeviceResource. That is the actual step 2 goal, and TopologyBuilder was the last builder without the seam.
• Ownership — scratch is freed by scope rather than by seven hand-written clear(stream) calls.
• Dropping unused dual-space state — each of these buffers carried a host pointer and a new void*[deviceCount] array it never used.

The builder's TempPool also becomes TempPool<ResourceT>, so cub scratch honors the injected resource too.

Why Data moved. TopologyBuilder::Data is hoisted to namespace scope as TopologyBuilderData<BuildT>, with an in-class alias so existing uses read unchanged. It carries no dependence on the resource, and leaving it nested would give every ResourceT instantiation its own incompatible Data type — which the device functors, templated on BuildT alone, could not name.

Compatibility. ResourceT defaults to nanovdb::cuda::DeviceResource, matching how PointsToGrid took its resource parameter in #2231, so all six existing consumers — DilateGrid, MergeGrids, MeshToGrid, PruneGrid, RefineGrid, CoarsenGrid — compile unchanged. Three of them reach into mBuilder.mUpperMasks / mLowerMasks directly and needed .deviceData() -> .data().

DeviceBuffer::clear(stream) takes a stream where cuda::Buffer frees on its retained stream. These are the same stream on every path: the builder is constructed with the operator's stream and every subsequent call uses that same stream, so the translation preserves ordering exactly.

2. cuda::Buffer naming alignment

Translating those clear(stream) calls surfaced a gap, so this PR also closes it. A member-by-member comparison against cuda::buffer is written up on #2232; the actionable parts:

• destroy() / destroy(stream) — cuda::buffer names the free operation destroy and provides both a retained-stream and an explicit-stream overload. We had only the no-argument form, under a different name, which is why the clear(stream) translation above needed hand-verification. clear() stays as a transitional delegate to destroy(), and is marked as such in the header. It exists only because GridHandle::reset() still calls clear() on its buffer; it goes away in step 3's removal phase, when the legacy dual buffers are deleted and GridHandle moves to destroy(). cuda::Buffer is the type those buffers are being replaced by, so its interface should be the one we want to keep, not an accretion of theirs.
• setStream -> set_stream — member names cannot be aliased, which is the stated reason the resource concept kept allocate_async in snake_case against house style; the same argument applies here.
• swap — the generic std::swap already behaves correctly through the move operations (verified: no allocation, no free, pointers and sizes exchanged), but std::vector and rmm::device_buffer both provide a dedicated one, and the B3 retrofit leans on swapping owners as its central idiom.

Note on set_stream semantics: ours deliberately does not synchronize — and that matches cuda::buffer::set_stream in both name and contract. NVIDIA/cccl#5697 removed the synchronization deliberately ("avoid implicit synchronization in fundamental primitives") but left the earlier "always synchronizes" @note in place; that stale doc is what NVIDIA/cccl#10649 reports. No divergence, no revisit needed.

Testing

Three new tests in TestBuffer.cu cover the added API: destroy() frees and is idempotent; destroy(stream) frees on the given stream rather than the retained one and retains it afterwards; swap exchanges pointers and sizes with no allocation or free. The TopologyBuilder conversion itself is behavior-preserving and is already covered by the existing suite through all six consumers.

Verified locally on an RTX 6000 Ada, CUDA 12.6:

• nanovdb_cuda_buffer_unit_test — 23/23 pass (20 existing + 3 new)
• nanovdb_test_cuda — 52/53 pass. The single failure, UnifiedBuffer_IO, is a missing test-data file (data/3_spheres.nvdb) in an unrelated component; the unmodified baseline fails identically, so it is pre-existing and environmental.

Per #2264 the CUDA tests are excluded from CI (ctest -E ".*cuda.*|.*mgpu.*"), so CI will build this but not run it — the numbers above are from a local GPU run.

Follow-ups

• B2 — TempPool's own bytes onto Buffer, and a SyncFromAsync mixin so a custom resource is two methods rather than four.
• B3 — PointsToGrid's 48 allocation sites, which need an ownership design first (device-visible raw pointers in an uploaded struct, ownership moved by std::swap, a goto retry loop).
• DistributedPointsToGrid stays deferred; NanoVDB: handle empty partitions in the distributed merge path search #2248 should land before any work there.

[harrism]

@kmuseth ready for review — B1 of the #2232 retrofit: TopologyBuilder allocates scratch through an injected resource, plus the cuda::Buffer member-name alignment with cuda::buffer (destroy/set_stream/swap; details and the CCCL comparison are on #2232). All CUDA suites pass on a local GPU — numbers in the description; per #2264 CI builds but does not run them. #2269 stacks on this, so landing order matters.

[harrism]

Ran Copilot review cycles on this PR via a fork mirror (requesting Copilot upstream requires write access). Net changes here: scratch is now released with destroy(stream) at all seven sites rather than relying on the retained stream matching, and four cudaGetDevice calls left dead by the conversion are removed (c7302bf). Alignment concerns it raised were incorrect (Buffer allocates at R::DEFAULT_ALIGNMENT, not alignof(T)) but are now also enforced by a static_assert in the stacked follow-up. Full threads: harrism#1

[kmuseth]

added two requests for changes but otherwise it looks good

[kmuseth]

shouldn't this have a deprecation warning?

[harrism]

Deprecation is part of step 3 of the plan in #2232, and this PR is part of Step 2. Adding [[deprecated]] to this now would cause warnings on internal calls of this function that users have no way to act on. And if they compile with -Werror their TU will fail to compile.

In step 3 we will be removing all internal calls to methods that will be removed, and then deprecate them.

[harrism]

I added Doxygen @deprecated line to the docs to cover this for now.

[kmuseth]

This is another one of these camel-case issues. Why the need for abandoning the camel-case which is used elsewhere in OpenVDB? I understand that CUDA itself uses the other case, but I don't understand why this means we need to adopt it as well. Even worse, currently we use both semantics, which is inconsistent and confusing.

[harrism]

This is as we discussed before, and documented in #2232:

Naming: types are CamelCase per OpenVDB style (cuda::Buffer, cuda::BufferView) — type names are aliasable, so the opt-in using Buffer = ::cuda::buffer<T> at CUDA ≥ 13.2 is preserved. Member names keep the standard spelling (data, size, size_bytes, allocate_async) because member matching is structural and cannot be aliased.

This is analogous to using std::vector rather than building your own vector class.

If we went with setStream then the above using alias will not work in the future. My goal, since long before I started working on VDB, is to maximize compatibility and shareability of resource management across GPU libraries in C++. CCCL is pushing this goal forward by standardizing memory resource and storage interfaces. I would like NanoVDB to be as compatible with that as possible. By matching the exact interface of CCCL types, they become drop-in compatible. So in the future (once our minimum CUDA version is high enough), we can eliminate this code (all buffer types) from our library and just rely on what comes with CUDA.

If you disagree with this change in favor of that goal, let me know.

[harrism]

Fixed the nvcc #177-D warnings in 7d58a33 — TestBuffer.cu now builds with zero warnings. Root cause: the trait static_asserts detect allocate/deallocate through unevaluated contexts, which never odr-use them, so every concept-required member of the file-local test resources was genuinely unreferenced. [[maybe_unused]] turns out not to help — nvcc's front end ignores it for this diagnostic (verified empirically) — so the members are now exercised instead: one test verifies the synchronous halves of the counting/stream-recording doubles behave like their stream-ordered halves, and one pins the trait probes' stub behavior. Real assertions, not just references. The other test TUs build #177-clean as well.

Note on provenance: these warnings are not introduced by this PR — they shipped with #2251 and current master produces the same twelve (verified). Building main shows them today; merging this PR clears them there.

[kmuseth]

thanks for the changes