NanoVDB: deprecate the dual-space handle transfers (CUDA) by harrism · Pull Request #2378 · AcademySoftwareFoundation/openvdb

harrism · 2026-10-21T00:32:12Z

Deprecates the GridHandle and NodeManagerHandle members that disappear with the dual-space buffers in 14.0. Users of the GPU tools' default buffer then get a warning in 13.2, before the 14.0 default change breaks their code.

The GPU tools default to cuda::DualDeviceBuffer, which is not deprecated by name. #2301 deprecated only the cuda::DeviceBuffer alias, so code like the example below never names a deprecated type and compiles cleanly today:

auto handle = nanovdb::tools::cuda::voxelsToGrid<float>(d_coords, n); // default buffer
handle.deviceDownload();
auto* grid = handle.grid<float>();

In 14.0 the default becomes cuda::Buffer<std::byte>, which holds the grids on the device only, and that code stops compiling at deviceDownload. This PR deprecates the members that won't exist then, and each warning names a replacement that exists today:

Deprecated	Replacement
`GridHandle::deviceUpload`	build a host handle and `cuda::copyTo` a device handle
`GridHandle::deviceDownload`	`cuda::copyTo` a host handle
`GridHandle::deviceData(int device)`	one device handle per GPU, or a managed `cuda::Buffer<std::byte, cuda::ManagedResource>`
`NodeManagerHandle::deviceUpload` / `deviceDownload`	`cuda::createNodeManager(deviceGrid, resource)` on the device, or `createNodeManager` on the host
`DualDeviceBuffer::deviceUpload` / `deviceDownload` (reached via `handle.buffer()`)	`cuda::Buffer<std::byte>` with `cuda::copyTo`
`DualUnifiedBuffer::deviceUpload` / `deviceDownload`	`cuda::Buffer<std::byte, cuda::ManagedResource>`

For multi-GPU handles, such as those DistributedPointsToGrid builds, the GridHandle messages also name cuda::Buffer<std::byte, cuda::ManagedResource>, since a device-only buffer can't replace them. The deprecated handle members forward to the buffer members inside a new NANOVDB_SUPPRESS_DEPRECATION_BEGIN/END pair, so a caller gets one warning at its own line and none from inside the headers. TestBuffer.DeprecatedDualTransfersWarnOnlyAtTheCall fails the -Werror build if one escapes.

Members that keep working after 14.0, such as deviceData() and deviceGrid<T>(), are not deprecated. Any host read of a device-built dual handle, or device use of a host-built one, has to go through upload or download, so those calls are where every affected caller gets the warning. A probe that never names a buffer type gets warnings at exactly those calls, whether NanoVDB is included with -I or with -isystem.

In-tree callers are migrated:

The GPU builder tests in TestNanoVDB.cu build into cuda::Buffer<std::byte> and read their results through cuda::copyTo. Every grid-content assertion is kept; checks that only described dual-buffer behaviour ("no host copy yet") were dropped.
The multi-GPU and memory-resource tests use single-space storage.
Tests whose subject is the dual-space buffers themselves still use them, and the paths that stay the default in 13.x keep a test each: the dual cuda::createNodeManager overload, the dual buffer's upload bookkeeping, and DistributedPointsToGrid::getHandle's default DualUnifiedBuffer.
The TEACHME lesson, cheat sheet and capstone now teach cuda::Buffer<std::byte> with cuda::copyTo, including the resource forms and build() from NanoVDB: resource forms of the GPU tools, and build() on the builder classes (CUDA) #2374. They also replace the deprecated math::ZeroCrossing.

Python is unchanged. DeviceGridHandle.deviceUpload/deviceDownload keep working without a warning; the binding calls the buffer's transfers with the deprecation suppressed locally. Python has no single-space API to move to yet, so a warning would leave users nothing to change their code to. How the Python side moves off the dual buffers before 14.0 is to be agreed with @swahtz alongside #2225/#2337.

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

