From 329fd03a829e31425d68b92c7543cd96d3dc5677 Mon Sep 17 00:00:00 2001 From: Efty Sifakis Date: Thu, 30 Jul 2026 15:05:02 -0500 Subject: [PATCH] Demonstrate that the accessor benchmark never reaches a leaf node This is a diagnostic change, not a merge candidate. It adds an assert immediately before every getValue() in both CUDA kernels, probing the same coordinate via grid->tree().root().probeLeaf() -- deliberately going through the root rather than the accessor, so accessor cache state cannot influence the result. Both benchmarks run to completion with zero assertion failures: bench_accessor_cuda_new : exit=0, 0 assertion failures bench_accessor_cuda_old : exit=0, 0 assertion failures That is, no lookup in any of the five workloads (Sequential, LeafJump, NodeJump, Random, 27-point Stencil) touches a leaf node, under either accessor type and either accessor mode. Why: createFogVolumeSphere is constant through its interior, so the interior is stored as active *tiles*, not leaves -- 82,608 leaves x 512 = 42.3M voxels against 523.6M active, i.e. ~92% of active voxels are tiles. The leaf shell lives only at radius ~[480,520), while DOMAIN=256 confines every sample to the tile interior. Measured termination levels are 100% upper-internal (L2) for Sequential and 99.8% L2 / 0.2% lower (L1) for Random; never level 0. Consequence: the reported figures compare a cached upper-node pointer comparison against a lookup in an 8-entry root table, both L1-resident, one hop below the root. They do not measure voxel access or leaf-cache reuse. This also explains why ReadAccessor<0> is nearly flat across Sequential/LeafJump/NodeJump regardless of stride -- its leaf cache can never hit, because no leaf exists in the sampled region. Note RESULTS.md:76-79 states the opposite ("every lookup lands on a real leaf and traverses the tree for real"). Verification: each assert was checked live by inverting its predicate and confirming it fires (the Release build defines NDEBUG, which is why the local #undef is required -- without it both asserts compile away silently and the run passes vacuously). Co-Authored-By: Claude Opus 5 Signed-off-by: Efty Sifakis --- nanovdb/nanovdb/benchmark/BenchAccessorCuda.cu | 18 +++++++++++++++--- 1 file changed, 15 insertions(+), 3 deletions(-) diff --git a/nanovdb/nanovdb/benchmark/BenchAccessorCuda.cu b/nanovdb/nanovdb/benchmark/BenchAccessorCuda.cu index 84bec13b7b..1f31f972d6 100644 --- a/nanovdb/nanovdb/benchmark/BenchAccessorCuda.cu +++ b/nanovdb/nanovdb/benchmark/BenchAccessorCuda.cu @@ -17,6 +17,9 @@ #include #include +#undef NDEBUG // Release defines NDEBUG, which would compile out the assert below +#include + static constexpr int CHUNK = 32; template @@ -32,8 +35,12 @@ __global__ void benchKernel(const nanovdb::NanoGrid* grid, AccT acc(grid->tree().root()); float sum = 0.0f; const int end = min(base + CHUNK, count); - for (int i = base; i < end; ++i) + for (int i = base; i < end; ++i) { + // Probe via the root, bypassing the accessor, so cache state cannot influence it. + assert(grid->tree().root().probeLeaf(coords[i]) == nullptr && + "Unexpected: Leaf node found at probed location"); sum += acc.getValue(coords[i]); + } out[tid] = sum; } @@ -54,8 +61,13 @@ __global__ void stencilKernel(const nanovdb::NanoGrid* grid, float s = 0.0f; for (int dz = -1; dz <= 1; ++dz) for (int dy = -1; dy <= 1; ++dy) - for (int dx = -1; dx <= 1; ++dx) - s += acc.getValue(c + nanovdb::Coord(dx, dy, dz)); + for (int dx = -1; dx <= 1; ++dx) { + const nanovdb::Coord n = c + nanovdb::Coord(dx, dy, dz); + // Probe via the root, bypassing the accessor, so cache state cannot influence it. + assert(grid->tree().root().probeLeaf(n) == nullptr && + "Unexpected: Leaf node found at probed location"); + s += acc.getValue(n); + } out[i] = s; } }