NanoVDB: correctness and determinism fixes

[swahtz]

Summary

Five correctness/determinism fixes to the NanoVDB CUDA tools, plus regression tests. Bug fixes only — no API changes, no behavior change for already-correct callers.

Fixes

1. GridChecksum — out-of-bounds device read + host stack corruption
evalChecksum copied sizeof(GridData)+sizeof(TreeData) (736 B) out of a device buffer that only holds a single 4-byte CRC, into the 4-byte head/tail of a stack Checksum. That's both an out-of-bounds device read and a host stack smash (compute-sanitizer flags it; some configs hard-fail with a CUDA "invalid argument"). Now copies sizeof(uint32_t) at all four sites, matching validateChecksum.

2. Stream propagation — work silently escaping to the default stream
The 26-neighbour (NN_FACE_EDGE_VERTEX) leaf-dilation kernel and SignedFloodFill's root pass launched and synchronized on the default stream instead of the caller's stream, silently corrupting output for callers on a non-blocking stream. Both now use the operation's stream.

3. Grid-name copy — up to 255 bytes of host heap over-read into the output
The name copy read a fixed MaxNameSize (256) bytes from a possibly shorter std::string — whose .data() is never null, so the old memset fallback was dead code — over-reading host heap into the grid and into any written .nvdb. Now zeroes the field and copies only min(name.size()+1, MaxNameSize). Applied to PointsToGrid, DistributedPointsToGrid, and MeshToGrid.

4. Deterministic output — recycled pool bytes leaking into grids/files
PointsToGrid and IndexToGrid left alignment padding and stats fields (absent when the source has no stats) holding whatever the recycled device pool last contained, making output nondeterministic and leaking host heap into files.

• PointsToGrid: the grid buffer is zeroed once after allocation, which also makes the three dense background-fill passes (upper/lower value tables, and inactive-leaf-values for every build type except Point) redundant, so they're dropped — net-neutral on runtime, measured against master (same harness, 30 iters / 10 warmup, RTX PRO 6000 sm_120, median):

  op	dataset	master	this PR	delta
  p2g_onindex_n23347893	dragon	10.700 ms	10.581 ms	−1.1%
  p2g_onindex_n93391572_dup4	dragon	42.949 ms	42.402 ms	−1.3%
  p2g_onindex_n373566288_dup16	dragon	169.011 ms	168.292 ms	−0.4%
  p2g_onindex_n96956688	emu	47.592 ms	46.996 ms	−1.3%

  (marginally favourable rather than exactly neutral, by about the width of run-to-run noise; the master side needs --checksum=skip to benchmark at all, since it trips the crash fixed in 1 above)

• IndexToGrid: the initialization is scoped to only the bytes actually left unwritten — a memset of the non-leaf region [0, node[0]) plus a per-leaf zeroing of the stats/padding gap in processLeafsKernel (every mValues[i] is written anyway). This is byte-identical to a full zero-init (all indextogrid goldens pass; compute-sanitizer memcheck clean) while avoiding the redundant multi-GB memset a whole-buffer zero-init would incur — which measures as a 10–22% regression vs. master on this op, eliminated here.

5. DeviceBuffer — stream-ordered free (latent use-after-free)
The destructor, move-assignment, and clear() freed device memory on the default stream regardless of what work was still outstanding on the buffer — a latent use-after-free for non-blocking-stream callers, since the default stream is only implicitly ordered after blocking streams.

An earlier revision of this PR fixed that by freeing on the stream the buffer was last used on. @matthewdcong pointed out that this is insufficient when a buffer is used on more than one stream: the free is then ordered after the last one only. That is reproducible — a device-only buffer plus a user kernel on a second blocking stream has the freed block recycled and fully overwritten by the late kernel, where freeing on the default stream passes.

The buffer now tracks its uses with a per-device event: each use waits on the event before issuing work and re-records it afterwards, so one event transitively covers every stream the buffer has been used on, and the frees wait on it. This is correct for any mix and number of blocking and non-blocking streams, and it removes any requirement that callers keep their streams alive until the buffer is destroyed, since the event is owned by the buffer.