All single-GPU ctest targets pass. Unit tests, tools, examples and the Python module build with -Werror and no new warnings. All 38 Python tests pass.
All 64 TEACHME snippets compile with no deprecation warnings, and both capstone programs build and render.

#2377 fixes a header include-order bug that the new docs' #include <nanovdb/cuda/HandleStorage.h> would otherwise run into. This PR doesn't depend on it.

🤖 Generated with Claude Code

The GPU tools default to the dual-space DualDeviceBuffer, which is not itself deprecated, so code that relies on that default never names a deprecated type. In 14.0 the default becomes cuda::Buffer<std::byte>, which holds the grids on the device only, and such code stops compiling at its deviceDownload/deviceUpload calls without having seen a warning. Deprecate the members that disappear with the dual-space buffers: GridHandle::deviceUpload, deviceDownload and deviceData(int device), and NodeManagerHandle::deviceUpload and deviceDownload. Each warning names its replacement (cuda::copyTo, cuda::createNodeManager, or one handle per device). Members that survive the switch, such as deviceData() and deviceGrid<T>(), are unchanged. Migrate the in-tree callers instead of suppressing the warnings: the GPU builder tests now build into cuda::Buffer<std::byte> and read the grids through cuda::copyTo, and the multi-GPU tests use a managed cuda::Buffer. Tests whose subject is the dual-space buffers keep using them. The Python DeviceGridHandle binding calls the buffer's transfers directly; its Python API is unchanged. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>

@deprecated

… (CUDA) Name the resource-taking cuda::createNodeManager overload in the NodeManagerHandle::deviceUpload warning (the default overload still returns a dual-space handle), give the headers for the managed-buffer replacement, and tag the deprecated members with @deprecated in their docs. The cuda::copyTo dual-space assertion now says how to build a single-space device handle. In the tests, move the MeshToGrid UDF sidecar to single-space storage, replace leftover GridHandle::size() calls with bufferSize(), read the dual family's device bytes with a checked copy, include the headers the multi-GPU tests use, and update comments that still described the dual-space flow. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>

The lesson, cheat sheet and capstone taught the dual-space cuda::DeviceBuffer with copy<DeviceBuffer>() and deviceUpload(), which are deprecated and removed in 14.0. Teach cuda::Buffer<std::byte> with cuda::copyTo in both directions instead, pass a cuda::Buffer<std::byte> prototype to the GPU builders and topology operators (noting that 14.0 makes it the default), build the device NodeManager through a memory resource, and allocate scratch arrays as typed cuda::Buffers. The mgr/deviceMgr table now covers host, device-only and managed storage. Also replace the deprecated math::ZeroCrossing with zeroCrossing. Every compiled snippet now builds without deprecation warnings. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>

swahtz

Thanks for this. The direction makes sense, and the migrated call sites read much more cleanly. My comments fall into two groups. The first is the deprecation messages and the lesson doc. In a couple of places, users either get pointed at the wrong replacement or get no warning at all before 14.0. The second is test coverage. Several tests were the only coverage for paths that stay the default in 13.x, and they now use the new storage. Those defaults still ship, so I'd like to keep at least one test on each of them until they are removed.

swahtz · 2026-10-21T01:33:22Z

    /// @brief Download the grid to from the device, e.g. from GPU to CPU
    /// @note This method is only available if the buffer supports devices
+    /// @deprecated cuda::copyTo a host handle to read the grids on the host.
    template<typename U = BufferT>


These messages send every dual handle to cuda::copyTo and cuda::Buffer<std::byte>. That doesn't work for GridHandle<DualUnifiedBuffer>, which is what DistributedPointsToGrid::getHandle returns by default. A device-only cuda::Buffer lives on one GPU, but the other GPUs in the mesh write into the grid. The replacement that works is cuda::Buffer<std::byte, cuda::ManagedResource>, which is what the TestMultiGPU migration in this PR uses. Right now only the deviceData(int) message mentions it.

