NanoVDB: two multi-GPU DistributedPointsToGrid failures (RadixSort merge sync; DenseLeaf illegal access)

[harrism]

Two pre-existing multi-GPU failures in nanovdb_test_mgpu (nanovdb/nanovdb/unittest/TestMultiGPU.cu), independent of the memory-resource work. Both reproduce on a 2-GPU host and are NanoVDB-side bugs in tools/cuda/DistributedPointsToGrid.cuh, not a CCCL/CUB regression — CUB's DeviceMerge::MergePairs and the CUB_LAUNCH two-phase pattern were verified correct in isolation under the vendored CCCL 3.5.

Environment

• 2× NVIDIA RTX 6000 Ada, driver 595.71.05, CUDA 12.6 (nvcc), CCCL/CUB 3.5.0 (vendored).
• Note: NanoVDB CUDA/mgpu tests are compiled but never executed on a GPU in CI (runners are GPU-less; every CUDA job runs ctest -E ".*cuda.*|.*mgpu.*"), so these only manifest on dev machines.

Bug 1 — TestNanoVDBMultiGPU.RadixSort (deterministic)

Sortedness checks (EXPECT_LE/EXPECT_GE) fail — the output is not globally sorted. Single-GPU (CUDA_VISIBLE_DEVICES=0) passes, isolating the fault to the 2-GPU merge. The right-device merge under-produces: the right half of the output shows inversions and many never-written slots.

Root cause: a cross-device event-ordering hazard in radixSortAsync (~L217–267). postEvents[deviceId] is reused to signal both "per-device sort done" and "mergePathKernel done", so the right-device DeviceMerge::MergePairs can launch before the peer device's sorted output is actually visible.

Bug 2 — TestNanoVDBMultiGPU.DenseLeaf_DistributedCudaPointsToGrid_UnifiedBuffer (flaky / order-dependent)

Illegal memory access (CUDA error 700). Passes in-suite, fails in isolation — order/state dependent. compute-sanitizer --tool memcheck pinpoints the real fault:

Invalid __global__ read of size 4 bytes
  at nanovdb::tools::cuda::VoxelKeyFunctor<...>::operator()(...)+0xe0 in tools/cuda/PointsToGrid.cuh:562
  Address 0x... is out of bounds, ~51 GB after the nearest allocation (512 bytes)
  Device Frame: offsetLambdaKernel<VoxelKeyFunctor<...>>(...) in util/cuda/Util.h:260

PointsToGrid.cuh:562 is Vec3T p = points[d_indx[tid]];. The index d_indx[tid] (mIndices[tid]) is uninitialized/garbage in this configuration, so points[garbage] reads far out of bounds. The kernel is launched from the distributed serial per-tile loop at DistributedPointsToGrid.cuh:747. The error is asynchronous, so it only surfaces at the next cudaEventSynchronize (~L705) — which is why it superficially looks like a copy/sync bug there. Root cause (why mIndices is uninitialized over the per-tile index range in the dense-single-tile path) is under investigation.

Repro

cmake --build build/nanovdb --target nanovdb_test_mgpu
cd build/nanovdb/unittest
./nanovdb_test_mgpu --gtest_filter='*RadixSort*'     # sortedness failures
./nanovdb_test_mgpu --gtest_filter='*DenseLeaf*'     # CUDA 700 (async)
/usr/local/cuda/compute-sanitizer/compute-sanitizer --tool memcheck \
    ./nanovdb_test_mgpu --gtest_filter='*DenseLeaf*'  # pinpoints PointsToGrid.cuh:562

Both fixes belong in NanoVDB (not upstream CCCL).

[harrism]

Root cause found — these are not NanoVDB bugs. Closing as environmental.

Correcting the earlier analysis in this issue: both failures are caused by silently-corrupting PCIe peer-to-peer transfers on the test system, not by anything in DistributedPointsToGrid. The two GPUs sit on different PCIe root complexes with the IOMMU enabled — a configuration where GPU↔GPU P2P is effectively unsupported and, on this driver, returns bad data instead of raising an error.

Evidence, reduced to plain CUDA (no managed memory, no NanoVDB):

• A kernel on GPU B reading a cudaMalloc buffer on GPU A through peer access returns all-ones (0xFFFFFFFF) for every element; cudaMemcpyPeer returns all-zeros — in both directions.
• cudaDeviceCanAccessPeer reports 1, and no CUDA error is raised — the corruption is entirely silent.
• The managed-memory (cudaMallocManaged) cross-GPU path fails the same way.

So NanoVDB's multi-GPU builder was a victim of corrupted peer transfers — garbage sort indices produced the DenseLeaf illegal-address crash, and stale data produced the RadixSort unsorted output. It is not the source. The "merge event-ordering hazard" and "empty-partition illegal access" hypotheses earlier in this thread are withdrawn; the merge choreography is correct CUDA.

This matches a public report of the same symptom on a different GPU generation:
https://forums.developer.nvidia.com/t/unified-memory-not-working-between-rtx6000pro-gpus/373966

Mitigation on affected bare-metal Linux systems: boot with iommu=pt (or disable the IOMMU / disable PCIe ACS). With that, GPU P2P — and the multi-GPU nanovdb_test_mgpu tests — behave correctly.

Environment: 2× RTX 6000 Ada on separate PCIe root complexes, driver 595.71.05, AMD IOMMU active. No NanoVDB code change is warranted, so closing as not planned.

[harrism]

Correcting the record on this closure now that the relevant fixes have merged.

The environmental diagnosis for the machine in this report stands: the plain-CUDA reproducers (no NanoVDB code involved) demonstrated silent PCIe peer-to-peer corruption on that box, and nothing since contradicts that. However, the implication "not a NanoVDB bug" was incomplete. Two real ordering bugs existed in DistributedPointsToGrid underneath the hardware fault, and both are now fixed:

1. processGridTreeRootEvent was recorded without a stream argument, i.e. on the legacy default stream rather than the stream carrying the root-processing work — so the other devices' waits on that event guarded nothing, and they could consume the grid tree root mid-write. Fixed in Fix race conditions and empty launches in DistributedPointsToGrid #2275.
2. leafPrefixSumKernel reads the prior device's populated range and needed an explicit cross-device event dependency in addition to implicit stream ordering. Fixed in Fix race conditions and empty launches in DistributedPointsToGrid #2275, along with guards for empty per-device ranges on small inputs.

(#2249 also fixed an adjacent latent hazard in the same code paths: cuda::DeviceBuffer freed device memory on the legacy default stream regardless of the stream it was last used on.)

Evidence these were real and machine-independent: on GPU CI hardware whose transport diagnostics run fully clean and which reports cudaDeviceCanAccessPeer = 0 for all pairs (so the failure mode in this issue cannot occur there), the DenseLeaf test failed intermittently with a structured signature — activeVoxelCount equal to exactly one device's partition — on identical code across consecutive runs, until the fixes above.

Why the original investigation missed this: the event-ordered reproducer written for this issue modeled the intended choreography, recording its events with explicit stream arguments — silently correcting, in transcription, the very omission the shipping code had. It tested the design rather than the code, and the design was sound.

Net: this machine's transport fault and the code bugs were both true at once. The issue remains correctly closed for the hardware component; the code component is resolved by #2275/#2249.