Worth noting for reviewers: clear() calls cudaFreeHost(mCpuData) before the device frees, and cudaFreeHost implicitly synchronizes. That masks this whole class of hazard for buffers that own pinned host memory — it is device-only buffers (create(size, nullptr, device, stream), e.g. in buildVoxelBlockManager) where mCpuData is null that are actually exposed. The fix here makes the ordering explicit rather than leaving correctness resting on that implicit synchronization, which is not in CUDA's documented list.

Tests

New gtest coverage in unittest/TestNanoVDB.cu:

• cross-stream dilation determinism (default stream vs. a non-blocking stream while the default stream is occupied)
• grid-name round-trip (empty / short / maximum-length)
• PointsToGrid byte-for-byte output determinism
• DeviceBuffer allocate/use/free lifetime on a non-blocking stream

[Copilot]

Pull request overview

This PR fixes several CUDA correctness and determinism issues in NanoVDB tooling (stream propagation, checksum safety, deterministic initialization, and safer grid-name copying) and adds targeted CUDA/gtest regressions to prevent recurrence.

Changes:

• Fix CUDA stream propagation so kernels and synchronizations honor the caller-provided stream (avoid default-stream leakage).
• Fix checksum and buffer-lifetime hazards (correct-sized checksum copies; stream-ordered frees for DeviceBuffer).
• Improve determinism and hygiene (zero-initialize grid buffers; safer grid-name copies) and add regression tests covering these scenarios.

Reviewed changes

Copilot reviewed 9 out of 9 changed files in this pull request and generated 4 comments.

Show a summary per file

File	Description
nanovdb/nanovdb/unittest/TestNanoVDB.cu	Adds CUDA regression tests for cross-stream correctness, grid-name round-trip safety, determinism, and DeviceBuffer lifetime behavior.
nanovdb/nanovdb/tools/cuda/SignedFloodFill.cuh	Ensures root processing honors the operation stream and synchronizes appropriately for host reads.
nanovdb/nanovdb/tools/cuda/PointsToGrid.cuh	Zero-initializes the output buffer for determinism; fixes grid-name copy to avoid host over-read; removes redundant background fill passes.
nanovdb/nanovdb/tools/cuda/MeshToGrid.cuh	Updates grid-name copy to avoid over-read by copying only the actual string (after clearing destination).
nanovdb/nanovdb/tools/cuda/IndexToGrid.cuh	Zero-initializes the output buffer to prevent nondeterminism/leakage from recycled allocator bytes.
nanovdb/nanovdb/tools/cuda/GridChecksum.cuh	Fixes incorrect copy size when reading back CRC values (avoid OOB device read / host corruption).
nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh	Updates distributed grid-name copy to avoid fixed-size over-reads and propagate deterministic name contents.
nanovdb/nanovdb/tools/cuda/DilateGrid.cuh	Fixes NN_FACE_EDGE_VERTEX dilation kernel launch to use mStream rather than the default stream.
nanovdb/nanovdb/cuda/DeviceBuffer.h	Tracks an associated stream and orders frees on it to avoid latent UAF for non-blocking-stream callers.

---

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

[Copilot]

Pull request overview

Copilot reviewed 9 out of 9 changed files in this pull request and generated 1 comment.

[Copilot]

Pull request overview

Copilot reviewed 9 out of 9 changed files in this pull request and generated no new comments.

[kmuseth]

it looks good but I would like Matthew Cong to take a look at it as well

[matthewdcong]

Would this have the potential to introduce bugs in existing code with blocking streams? For example, if you download the same buffer to multiple blocking streams, the clear operation would now only synchronize on the last blocking stream as opposed to all of the blocking streams (via the default stream) leading to early destruction.

[swahtz]

Good catch — you're right, and it reproduces. I wrote a repro and ran it against three states of the code:

scenario	master	this PR (as reviewed)	with fix
A) device-only buffer, user kernel on a second blocking stream	pass	FAIL — 64 MB clobbered	pass
B) host buffer, upload on a non-blocking stream	pass	pass	pass

Scenario A is exactly the mechanism you described: the free is ordered on the last-used stream only, the block is recycled, and the other stream's still-queued work lands in the next allocation. Two things were needed to observe it at all — a warm-up pass (on a cold context the first launches serialize and hide the race) and a spin loop the compiler can't elide.