Could the deviceUpload/deviceDownload messages mention the managed buffer for multi-GPU handles? Alternatively, the DualUnifiedBuffer case could get its own message.

Fixed in 89f321b: the deviceUpload/deviceDownload messages (and their @deprecated docs) now add "for a multi-GPU handle, use cuda::Buffer<std::byte, cuda::ManagedResource>". The DualUnifiedBuffer transfers, now deprecated too (see the next thread), point straight at the managed buffer.

swahtz · 2026-10-21T01:33:24Z

+        // instead of the deprecated GridHandle members.
        .def(
-            "deviceUpload", [](GridHandle<BufferT>& handle, bool sync) { handle.deviceUpload(nullptr, sync); }, "sync"_a = true)
+            "deviceUpload", [](GridHandle<BufferT>& handle, bool sync) { handle.buffer().deviceUpload(cudaStream_t(0), sync); }, "sync"_a = true)


This binding shows a gap in the deprecation. GridHandle::deviceDownload is deprecated, but DualDeviceBuffer::deviceDownload/deviceUpload are not. So user code like

auto h = cuda::voxelsToGrid<float>(d_coords, n); h.buffer().deviceDownload(); auto* grid = h.grid<float>();

compiles in 13.x with no warning. It then stops compiling in 14.0, when the default becomes cuda::Buffer<std::byte>, which has no deviceDownload. That is the break this PR is meant to warn about.

I'd suggest deprecating the transfer members on the dual buffers too. This binding could then suppress the warning locally for as long as the Python API keeps deviceUpload/deviceDownload.

Agreed, fixed in 89f321b. DualDeviceBuffer and DualUnifiedBuffer now deprecate their own deviceUpload/deviceDownload. The bodies moved into private members so the classes' own overloads don't warn. The deprecated GridHandle/NodeManagerHandle members forward through a new NANOVDB_SUPPRESS_DEPRECATION_BEGIN/END pair (in util/Util.h), so a caller gets exactly one warning at their own line and none from inside the headers. A probe confirms one warning each for h.deviceDownload(), h.buffer().deviceDownload(), h.buffer().deviceUpload(...) and both on a DistributedPointsToGrid handle, under -I and -isystem. This binding uses the same macros locally, as you suggested. TestBuffer.DeprecatedDualTransfersWarnOnlyAtTheCall lives in a file without file-wide suppression, so a warning that escapes from a header fails the -Werror build; I checked that it does when the macros are removed.

swahtz · 2026-10-21T01:33:25Z

+|----------------------------------------------------|-----------------------------------|----------------|------------------|
+| `HostBuffer`                                       | host only                         | host pointer   | not available    |
+| `cuda::Buffer<std::byte>`                          | device only                       | not available  | device pointer   |
+| `cuda::Buffer<std::byte, cuda::ManagedResource>`   | host **and** device (managed)     | host pointer   | device pointer   |


The managed row says mgr<T>() returns a usable host pointer. That holds only if the grid is managed too. The node manager stores a pointer to the grid it was built from. So if the grid comes from copyTo<cuda::Buffer<std::byte>>, as this lesson teaches, and the manager is then built with a ManagedResource, host traversal through mgr<float>() dereferences device memory and segfaults. Could the table or the paragraph below it say that the grid must be managed as well?

Fixed in 89f321b: the paragraph under the table now says the node manager stores a pointer to the grid it was built from, so mgr<T>() is usable on the host only if that grid is managed too. A manager built over a device-only grid points into device memory, and traversing it on the host faults.

swahtz · 2026-10-21T01:33:27Z

        //gpuTimer.start("GPU NodeManager");
-        auto handle2 = nanovdb::cuda::createNodeManager(d_grid);
+        nanovdb::cuda::DeviceResource resource;
+        auto handle2 = nanovdb::cuda::createNodeManager(d_grid, resource);


