NanoVDB: correctness and determinism fixes#2249
Open
swahtz wants to merge 8 commits into
Open
Conversation
Five correctness/determinism fixes to the NanoVDB CUDA tools: * GridChecksum: evalChecksum copied sizeof(GridData)+sizeof(TreeData) (736 B) from a device buffer holding a single uint32 CRC into the 4-byte head/tail of a stack Checksum -- an out-of-bounds device read plus host stack corruption (compute-sanitizer visible; a hard CUDA "invalid argument" on some configs). Copy sizeof(uint32_t) at all four sites, matching validateChecksum. * Stream propagation: the NN_FACE_EDGE_VERTEX leaf-dilation kernel and SignedFloodFill's root pass ran on the default stream instead of the caller's stream, silently corrupting output for callers using a non-blocking stream. Launch and synchronize on the operation's stream. * Bounded grid-name copy: the grid-name copy read a fixed MaxNameSize bytes from a possibly shorter std::string (whose .data() is never null, making the memset branch dead), over-reading up to 255 bytes of host heap into the grid and any written .nvdb file. Zero the field, then copy only min(name.size()+1, MaxNameSize) bytes. Applied to PointsToGrid, DistributedPointsToGrid, MeshToGrid. * Deterministic output initialization: PointsToGrid and IndexToGrid left alignment padding and stats fields carrying recycled pool bytes, making output nondeterministic and leaking host heap into files. Zero the whole grid buffer once after allocation. For PointsToGrid this makes the three dense background-fill passes (upper/lower value tables, and the inactive-leaf-values pass for all but the Point build type) redundant, so they are dropped. * DeviceBuffer stream-ordered lifetime: the destructor, move-assignment and clear() freed device allocations on the default stream regardless of the stream the buffer was last used on -- a latent use-after-free for non-blocking-stream callers, masked today by legacy-stream semantics. Track the allocation stream and free on it, matching the stream-ordered free TempPool now performs. Adds gtest coverage in unittest/TestNanoVDB.cu: cross-stream dilation determinism, grid-name round-trip (empty/short/max-length), PointsToGrid output determinism, and DeviceBuffer non-blocking-stream lifetime. No performance claims. Signed-off-by: Jonathan Swartz <jonathan@jswartz.info> Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
There was a problem hiding this comment.
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.
…eeds
The Tranche-1 determinism fix zero-initialized the *entire* IndexToGrid output
buffer so that bytes the build kernels never write - stats fields absent from
the source grid, and struct alignment padding - are deterministic instead of
carrying recycled pool bytes (which otherwise leak into written .nvdb files).
Measured against master this cost 10-22% on indextogrid across dragon/emu/
crawler/wdas_cloud: the leaf array dominates the buffer and was memset to zero
and then immediately overwritten value-by-value - a redundant multi-GB pass.
Scope the initialization to the bytes that are genuinely left unwritten:
* getBuffer() memsets only the non-leaf region [0, node[0]) - grid, tree,
root, root tiles and the internal nodes - a small fraction of the buffer.
* processLeafsKernel zeroes each leaf's own gap [&mMinimum, mValues): the
stats fields (skipped when the source has no stats) and any padding before
the 32-aligned mValues array. Every mValues[i] is written anyway, so this
reproduces the former full zero-init byte-for-byte.
Output is bit-identical to the full zero-init - all indextogrid goldens pass on
dragon/emu/crawler/wdas_cloud, compute-sanitizer memcheck is clean and initcheck
shows no new reports - and the regression vs master is gone (NEUTRAL on all four
datasets, RTX PRO 6000 / sm_120).
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
…ce in test Addresses PR review: - The grid-name copy could leave GridData::mGridName without a null terminator when the name length >= MaxNameSize: min(size+1, MaxNameSize) == MaxNameSize overwrote the entire zeroed field, and gridName() (a C-string accessor) would then read past it. Copy at most MaxNameSize-1 bytes and rely on the preceding memset for termination, truncating over-long names instead. Fixed in PointsToGrid, MeshToGrid and DistributedPointsToGrid. - GridName_CudaPointsToGrid gains an over-long (>= MaxNameSize) case asserting the result is truncated to MaxNameSize-1 chars and stays null-terminated. - NonBlockingStreamDilate_ValueOnIndex queried device 0 for the clock rate; query the current device via cudaGetDevice() instead. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
…ream The Tranche-1 stream-correctness fix added an unconditional cudaStreamSynchronize(stream) before processRoot's synchronous default-stream cudaMemcpy of the tree. That sync is only needed when the caller's node passes ran on a non-default stream; on the default stream (0) the blocking copy already orders after prior stream-0 work, so the extra sync there was pure overhead - a measured ~4-11% regression on floodfill/dragon vs master. Guard it with `if (stream != 0)`. Correctness for non-blocking-stream callers is unchanged. floodfill is now NEUTRAL vs master on dragon/emu/crawler (0 regressions), and byte-exact validation is unaffected. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
swahtz
changed the title
Addresses PR review: the Tranche-1 stream-ordered-free change stored a single mStream, overwritten by every DeviceBuffer::deviceUpload. A DeviceBuffer can hold an allocation on each device (mGpuData is per-device), and the multi-GPU example (ex_make_mgpu_nanovdb) uploads one handle to every device with that device's own stream. The destructor/move-assign then freed *every* device's allocation on the last device's stream - a cross-device cudaFreeAsync(ptr@devA, stream@devB), which is invalid (the free stream must belong to the allocation's device). Pre-Tranche-1 code freed on stream 0 uniformly, so this was a regression for the multi-GPU path. Track the allocation stream per device in a cudaStream_t array parallel to mGpuData, set in init()/deviceUpload() at mStreams[device], and free each mGpuData[i] on its own mStreams[i] in clear()/move-assignment. The single-GPU path is unchanged (one device, one stream); the multi-GPU path now frees each device's allocation on its own device-bound stream. Single-GPU verified (DeviceBuffer lifetime gtest + byte-exact validation); the multi-GPU free path is correct by construction but not runtime-tested here (single-GPU environment). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
…-wait (test) Addresses PR review: the test launches streamBusyWaitKernel on the default stream to occupy it during the candidate dilation, but only synchronized the non-blocking stream afterward. The busy-wait could remain queued on stream 0 and leak into subsequent tests, causing timing-dependent flakes. Drain it with a cudaDeviceSynchronize() after the candidate run. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Signed-off-by: Jonathan Swartz <jonathan@jswartz.info>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
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, and no performance claims.
Fixes
1. GridChecksum — out-of-bounds device read + host stack corruption
evalChecksumcopiedsizeof(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 stackChecksum. 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 copiessizeof(uint32_t)at all four sites, matchingvalidateChecksum.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 shorterstd::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 onlymin(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.
[0, node[0])plus a per-leaf zeroing of the stats/padding gap inprocessLeafsKernel(everymValues[i]is written anyway). This is byte-identical to a full zero-init (allindextogridgoldens 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 the stream the buffer was last used on — a latent use-after-free for non-blocking-stream callers, currently masked by legacy-stream semantics. The buffer now tracks its allocation stream and frees on it, matching the stream-ordered frees TempPool already performs.Tests
New gtest coverage in
unittest/TestNanoVDB.cu:🤖 Generated with Claude Code