A second finding while chasing this, which I think reframes the change. clear() calls cudaFreeHost(mCpuData) before the device frees, and cudaFreeHost implicitly synchronizes — I verified directly that a busy stream goes to done across it, while cudaFreeAsync does not. So for any buffer that owns pinned host memory that call already masks the hazard, including the non-blocking-stream case this hunk was written to fix, which is why scenario B passes on master too. The regression only escapes because device-only buffers (create(size, nullptr, device, stream), e.g. in buildVoxelBlockManager) leave mCpuData null, so cudaFreeHost(nullptr) is a no-op.

So as written, the hunk fixed a hazard that is currently masked and introduced one that isn't. Rather than revert it, I've made the ordering explicit instead of leaving correctness resting on cudaFreeHost's implicit synchronization, which isn't in CUDA's documented implicit-synchronization list.

The fix replaces the per-device stream with a per-device event: each use waits on the event before issuing work and re-records it afterwards, so a single event transitively covers every stream the buffer has been used on, and the frees wait on it. That is correct for any mix and number of blocking/non-blocking streams, and it also drops the "caller must keep these streams alive until the buffer is destroyed" requirement the old comment conceded, since the event is owned by the buffer.

Both scenarios now pass, and the VBM/topology/maintenance goldens are unchanged (46/46 on dragon and emu). Pushed, along with a corrected pendingchanges entry — the old wording described the superseded "free on the last-used stream" mechanism.

Happy to add the repro as a regression test if you'd like it, though I left it out deliberately: it needs a ~350 ms spin, and if the allocator declines to recycle the block it would pass vacuously, so it risks being a slow test that can silently stop testing anything.

[swahtz]

Follow-up: I went ahead and added the regression test after all — DeviceBufferMultiStreamFreeOrdering in TestNanoVDB.cu.

The concern about it passing vacuously turned out to be fixable. My first attempt gated the assertion on having staged the race (work still pending and the block recycled), which was backwards: with a correctly ordered free the allocator cannot hand the block out again until the other stream drains, so both of those observations are the fix working rather than a reason to skip. They are now reported in the failure message instead of gating the check.

Verified in both directions — it fails against the pre-fix header (64 MB clobbered) and passes with the fix. It uses a device-only buffer deliberately, so clear()'s cudaFreeHost cannot mask the problem, and it needs the warm-up pass because on a cold context the first launches serialize and the race never forms.

Also relabelled the pre-existing DeviceBufferNonBlockingStreamLifetime as a smoke test: it only asserts cudaSuccess, so it passes with or without a correctly ordered free, and its comment still described the superseded "free on the stream it was allocated on" mechanism.

[matthewdcong]

Looks good to me for the most part, minus the change to the DeviceBuffer which I raised a question about above.

[harrism]

All five are real bugs, and the fixes are the right shape. Three brief inline notes, none blocking.

On the buffer stream fixes specifically:

• The freeAsync(ptr, 0) free is a genuine use-after-free for non-blocking-stream callers — the legacy default stream does not order against them — so this is a correctness fix, not defensive extra synchronization.
• Tracking the allocation/last-use stream and freeing on it is the same retained-stream contract cuda::Buffer (#2251) uses and the roadmap (#2232) records. The legacy buffer converging on it independently is good validation of the rule, and the per-device stream array is the right extension for a buffer that can hold allocations on several devices.
• I swept tools/cuda/ on this head for kernel launches without an explicit stream argument: none remain — the two fixed here (DilateGrid 26-neighbour pass, SignedFloodFill root pass) were the last.

The IndexToGrid scoped zero-init — measured against the naive whole-buffer version and engineered to be byte-identical without the 10–22% cost — is exactly how a change like that should be justified.

Sequencing note: this conflicts textually with #2269/#2270 in MeshToGrid.cuh and PointsToGrid.cuh. This PR should land first; I'll rebase that stack over it.

[Copilot]

Pull request overview

Copilot reviewed 10 out of 10 changed files in this pull request and generated 1 comment.

Suppressed comments (3)

nanovdb/nanovdb/cuda/DeviceBuffer.h:80

• This reuses one stream for every device allocation. A CUDA stream belongs to one device, but DeviceBuffer can own allocations on several devices (for example, ex_make_mgpu_nanovdb/main.cu:197-200 uploads one handle to each device). Waiting and calling cudaFreeAsync for another device's pointer on this stream can fail or free with the wrong ordering. Select a stream on each allocation's device (and switch the current device as needed) rather than applying the caller/current device's stream to the entire array.

        for (int i = 0; i < mDeviceCount; ++i) {
            this->orderAfterPriorUses(i, stream);
            cudaCheck(util::cuda::freeAsync(mGpuData[i], stream));

nanovdb/nanovdb/tools/cuda/SignedFloodFill.cuh:217

• The SignedFloodFill stream-propagation fix has no non-blocking-stream regression coverage. The existing CUDA SignedFloodFill test invokes only the default stream (TestNanoVDB.cu:1219), so it would not detect this code reverting to a default-stream launch. Add a test analogous to the dilation regression that occupies the default stream and compares results produced on a non-blocking stream.

    kernels::processRoot<BuildT>(d_tree, mStream);

nanovdb/nanovdb/unittest/TestNanoVDB.cu:3771

• Because user is a blocking stream, a free enqueued on the legacy default stream is implicitly ordered after this stream's pending kernel. Consequently, the original default-stream implementation can pass this supposed regression test—the adjacent comment even identifies this synchronization. Make the stream carrying the late write non-blocking so the test actually removes default-stream ordering and exercises the lifetime fix.

    cudaCheck(cudaStreamCreate(&user));// blocking streams, i.e. the ones the legacy
    cudaCheck(cudaStreamCreate(&other));// default stream implicitly synchronizes with

[swahtz]

All four points from this review (the comment above plus the three suppressed ones) evaluated and addressed in 457b671. For the humans following along:

Raw-pointer uses are untracked (this comment) — correct, and the most important of the four. The event only covers deviceUpload/deviceDownload; kernels launched against deviceData()'s raw pointer are invisible to it. Two clarifications on scope: work on blocking streams is still covered (the free is issued on the default stream, which is implicitly ordered after them — the same coverage master provided), and the uncovered case — raw-pointer work on a non-blocking stream — was equally uncovered on master, where no API could even express the dependency. Fixed as suggested by making the contract explicit: recordUse(device, stream) is now public, and deviceData() documents that raw-pointer work on non-blocking streams must either be registered through it or synchronized before the buffer is cleared/destroyed.

Test's user stream is blocking (suppressed, TestNanoVDB.cu) — deliberate, but the observation exposed a gap. The blocking scenario is what discriminates the last-used-stream revision this thread started with; simply flipping it to non-blocking would create a test no implementation can pass, because the late write never goes through the buffer's API. With recordUse public, the constructive version exists, so the test is now two scenarios, verified per revision:

revision	blocking + unregistered write	non-blocking + recordUse
this PR (event fix)	pass	pass
earlier revision (free on last-used stream)	FAIL	FAIL
master (free on stream 0, no event)	pass	FAIL

One stream for every device's free (suppressed, DeviceBuffer.h) — agreed in principle. A stream belongs to one device, and nothing in the SOMA documentation blesses freeing one device's pool allocation on another device's stream ("deallocation can be performed in any stream" is stated in a same-device context). freeDeviceBuffers now frees each allocation on its own device — the caller's stream for the current device, the owning device's default stream otherwise, switching devices as needed. Honesty note: master's loop had the same single-stream shape, and this machine has one GPU, so the multi-device path is verified by inspection and compilation only.

SignedFloodFill test coverage (suppressed, SignedFloodFill.cuh) — agreed, added as NonBlockingStreamSignedFloodFill, mirroring the dilation regression (occupied default stream, non-blocking run, readback on the producing stream, identical output required). One scope caveat, verified rather than assumed: it locks the cross-stream contract and would catch node passes escaping to the default stream, but it cannot discriminate processRoot's internal ordering on constructible inputs — I swapped the pre-fix header in and it passes (3/3), because that path only does work when interior root-level tiles exist (a level set thousands of voxels across) and the pre-fix code was accidentally host-synchronous otherwise via its blocking legacy-stream memcpy. The test comment says so.

Goldens unchanged (46/46), full test file compiles, self-move check still passes.