This was the only test of the dual cuda::createNodeManager(const NanoGrid*, const BufferT& = DualDeviceBuffer(), stream) overload. That overload is still the default in 13.x and isn't deprecated. After this change, every caller in unittest/ and examples/ uses the resource form. Could this test keep a dual-overload case next to the resource one, so the default path stays covered until it's removed?

Fixed in 89f321b: NodeManager keeps a dual-overload block. It builds a DualDeviceBuffer sphere, uploads it, calls cuda::createNodeManager(d_grid) with no resource, pins the result type to NodeManagerHandle<DualDeviceBuffer> with a static_assert, and checks the active-voxel count against the host manager.

swahtz · 2026-10-21T01:33:29Z

    nanovdb::cuda::DeviceMesh deviceMesh;
    nanovdb::tools::cuda::DistributedPointsToGrid<BuildT> converter(deviceMesh);
-    auto handle = converter.getHandle(voxels, voxelCount);
+    auto handle = converter.getHandle(voxels, voxelCount, BufferT());


FewTiles and ManyTiles were the last tests calling getHandle with its default DualUnifiedBuffer. Now every DistributedPointsToGrid test passes a managed cuda::Buffer, so nothing covers what users get from a plain getHandle(points, n). That includes the DualUnifiedBuffer storage path (create(size, pool) with cudaCpuDeviceId and host-preferred pages). Could one of these two keep the default buffer?

Fixed in 89f321b: FewTiles_DistributedCudaPointsToGrid is back on plain getHandle(points, n), with a static_assert that the handle is GridHandle<DualUnifiedBuffer>. It reads the grid on the host in place (unified memory), so it needs no deprecated transfer. ManyTiles keeps the managed cuda::Buffer.

swahtz · 2026-10-21T01:33:31Z

-    EXPECT_EQ(2u, idxHdl.buffer().bufferCount());
-    EXPECT_TRUE(idxHdl.buffer().data(0, cudaCpuDeviceId));
-    EXPECT_TRUE(idxHdl.buffer().data(0, 0));
+    auto devIdxHdl = nanovdb::cuda::copyTo<SingleBufT>(idxHdl, cudaStream_t(0));


The migration dropped the DualDeviceBuffer bookkeeping checks from this test. Those checks confirmed that bufferCount() goes from 1 to 2 after upload, and that data(0, cudaCpuDeviceId) and data(0, 0) are both present. No other test checks this now (the remaining bufferCount checks near line 217/226 cover only the empty and host-only states). Since DualDeviceBuffer is still the default, I'd keep these checks somewhere, either here or in a dual-buffer test in TestBuffer.cu.

Fixed in 89f321b: the new TestNanoVDBCUDA.DualDeviceBufferUploadBookkeeping checks that bufferCount() goes from 1 to 2 on upload, that data(0, cudaCpuDeviceId) and data(0, 0) are both present afterwards, and that deviceGrid appears with the right type only.

…aults tested (CUDA) handle.buffer().deviceDownload() reached the same dual-space transfer as the deprecated GridHandle::deviceDownload without a warning, and breaks the same way in 14.0. Deprecate deviceUpload and deviceDownload on cuda::DualDeviceBuffer and cuda::DualUnifiedBuffer too, with their bodies moved into private members so the classes' own overloads stay quiet. NANOVDB_SUPPRESS_DEPRECATION_BEGIN/END bracket the deprecated GridHandle and NodeManagerHandle members' forwarding calls, so a caller gets one warning at its own line and none from inside the headers; a test without file-wide suppression, built with warnings as errors, fails if one escapes. The current-device overloads of DualDeviceBuffer now check the device lookup. The upload/download messages name cuda::Buffer<std::byte, cuda::ManagedResource> for multi-GPU handles, which a device-only cuda::Buffer cannot replace. Keep covering the paths that stay the default in 13.x: the dual cuda::createNodeManager overload, the dual-space buffer's upload bookkeeping, and DistributedPointsToGrid::getHandle's default DualUnifiedBuffer. The lesson notes that a node manager's host pointer is usable only if its grid is managed too. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>

…CCL does nvcc with an MSVC host compiler reported every NANOVDB_SUPPRESS_DEPRECATION_END as a diagnostic pop without a matching push, failing the windows-nanovdb Python build. Under an MSVC host, suppress the warnings with __pragma("nv_diag_suppress N") and restore them with nv_diag_default, the form CCCL uses there. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Mark Harris <mharris@nvidia.com>

swahtz

Great, thanks for addressing my notes

Part of #785. ## Summary Advances the NanoVDB pin from `0ab0a81e` to `71ae97d4`, the head of openvdb's `develop/nanovdb` branch. That branch carries the batch-capable `TopologyBuilder` (AcademySoftwareFoundation/openvdb#2336), which fvdb's batched topology construction will move onto under #785. Going forward we pin to a specific commit on `develop/nanovdb` rather than `master`. `develop/nanovdb` is two commits ahead of openvdb `master`: the `TopologyBuilder` change and a `cuda::SignedFloodFill` fix (AcademySoftwareFoundation/openvdb#2358). From the old pin the bump crosses 23 NanoVDB commits. No fvdb behavior changes. ## Changes fvdb builds with `-Werror`, so three upstream changes need small adaptations: | upstream change | fvdb change | |---|---| | AcademySoftwareFoundation/openvdb#2373 deprecates memory resources without `operator==` | Add `operator==`/`operator!=` to `TorchResource` and `TorchDeviceResource` | | AcademySoftwareFoundation/openvdb#2378 deprecates `cuda::detail::HandleFactory` | Call `GridHandle::adopt` and `GridHandle::metaData` directly in `GridStorage.cu` and `DeviceGridHandleUtils.cuh` | | AcademySoftwareFoundation/openvdb#2207 makes the `Vec3` constructor `constexpr` | Remove an unused `Vec3d half` in `VoxelCoordTransform.h`, which nvcc now reports as unused (177-D) | Equality follows NanoVDB's definition: two resources are equal when memory allocated through one can be freed through the other. - `TorchResource` is stateless, so any two instances compare equal. - `TorchDeviceResource` instances compare equal when they share a device type. Torch's CUDA allocator frees a block on whichever device holds it, and PrivateUse1 memory frees only through the PrivateUse1 allocator. `HandleFactory::make` and `HandleFactory::meta` already forwarded to `adopt` and `metaData` at this pin, so that change is mechanical. ## Test plan - [x] Build with no warnings: `./build.sh wheel` (CUDA 13.4, sm_120) - [x] Full unit suite from `tests/`: `cd tests && pytest unit -q` gives 3109 passed, 85 skipped - [x] `clang-format-18 --style=file --dry-run` reports nothing on the changed files 🤖 Generated with [Claude Code](https://claude.com/claude-code) Signed-off-by: Jonathan Swartz <jonathan@jswartz.info> Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>

harrism and others added 3 commits October 21, 2026 00:20

harrism requested review from Idclip, apradhana, danrbailey, jmlait, kmuseth and richhones as code owners October 21, 2026 00:32

swahtz reviewed Oct 21, 2026

View reviewed changes

harrism and others added 2 commits October 21, 2026 02:10

swahtz approved these changes Oct 21, 2026

View reviewed changes

swahtz merged commit a2eb28d into AcademySoftwareFoundation:master Oct 21, 2026
16 checks passed

harrism mentioned this pull request Oct 21, 2026

NanoVDB Python: back DeviceGridHandle with single-space buffers (CUDA) #2384

Merged

swahtz mentioned this pull request Oct 22, 2026

Advance NanoVDB pin to develop/nanovdb (openvdb #2336) openvdb/fvdb-core#834

Merged

3 tasks done

Conversation

harrism commented Oct 21, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

swahtz left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

swahtz left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

harrism commented Oct 21, 2026 •

edited

Loading