diff --git a/crypto/math-cuda/kernels/rpx.cu b/crypto/math-cuda/kernels/rpx.cu index 5c6a871af..1863c65d9 100644 --- a/crypto/math-cuda/kernels/rpx.cu +++ b/crypto/math-cuda/kernels/rpx.cu @@ -944,3 +944,129 @@ extern "C" __global__ void rpx_grind_search(const uint64_t *inner_felts, } } } + +// --------------------------------------------------------------------------- +// ⛔ A DIAGNOSTIC TWIN, NEVER A PROVING PATH. +// +// `rpx_grind_search` above is the shipped kernel and this file does not change +// it. This twin is the same search with three counters, and it exists to answer +// ONE question that no timing can: when a launch is slow, does it EXECUTE MORE +// PERMUTATIONS, or the same ones more slowly? +// +// The question is not idle. The paired microbench reads a mean 20% above the +// model on the record posture while the MEDIAN seed sits on it, and the excess +// saturates in `count` rather than growing with it: converted to iterations per +// thread the excess fits `T·(1 − e^(−(N−8)/τ))` with τ ≈ 19 on three +// independent points. That is the shape of a thread that keeps scanning for a +// bounded TIME after the answer is known — a poll of `*result` served stale — +// and not the shape of a thread running to the loop bound. `executed` settles +// it: stale polls mean real extra permutations, bounded by time and therefore +// the SAME at scan 8 and scan 64; anything else means the work is unchanged and +// the cost is outside this loop. +// +// ⚠ WASTED WORK, NEVER A WRONG ANSWER. The search returns the globally smallest +// valid nonce whatever any thread does after the `atomicMin`, which is why the +// counters can be read at leisure while the answer stays pinned by +// `gpu_grind_returns_smallest_valid_nonce`. Nothing here is a soundness matter. +// +// ⛔ `#if defined(__CUDACC__)`, and that is deliberate rather than defensive: +// the host-KAT compiles this file through `cuda_host_shim.h`, which supplies +// `gridDim`, `blockDim` and `atomicMin` but NOT `__shfl_down_sync`. A KAT has +// no answer to check here anyway — this kernel produces no digest — and its +// nonce agreement with the shipped kernel is asserted on the device, seed by +// seed, by the bench that launches it. +// +// SIZING, before the first cubin, by the rule this file learnt the hard way +// (`permute` stays a called function; see its CODE SHAPE note): one more +// `permute` CALL SITE, not one more inlined copy, so this adds an entry of +// order 150-250 PTX lines against a file of 6,241 — not a duplicated +// permutation body. If the `.ptx` grows by thousands, something inlined that +// must not. +#if defined(__CUDACC__) + +// One warp's reduction of a sum and a max, so the counters cost four atomics a +// warp instead of one per thread. +// +// ⛔ A per-thread `atomicAdd` would be 131,072 serialised updates of one L2 +// line per launch — the same order as the 5-11 ms effect being measured. An +// instrument that manufactures its own signal answers a different question. +// Every thread in a warp reaches this (the loop's `break`s leave the loop, not +// the function), so the full mask is correct. +// ⛔ `unsigned long long`, NOT `uint64_t`, and that is not a style choice. The +// shuffle intrinsics are overloaded on `int`, `unsigned int`, `long long`, +// `unsigned long long`, `float` and `double`. On an LP64 host `uint64_t` is +// `unsigned long`, which is NONE of them — the call would be ambiguous or +// absent rather than wrong, so it fails at compile time on the box and not +// here, where no nvcc runs. The file already casts for exactly this reason +// where it calls `atomicMin`. +__device__ __forceinline__ void warp_reduce_counts(unsigned long long &sum, + unsigned long long &max_v, + unsigned long long &ends) { +#pragma unroll + for (int off = 16; off > 0; off >>= 1) { + sum += __shfl_down_sync(0xffffffffu, sum, off); + ends += __shfl_down_sync(0xffffffffu, ends, off); + const unsigned long long other = __shfl_down_sync(0xffffffffu, max_v, off); + if (other > max_v) max_v = other; + } +} + +// The three counters live in ONE device array so a search reads them back in a +// single copy: `counts[COUNT_EXECUTED]` and `counts[COUNT_RAN_TO_END]` are +// accumulated with `atomicAdd` across every launch of the search, +// `counts[COUNT_MAX_ITERS]` with `atomicMax`. The host mirrors these names. +__device__ constexpr int COUNT_EXECUTED = 0; +__device__ constexpr int COUNT_MAX_ITERS = 1; +__device__ constexpr int COUNT_RAN_TO_END = 2; + +extern "C" __global__ void rpx_grind_search_counted(const uint64_t *inner_felts, + uint64_t limit, + uint64_t base, + uint64_t count, + volatile unsigned long long *result, + unsigned long long *counts) { + uint64_t tid = (uint64_t)blockIdx.x * blockDim.x + threadIdx.x; + uint64_t stride = (uint64_t)gridDim.x * blockDim.x; + const uint64_t f0 = inner_felts[0], f1 = inner_felts[1], f2 = inner_felts[2], + f3 = inner_felts[3]; + // Permutations THIS thread ran. Counted after both exit tests and before + // the sponge, so it counts work done and never work declined. + uint64_t iters = 0; + // Did this thread leave by the loop bound rather than by an early exit? + // Only meaningful for a thread that did some work: a thread whose `tid` is + // past `count` never enters the loop and must not be scored as having + // scanned to the end. + uint64_t to_end = 0; + uint64_t i = tid; + for (; i < count; i += stride) { + uint64_t nonce = base + i; + if (nonce < base) break; + if (nonce >= (uint64_t)*result) break; + ++iters; + rpx::Sponge sp; + sp.init(GRIND_FELTS); + sp.absorb(f0); + sp.absorb(f1); + sp.absorb(f2); + sp.absorb(f3); + sp.absorb(goldilocks::canonical(nonce)); + uint64_t digest[rpx::DIGEST_FELTS]; + sp.finalize(digest); + if (digest[0] < limit) { + atomicMin((unsigned long long *)result, (unsigned long long)nonce); + } + } + if (i >= count && iters > 0) to_end = 1; + + unsigned long long sum = (unsigned long long)iters; + unsigned long long max_v = (unsigned long long)iters; + unsigned long long ends = (unsigned long long)to_end; + warp_reduce_counts(sum, max_v, ends); + if ((threadIdx.x & 31u) == 0u) { + atomicAdd(&counts[COUNT_EXECUTED], sum); + atomicAdd(&counts[COUNT_RAN_TO_END], ends); + atomicMax(&counts[COUNT_MAX_ITERS], max_v); + } +} + +#endif // __CUDACC__ diff --git a/crypto/math-cuda/src/columns.rs b/crypto/math-cuda/src/columns.rs index 8e368c4c6..ab25c0e37 100644 --- a/crypto/math-cuda/src/columns.rs +++ b/crypto/math-cuda/src/columns.rs @@ -66,7 +66,10 @@ impl DeviceColumns { } let total: usize = columns.iter().map(|c| c.len()).sum(); let be = backend().ok()?; - let room = be.reserve(total as u64 * 8)?; + let Some(room) = be.reserve(total as u64 * 8) else { + crate::device::note_device_fallback(); + return None; + }; let stream = be.next_stream(); // SAFETY: every element is written by the copies below. let mut buffer = unsafe { alloc_or_trim::(&stream, total) }.ok()?; diff --git a/crypto/math-cuda/src/device.rs b/crypto/math-cuda/src/device.rs index b71695a2e..9d010151a 100644 --- a/crypto/math-cuda/src/device.rs +++ b/crypto/math-cuda/src/device.rs @@ -258,6 +258,11 @@ pub struct Backend { pub rpx_merkle_tail: CudaFunction, pub rpx_permute_probe: CudaFunction, pub rpx_grind_search: CudaFunction, + /// ⛔ DIAGNOSTIC ONLY — the grind search with its executed-permutation + /// counters. Nothing on a proving path launches it; its one caller is + /// [`crate::grinding::search_counted`], which reads whether a slow launch + /// does MORE work or the same work more slowly. + pub rpx_grind_search_counted: CudaFunction, // rpx.cubin — the algebraic hash's twins of the keccak entries above. // Only the ones the WHIR path reaches are bound: the coset leaves, the two @@ -473,6 +478,7 @@ impl DeviceReservation { ) { Ok(_) => { self.bytes.fetch_add(extra, Ordering::Relaxed); + note_reserved(held + extra); return true; } Err(seen) => held = seen, @@ -551,6 +557,111 @@ pub fn reserve(bytes: u64) -> Option { backend().ok()?.reserve(bytes) } +/// Argue-surface device fallbacks: the reservation refusals in math-cuda's +/// `sumcheck`, `gkr` and `columns`, counted where each one's `reserve` returns +/// `None` and its work moves to the host. +/// +/// ⛔ WHY THIS EXISTS. Until this counter the only fallback number the campaign +/// read was `multilinear::gpu::host_fallbacks()`, which has ONE caller — the +/// COMMIT path (`multilinear/src/whir_chain.rs`) — so every `host fallbacks 0` +/// certified that no COMMITMENT fell back and said NOTHING about the per-table +/// ARGUMENT. wt16 was net-negative for exactly that blind spot: the leaf-layer +/// retention grew `be.reserved`, argue's `reserve` then refused and moved to +/// the host UNCOUNTED, and the slot-level reading looked like a clean win. +/// Read beside `host_fallbacks()`, this makes "the device did the work" +/// distinguishable from "it quietly did not" on the argue surface. +/// +/// SCOPE, stated precisely. The FIVE argue-side `reserve`→`None` sites in +/// `crypto/math-cuda/src`: `sumcheck.rs` (×3), `gkr.rs` (×1), `columns.rs` +/// (×1). This is NOT the whole device surface, and it does not claim to be: +/// `multilinear/src/gpu.rs` holds two further argue-side sites — `:177` +/// (`reserve_room`) and `:1572` (the GKR tree) — whose `None` still falls to +/// the host uncounted. Those are a documented FOLLOW-UP, out of this counter's +/// scope, because each needs its caller traced before it can honestly be +/// labelled a fallback. A THIRD site there, `:1551`, is a SPECULATIVE reserve +/// whose `None` selects a lazy path that is STILL on the device — NOT a +/// fallback, and it must never be counted. Putting a wrong site into the very +/// counter meant to end false numbers is the one thing to avoid. +static DEVICE_FALLBACKS: AtomicU64 = AtomicU64::new(0); + +/// Argue-surface device fallbacks this process has taken — see +/// [`DEVICE_FALLBACKS`] for the enumerated sites and the scope it does not +/// cover. Read alongside `multilinear::gpu::host_fallbacks()` (the commit-side +/// count) for both surfaces. +pub fn device_fallbacks() -> u64 { + DEVICE_FALLBACKS.load(Ordering::Relaxed) +} + +/// Zero the process-wide counter. For a test that wants to assert a delta, and +/// for a harness that reads one prove's worth from a reused process. +pub fn reset_device_fallbacks() { + DEVICE_FALLBACKS.store(0, Ordering::Relaxed); +} + +/// Record one argue-side reservation refusal — bumped at each of the five +/// sites [`DEVICE_FALLBACKS`] enumerates, and nowhere else. +pub(crate) fn note_device_fallback() { + DEVICE_FALLBACKS.fetch_add(1, Ordering::Relaxed); +} + +/// ★ The high-water mark of `be.reserved` — the PEAK simultaneous device +/// reservation this process reached, updated wherever the total rises +/// ([`Backend::reserve`] and [`DeviceReservation::grow`]). +/// +/// This is the reservation quantity argue's `reserve` is checked against, which +/// the raw device trace cannot report: the raw peak includes the never-purge +/// pool's retained blocks and sits above the reservation budget, while THIS is +/// exactly what the budget gates. A control run reads argue's peak reservation +/// demand here; and while the evictable retention holds only spare bytes, this +/// stays below the budget by construction. +static RESERVED_HIGH_WATER: AtomicU64 = AtomicU64::new(0); + +/// Note that `be.reserved` just rose to `now`, keeping the peak. +fn note_reserved(now: u64) { + RESERVED_HIGH_WATER.fetch_max(now, Ordering::Relaxed); +} + +/// The peak simultaneous device reservation this process reached — see +/// [`RESERVED_HIGH_WATER`]. +pub fn reserved_high_water() -> u64 { + RESERVED_HIGH_WATER.load(Ordering::Relaxed) +} + +/// Zero the reservation high-water. For a test asserting a delta. +pub fn reset_reserved_high_water() { + RESERVED_HIGH_WATER.store(0, Ordering::Relaxed); +} + +/// ★ THE RETENTION EVICTOR — the callback the WHIR leaf-layer retention installs +/// so the allocator can reclaim spare retained bytes UNDER PRESSURE, on a real +/// caller's behalf, without `device.rs` depending on `whir.rs`. +/// +/// This is what makes the retention a PRIORITY scheme, not a byte race +/// ([[gpu-two-vramgates-overlap]]): it holds only genuinely-spare bytes, and the +/// instant a real caller (argue) cannot get its reservation, the layers are +/// given back. A plain `fn` pointer — Send + Sync, no allocation — installed once +/// via `OnceLock`; it takes a byte TARGET and returns how many it actually freed. +/// +/// ⛔ CONTRACT the evictor MUST honour (see `whir::evict_retained_layers`): it +/// frees by dropping retained layers and calling `DeviceReservation::shrink`, and +/// it MUST NOT call `reserve` or `grow` — those would re-enter the allocator that +/// called it and deadlock. +static EVICTOR: OnceLock u64> = OnceLock::new(); + +/// Install the retention evictor. Idempotent — only the first install sticks. +pub fn set_retention_evictor(evictor: fn(u64) -> u64) { + let _ = EVICTOR.set(evictor); +} + +/// Ask the retention to give back at least `target` bytes of budget; returns how +/// many it freed (0 if none is installed or nothing was reclaimable). +fn try_evict_retained(target: u64) -> u64 { + match EVICTOR.get() { + Some(evict) => evict(target), + None => 0, + } +} + /// Allocates on `stream`, and if the device says no, gives the pool's retained /// blocks back and asks once more. /// @@ -810,6 +921,7 @@ impl Backend { rpx_leaves_base_coset: rpx.load_function("rpx_leaves_base_coset")?, rpx_leaves_ext3_coset: rpx.load_function("rpx_leaves_ext3_coset")?, rpx_grind_search: rpx.load_function("rpx_grind_search")?, + rpx_grind_search_counted: rpx.load_function("rpx_grind_search_counted")?, barycentric_base_batched: bary.load_function("barycentric_base_batched")?, barycentric_ext3_batched: bary.load_function("barycentric_ext3_batched")?, barycentric_base_batched_strided: bary @@ -954,9 +1066,25 @@ impl Backend { /// has a host path. The budget binds only when several proofs share a /// card: one of them is enough to fill it. pub fn reserve(&self, bytes: u64) -> Option { + let mut evicted = false; let mut held = self.reserved.load(Ordering::Relaxed); loop { if held.saturating_add(bytes) > self.vram_budget_bytes { + // ★ Before giving up, ask the retention for spare bytes ONCE. The + // retention holds only genuinely-spare leaf layers and gives them + // back here, so a real caller (argue) is never displaced by a + // cache — the fix for wt16, where the retention won the shared + // budget and argue fell to the host. Bounded to one pass by + // `evicted`, so a persistent miss returns None (and the call + // site's `note_device_fallback` counts it) rather than spinning. + if !evicted { + evicted = true; + let deficit = held.saturating_add(bytes) - self.vram_budget_bytes; + if try_evict_retained(deficit) > 0 { + held = self.reserved.load(Ordering::Relaxed); + continue; + } + } return None; } match self.reserved.compare_exchange_weak( @@ -966,6 +1094,7 @@ impl Backend { Ordering::Relaxed, ) { Ok(_) => { + note_reserved(held + bytes); return Some(DeviceReservation { bytes: AtomicU64::new(bytes), }); @@ -1384,3 +1513,82 @@ impl Backend { Ok(PooledEvent { event: Some(ev) }) } } + +#[cfg(test)] +mod device_fallback_counter_tests { + use super::{ + device_fallbacks, note_device_fallback, note_reserved, reserved_high_water, + reset_device_fallbacks, reset_reserved_high_water, + }; + use std::sync::Mutex; + + /// The counter is process-wide, so a test asserting an absolute value + /// serialises against anything else in this binary that might move it. + static COUNTER: Mutex<()> = Mutex::new(()); + + /// The plumbing, card-free: a bump reads as one, bumps accumulate, and a + /// reset reads as zero. This is the CONTROL for the box test — if + /// `note`/`device_fallbacks`/`reset` did not agree here, no site test could + /// be trusted. + #[test] + fn note_bumps_read_and_reset_zeroes() { + let _g = COUNTER.lock().unwrap_or_else(|e| e.into_inner()); + reset_device_fallbacks(); + assert_eq!(device_fallbacks(), 0, "reset must zero the counter"); + note_device_fallback(); + assert_eq!(device_fallbacks(), 1, "one note reads one"); + note_device_fallback(); + assert_eq!(device_fallbacks(), 2, "notes accumulate"); + reset_device_fallbacks(); + assert_eq!(device_fallbacks(), 0, "reset must zero it again"); + } + + /// ⛔ THE FOUR UNDRIVEN SITES' FALSIFIER. The box test drives ONE site + /// (`DeviceColumns::upload`); this arm is what lets the other four fail + /// without four card fixtures. `note_device_fallback` must be CALLED at + /// exactly the five argue-surface sites the counter's doc enumerates — + /// three in `sumcheck.rs`, one in `gkr.rs`, one in `columns.rs`. Removing + /// the call at ANY site changes the tuple and reddens this test by name, + /// which localises the loss to the file it happened in. + /// + /// The pattern carries the `crate::device::` prefix so it counts CALLS and + /// never the definition (`pub(crate) fn note_device_fallback`). + #[test] + fn note_device_fallback_is_called_at_exactly_the_five_argue_sites() { + const PATTERN: &str = "crate::device::note_device_fallback()"; + let sumcheck = include_str!("sumcheck.rs").matches(PATTERN).count(); + let gkr = include_str!("gkr.rs").matches(PATTERN).count(); + let columns = include_str!("columns.rs").matches(PATTERN).count(); + assert_eq!( + (sumcheck, gkr, columns), + (3, 1, 1), + "the argue-surface device-fallback counter must be bumped at exactly \ + the five sites the counter's doc names: sumcheck ×3, gkr ×1, \ + columns ×1 (found sumcheck {sumcheck}, gkr {gkr}, columns {columns})" + ); + } + + /// The reservation high-water is a MONOTONE peak: it takes the max, so a + /// smaller rise never lowers it and a larger one raises it. Card-free — it + /// exercises `note_reserved` directly (the same call `reserve`/`grow` make on + /// every successful rise of `be.reserved`), so a peak that failed to track + /// would show here before any card run relied on the number. + #[test] + fn reserved_high_water_keeps_the_max() { + let _g = COUNTER.lock().unwrap_or_else(|e| e.into_inner()); + reset_reserved_high_water(); + assert_eq!(reserved_high_water(), 0, "reset must zero the high-water"); + note_reserved(100); + assert_eq!(reserved_high_water(), 100, "the first rise sets the peak"); + note_reserved(50); + assert_eq!( + reserved_high_water(), + 100, + "a smaller rise must not lower it" + ); + note_reserved(200); + assert_eq!(reserved_high_water(), 200, "a larger rise raises it"); + reset_reserved_high_water(); + assert_eq!(reserved_high_water(), 0, "reset must zero it again"); + } +} diff --git a/crypto/math-cuda/src/gkr.rs b/crypto/math-cuda/src/gkr.rs index 544d4de4a..49acf99b0 100644 --- a/crypto/math-cuda/src/gkr.rs +++ b/crypto/math-cuda/src/gkr.rs @@ -130,6 +130,7 @@ impl DeviceFractionTree { // The levels above the input layer halve, so the whole tree is twice // it — and the input layer is `p` and `q` together. let Some(room) = be.reserve(p.len() as u64 * 8 * 4) else { + crate::device::note_device_fallback(); return Err(cudarc::driver::DriverError( cudarc::driver::sys::CUresult::CUDA_ERROR_OUT_OF_MEMORY, )); diff --git a/crypto/math-cuda/src/grinding.rs b/crypto/math-cuda/src/grinding.rs index 1a83a859a..31efb03f3 100644 --- a/crypto/math-cuda/src/grinding.rs +++ b/crypto/math-cuda/src/grinding.rs @@ -25,8 +25,12 @@ use cudarc::driver::{LaunchConfig, PushKernelArg}; use crate::device::backend; -const BLOCK_DIM: u32 = 256; -const GRID_DIM: u32 = 1024; +/// Threads per block for the keccak arm. +/// +/// ⛔ NOT a knob, and neither is its RPX twin. The two block dims are tuned per +/// kernel against register pressure, which is a property of the kernel body; +/// the GRID is what a launch chooses, and that is [`GRID_ENV`]. +pub const BLOCK_DIM: u32 = 256; /// Threads per block for the RPX arm. /// @@ -34,7 +38,7 @@ const GRID_DIM: u32 = 1024; /// narrow: a thread carries a twelve-lane `u64` state plus the inverse S-box's /// live temporaries across a non-inlined `permute` call, so occupancy is bought /// with registers rather than threads. -const RPX_BLOCK_DIM: u32 = 128; +pub const RPX_BLOCK_DIM: u32 = 128; /// Below this grinding factor the CPU search finds a valid nonce in well under /// a microsecond, so a device launch + shared-stream `synchronize` (which also @@ -42,6 +46,241 @@ const RPX_BLOCK_DIM: u32 = 128; /// those to the CPU. The production factor is 20; only tests use tiny factors. pub const GRIND_MIN_FACTOR: u8 = 12; +/// Smallest per-launch block, so a tiny grinding factor still fills the grid. +const MIN_BLOCK: u64 = 1 << 18; + +/// Largest per-launch block, so a huge grinding factor does not ask for an +/// absurd single launch. A miss just advances `base` and relaunches. +const MAX_BLOCK: u64 = 1 << 28; + +/// ★ `LAMBDA_VM_GRIND_SCAN_FACTOR` — how many expected hit distances one launch +/// covers. Default [`SCAN_FACTOR_DEFAULT`]. +/// +/// # ⛔ This is a CEILING on the block, not a multiplier on the work +/// +/// Both kernels carry an early exit — `if (nonce >= *result) break;` against a +/// `volatile` result the `atomicMin` writes through L2 — and the stride walk +/// `for (i = tid; i < count; i += gridDim*blockDim)` gives every nonce in +/// `[base, base+count)` exactly one owner. So once the first valid nonce `h` is +/// recorded, every thread stops within one stride round: the permutations +/// actually executed are `h + stride`. Launches before the hitting one cover +/// exactly the part of `[0, h)` below it, so the total over the whole search is +/// `h + stride` **for any scan factor**. +/// +/// What the factor buys is only the probability that one launch suffices, +/// `P = 1 − e^−k`: 99.97% at the default 8, 63% at 1. Lowering it removes no +/// permutations — they were never executed — and adds `1/(1 − e^−k)` expected +/// launches, each a sentinel H2D, a launch, an 8-byte D2H and a stream +/// synchronize. +/// +/// ⇒ The knob exists so that reading is measurable on the card rather than +/// argued from the source. Nothing here touches the grinding factor itself +/// (the security parameter), the kernels, or the host re-validation of every +/// nonce the device returns. +pub const SCAN_FACTOR_ENV: &str = "LAMBDA_VM_GRIND_SCAN_FACTOR"; + +/// The default, and the posture every recorded measurement was taken under. +pub const SCAN_FACTOR_DEFAULT: u32 = 8; + +/// What [`SCAN_FACTOR_ENV`] accepts, inclusive, and what its error names. +/// +/// 0 is refused rather than clamped: `0 * expected` is 0, which the clamp would +/// turn into a fixed [`MIN_BLOCK`] block — a different search, not a smaller +/// one, and silently so. +pub const SCAN_FACTOR_RANGE: std::ops::RangeInclusive = 1..=64; + +/// ★ `LAMBDA_VM_GRIND_GRID` — blocks per launch. Default [`GRID_DEFAULT`]. +/// +/// # Why the GRID and not the block dims +/// +/// `stride = grid × block_dim` is the number of nonces the walk advances per +/// iteration, and it is the term the early exit above leaves behind: the search +/// executes `h + stride` permutations, so the overshoot past the first hit is +/// `stride/h` — 12.5% at the default against `h = 2^20`. +/// +/// That gives the knob two opposite edges and the sweep has to run BOTH ways: +/// - while the card is **not** filled, a wider grid raises throughput faster +/// than it raises the overshoot, and the wall falls; +/// - once the card **is** filled, extra blocks only queue, and the wider stride +/// is pure added work — narrower is then strictly better. +/// +/// Which edge the default sits on is a residency question, and residency is +/// decided by registers per thread, which is why [`BLOCK_DIM`] and +/// [`RPX_BLOCK_DIM`] are NOT knobs: they are tuned per kernel body (lane K's +/// territory), and moving them changes what an occupancy reading means. +/// [`device_fill`] reads the answer off the driver instead of estimating it. +pub const GRID_ENV: &str = "LAMBDA_VM_GRIND_GRID"; + +/// The default, and the posture every recorded measurement was taken under. +pub const GRID_DEFAULT: u32 = 1024; + +/// What [`GRID_ENV`] accepts, inclusive. +/// +/// The top is where the knob stops meaning anything rather than where CUDA +/// stops accepting it (the driver allows 2^31−1 blocks in x): at 65,536 blocks +/// the RPX stride is 8.4 M against an expected hit at 2^20, so the search would +/// be overshoot and nothing else. 0 is refused because a launch of no blocks +/// scans nothing and the retry loop would spin forever. +pub const GRID_RANGE: std::ops::RangeInclusive = 1..=65_536; + +/// ★ The two launch knobs, read together so one line can print both. +#[derive(Clone, Copy, Debug, PartialEq, Eq)] +pub struct Knobs { + /// How many expected hit distances one launch covers. + pub scan: u32, + /// Blocks per launch. + pub grid: u32, +} + +impl Knobs { + /// The record posture: what an unset environment reads. + pub const DEFAULT: Self = Self { + scan: SCAN_FACTOR_DEFAULT, + grid: GRID_DEFAULT, + }; + + /// Nonces the walk advances per iteration — `grid × block_dim`, and the + /// term the early exit leaves behind as overshoot. + pub const fn stride(self, block_dim: u32) -> u64 { + self.grid as u64 * block_dim as u64 + } + + /// Nonces one launch covers: the expected hit distance `2^grinding_factor` + /// times the scan factor, clamped to `[MIN_BLOCK, MAX_BLOCK]`. + /// + /// `2^grinding_factor` can overflow u64 (factor 64), so saturate. + pub fn block(self, grinding_factor: u8) -> u64 { + let expected = 1u64.checked_shl(grinding_factor as u32).unwrap_or(u64::MAX); + expected + .saturating_mul(self.scan as u64) + .clamp(MIN_BLOCK, MAX_BLOCK) + } +} + +/// ★ The knobs for this process, read once and cached. +/// +/// Prints ONE line on every setting including the defaults, carrying both +/// knobs AND the stride each arm gets — the stride is the mechanism, so it is +/// printed as a read rather than left for the reader to multiply. A log with no +/// `★ GRIND KNOBS:` line is a run that never reached the device search. +/// +/// Aborts on a value outside its range for the reason `LAMBDA_VM_WHIR_HASH` +/// aborts on an unknown hash: a measurement taken under a silently ignored knob +/// is worse than no measurement. +pub fn knobs_in_effect() -> Knobs { + static KNOBS: std::sync::OnceLock = std::sync::OnceLock::new(); + *KNOBS.get_or_init(|| { + let knobs = Knobs { + scan: read_knob(SCAN_FACTOR_ENV, SCAN_FACTOR_DEFAULT, &SCAN_FACTOR_RANGE), + grid: read_knob(GRID_ENV, GRID_DEFAULT, &GRID_RANGE), + }; + println!( + "★ GRIND KNOBS: scan {} · grid {} · stride rpx {} / keccak {}", + knobs.scan, + knobs.grid, + knobs.stride(RPX_BLOCK_DIM), + knobs.stride(BLOCK_DIM), + ); + knobs + }) +} + +/// ★ The nonces one launch would cover at `grinding_factor` under the knobs in +/// effect — what [`search`] passes the kernel as `count`. +pub fn per_launch_block(grinding_factor: u8) -> u64 { + knobs_in_effect().block(grinding_factor) +} + +/// One knob, parsed from the environment or defaulted, refused outside its +/// range with the offending value and the range both named. +fn read_knob(env: &str, default: u32, range: &std::ops::RangeInclusive) -> u32 { + match std::env::var(env) { + Err(_) => default, + Ok(raw) => parse_knob(raw.trim(), range).unwrap_or_else(|| { + // eprintln then abort rather than a panic: a configuration error at + // startup, where the operator needs the accepted range and not a + // backtrace through the prover. + eprintln!( + "{env}={raw:?} is not a value this path accepts. Accepted: an integer in {}..={}.", + range.start(), + range.end() + ); + std::process::abort() + }), + } +} + +/// The accepted spellings: a plain decimal integer inside `range`. +fn parse_knob(raw: &str, range: &std::ops::RangeInclusive) -> Option { + let value: u32 = raw.parse().ok()?; + range.contains(&value).then_some(value) +} + +/// ★ What the driver says about filling this card with the RPX grind kernel. +/// +/// Every field is READ, not estimated: the residency question the grid knob +/// turns on is decided by registers per thread, and guessing that is how a +/// sweep gets sized against a card nobody measured. +#[derive(Clone, Copy, Debug)] +pub struct DeviceFill { + /// Multiprocessors on the device. + pub sm_count: u32, + /// The device's own ceiling on resident threads per multiprocessor. + pub max_threads_per_sm: u32, + /// Registers the RPX grind kernel uses per thread. + pub rpx_regs_per_thread: i32, + /// Blocks of [`RPX_BLOCK_DIM`] the driver will keep resident per + /// multiprocessor — the occupancy the register count actually buys. + pub rpx_blocks_per_sm: u32, + /// The block dim those blocks carry. + pub rpx_block_dim: u32, +} + +impl DeviceFill { + /// ⭐ Blocks that can be resident at once. A grid ABOVE this queues: the + /// extra blocks buy no parallelism and their stride is pure overshoot. + pub const fn resident_blocks(&self) -> u64 { + self.sm_count as u64 * self.rpx_blocks_per_sm as u64 + } + + /// Threads that can be resident at once, by the same reading. + pub const fn resident_threads(&self) -> u64 { + self.resident_blocks() * self.rpx_block_dim as u64 + } + + /// What fraction of the resident ceiling a grid of `grid` blocks asks for. + /// Above 1.0 the surplus queues. + pub fn fill(&self, grid: u32) -> f64 { + grid as f64 / self.resident_blocks() as f64 + } +} + +/// Reads [`DeviceFill`] off the driver, or `None` where there is no device. +pub fn device_fill() -> Option { + use cudarc::driver::sys::CUdevice_attribute; + let be = backend().ok()?; + let sm_count = be + .ctx + .attribute(CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT) + .ok()?; + let max_threads_per_sm = be + .ctx + .attribute(CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_MULTIPROCESSOR) + .ok()?; + let rpx_regs_per_thread = be.rpx_grind_search.num_regs().ok()?; + let rpx_blocks_per_sm = be + .rpx_grind_search + .occupancy_max_active_blocks_per_multiprocessor(RPX_BLOCK_DIM, 0, None) + .ok()?; + Some(DeviceFill { + sm_count: sm_count.max(0) as u32, + max_threads_per_sm: max_threads_per_sm.max(0) as u32, + rpx_regs_per_thread, + rpx_blocks_per_sm, + rpx_block_dim: RPX_BLOCK_DIM, + }) +} + /// Which outer hash the search runs. #[derive(Clone, Copy, Debug, PartialEq, Eq)] enum Arm { @@ -55,7 +294,30 @@ enum Arm { /// `inner_lanes` are the four **little-endian**-read `u64` lanes of the 32-byte /// inner hash — build them with `crypto::grinding::inner_hash_lanes`. pub fn generate_nonce_gpu(inner_lanes: &[u64; 4], grinding_factor: u8) -> Option { - search(Arm::Keccak256, inner_lanes, grinding_factor) + search( + Arm::Keccak256, + inner_lanes, + grinding_factor, + knobs_in_effect(), + ) +} + +/// [`generate_nonce_gpu`] at knobs given here rather than read from the +/// environment. +/// +/// ⛔ Not a second policy — the SAME [`search`], with the one thing the +/// environment would have decided passed in. It exists because the knobs are +/// cached in a `OnceLock`, so a process cannot compare two settings through the +/// environment: a sweep would need one process per arm, and two processes are +/// two device contexts, two cubin loads and two clock domains. The arms that +/// matter (the nonce is unmoved by the block size; what a grind costs at each +/// setting) are exact only when they share a process. +pub fn generate_nonce_gpu_at( + inner_lanes: &[u64; 4], + grinding_factor: u8, + knobs: Knobs, +) -> Option { + search(Arm::Keccak256, inner_lanes, grinding_factor, knobs) } /// Smallest nonce whose RPX grind head is `< limit`, or `None` when the CUDA @@ -79,7 +341,17 @@ pub fn generate_nonce_gpu(inner_lanes: &[u64; 4], grinding_factor: u8) -> Option /// algebraic digest's own output, which `digest_to_commitment` writes as four /// canonical big-endian `u64`s. pub fn generate_nonce_rpx_gpu(inner_felts: &[u64; 4], grinding_factor: u8) -> Option { - search(Arm::Rpx256, inner_felts, grinding_factor) + search(Arm::Rpx256, inner_felts, grinding_factor, knobs_in_effect()) +} + +/// [`generate_nonce_rpx_gpu`] at knobs given here rather than read from the +/// environment. See [`generate_nonce_gpu_at`] for why this exists. +pub fn generate_nonce_rpx_gpu_at( + inner_felts: &[u64; 4], + grinding_factor: u8, + knobs: Knobs, +) -> Option { + search(Arm::Rpx256, inner_felts, grinding_factor, knobs) } /// The range walk both arms share. @@ -89,7 +361,13 @@ pub fn generate_nonce_rpx_gpu(inner_felts: &[u64; 4], grinding_factor: u8) -> Op /// each launch scans a contiguous block several times that, from 0 upward, and /// the first block that hits yields the globally smallest valid nonce (the /// kernels `atomicMin` it). -fn search(arm: Arm, inner: &[u64; 4], grinding_factor: u8) -> Option { +/// +/// ★ **The returned nonce is therefore a function of the inner hash and the +/// grinding factor alone** — the blocks are contiguous from 0 and the first one +/// to hit returns its minimum, so the block SIZE cannot move it. That is what +/// `tests/grinding.rs::gpu_grind_returns_smallest_valid_nonce` pins, and it is +/// why sweeping [`SCAN_FACTOR_ENV`] moves no proof byte. +fn search(arm: Arm, inner: &[u64; 4], grinding_factor: u8, knobs: Knobs) -> Option { if !(GRIND_MIN_FACTOR..=64).contains(&grinding_factor) { return None; } @@ -103,15 +381,10 @@ fn search(arm: Arm, inner: &[u64; 4], grinding_factor: u8) -> Option { let stream = be.next_stream(); let inner_dev = stream.clone_htod(inner.as_slice()).ok()?; - // Per-launch block size: ~8× the expected hit distance, clamped so tiny - // factors still launch a full grid and huge factors don't ask for an - // absurd single block. `2^grinding_factor` can overflow u64 (factor 64), so - // saturate. - let expected = 1u64.checked_shl(grinding_factor as u32).unwrap_or(u64::MAX); - let count = expected.saturating_mul(8).clamp(1 << 18, 1 << 28); + let count = knobs.block(grinding_factor); let cfg = LaunchConfig { - grid_dim: (GRID_DIM, 1, 1), + grid_dim: (knobs.grid, 1, 1), block_dim: (block_dim, 1, 1), shared_mem_bytes: 0, }; @@ -146,3 +419,273 @@ fn search(arm: Arm, inner: &[u64; 4], grinding_factor: u8) -> Option { base = base.checked_add(count)?; } } + +/// ⛔ DIAGNOSTIC: what one grind EXECUTED, beside what it returned. +/// +/// Every field is counted ON THE DEVICE by the threads that did the work, so +/// `executed - (h + stride)` is a READ rather than a model. See +/// [`search_counted`]. +#[derive(Clone, Copy, Debug, Default, PartialEq, Eq)] +pub struct GrindCounts { + /// The nonce the search returned — must equal the shipped kernel's. + pub nonce: u64, + /// Permutations every thread of every launch actually ran. + pub executed: u64, + /// The most iterations any single thread ran, over every launch. + pub max_iters: u64, + /// Threads that left by the loop bound rather than by the early exit. + pub ran_to_end: u64, + /// Launches this search made: the hit's block, plus every miss before it. + pub launches: u64, +} + +/// ⛔ THE DIAGNOSTIC TWIN OF [`search`], RPX ONLY, ON NO PROVING PATH. +/// +/// Identical walk, identical exits, identical `atomicMin` — plus three device +/// counters. It exists to separate two explanations of the same slow launch +/// that no stopwatch can tell apart: MORE PERMUTATIONS (a thread that did not +/// observe the early exit and kept hashing) against THE SAME PERMUTATIONS MORE +/// SLOWLY (a cost outside this loop entirely). +/// +/// ⚠ ITS OWN ADMISSIBILITY IS THE CALLER'S JOB, and it is not optional: a +/// counted kernel with different register pressure has different occupancy and +/// therefore measures a different kernel. The bench asserts this twin +/// reproduces the shipped kernel's milliseconds per seed within the measured +/// noise floor, and reports the instrument as having changed the phenomenon if +/// it does not. +/// +/// The counters accumulate ACROSS the launches of one search — `atomicAdd` for +/// the two sums, `atomicMax` for the deepest thread — so a miss-and-relaunch +/// seed reports the whole search and not only its last block. `launches` says +/// how many blocks that was. +/// +/// Returns `None` for the same reasons [`search`] does, plus a factor outside +/// the supported range. +pub fn search_counted(inner: &[u64; 4], grinding_factor: u8, knobs: Knobs) -> Option { + if !(GRIND_MIN_FACTOR..=64).contains(&grinding_factor) { + return None; + } + let limit: u64 = 1u64 << (64 - grinding_factor); + + let be = backend().ok()?; + let stream = be.next_stream(); + let inner_dev = stream.clone_htod(inner.as_slice()).ok()?; + + let count = knobs.block(grinding_factor); + let cfg = LaunchConfig { + grid_dim: (knobs.grid, 1, 1), + block_dim: (RPX_BLOCK_DIM, 1, 1), + shared_mem_bytes: 0, + }; + + let sentinel = [u64::MAX]; + let mut result_dev = stream.clone_htod(&sentinel).ok()?; + // One slot each, zeroed once and accumulated into by every launch. + let zeros = [0u64; 3]; + let mut counts_dev = stream.clone_htod(&zeros).ok()?; + + let mut base: u64 = 0; + let mut launches: u64 = 0; + loop { + stream.memcpy_htod(&sentinel, &mut result_dev).ok()?; + launches += 1; + // SAFETY: the same contract as `search`'s launch, with three more + // device words the kernel only ever adds into or maxes against. + unsafe { + stream + .launch_builder(&be.rpx_grind_search_counted) + .arg(&inner_dev) + .arg(&limit) + .arg(&base) + .arg(&count) + .arg(&mut result_dev) + .arg(&mut counts_dev) + .launch(cfg) + .ok()?; + } + let host = stream.clone_dtoh(&result_dev).ok()?; + stream.synchronize().ok()?; + if host[0] != u64::MAX { + let counts = stream.clone_dtoh(&counts_dev).ok()?; + stream.synchronize().ok()?; + return Some(GrindCounts { + nonce: host[0], + executed: counts[0], + max_iters: counts[1], + ran_to_end: counts[2], + launches, + }); + } + base = base.checked_add(count)?; + } +} + +#[cfg(test)] +mod tests { + //! The knobs' parser and their launch arithmetic, card-free: nothing here + //! touches `backend()`, so these run wherever the crate compiles. + //! + //! `knobs_in_effect()` itself is deliberately NOT exercised — it caches in + //! a `OnceLock`, so a test that set a variable would fix the value for + //! every other test in the binary and the second assertion would pass on + //! the first one's cache. That is also why the sweep goes through + //! `generate_nonce_rpx_gpu_at` rather than through the environment. + + use super::{ + BLOCK_DIM, GRID_DEFAULT, GRID_RANGE, Knobs, MAX_BLOCK, MIN_BLOCK, RPX_BLOCK_DIM, + SCAN_FACTOR_DEFAULT, SCAN_FACTOR_RANGE, parse_knob, + }; + + /// A knob at a named scan factor, the grid left at the record posture. + fn at_scan(scan: u32) -> Knobs { + Knobs { + scan, + grid: GRID_DEFAULT, + } + } + + /// Every arm of the scan sweep parses, and so does every arm of the grid + /// sweep — both directions, since the grid's optimum can sit on either side + /// of the default. + #[test] + fn the_parser_accepts_every_arm_of_both_sweeps() { + for k in [8u32, 4, 2, 1] { + assert_eq!( + parse_knob(&k.to_string(), &SCAN_FACTOR_RANGE), + Some(k), + "scan factor {k} is a sweep arm and must parse" + ); + } + for g in [256u32, 512, 1024, 2048, 4096] { + assert_eq!( + parse_knob(&g.to_string(), &GRID_RANGE), + Some(g), + "grid {g} is a sweep arm and must parse" + ); + } + assert_eq!(parse_knob("64", &SCAN_FACTOR_RANGE), Some(64), "scan top"); + assert_eq!(parse_knob("65536", &GRID_RANGE), Some(65_536), "grid top"); + } + + /// The refusing half, executed on both ranges. Each of these would + /// otherwise reach the launch and become a DIFFERENT search that no log + /// could distinguish from the one the operator asked for. + #[test] + fn the_parser_refuses_what_would_silently_change_the_search() { + for raw in ["0", "abc", "", "-1", "8.0", "0x8", " ", "4294967296"] { + assert_eq!( + parse_knob(raw, &SCAN_FACTOR_RANGE), + None, + "scan {raw:?} must be refused, not defaulted" + ); + assert_eq!( + parse_knob(raw, &GRID_RANGE), + None, + "grid {raw:?} must be refused, not defaulted" + ); + } + // Each range refuses just past its own top, and the two tops differ — + // so a range accidentally shared between the knobs reddens here. + assert_eq!(parse_knob("65", &SCAN_FACTOR_RANGE), None, "scan past top"); + assert_eq!(parse_knob("65537", &GRID_RANGE), None, "grid past top"); + assert_eq!( + parse_knob("1024", &SCAN_FACTOR_RANGE), + None, + "a grid value is not a scan factor" + ); + } + + /// The block size at the production grinding factor, written out as + /// arithmetic rather than as the expression under test. + /// + /// ⇒ A drifted default or a dropped multiply reddens here by value. + #[test] + fn the_block_is_the_hit_distance_times_the_scan_factor() { + assert_eq!(at_scan(8).block(20), 8 * 1_048_576, "2^20 * 8 = 2^23"); + assert_eq!(at_scan(4).block(20), 4 * 1_048_576, "2^20 * 4 = 2^22"); + assert_eq!(at_scan(2).block(20), 2 * 1_048_576, "2^20 * 2 = 2^21"); + assert_eq!(at_scan(1).block(20), 1_048_576, "2^20 * 1 = 2^20"); + } + + /// ★ The stride — the term the kernels' early exit leaves behind as + /// overshoot, and the whole mechanism of the grid knob. Written as the + /// products themselves, so a changed block dim reddens rather than being + /// absorbed by the formula under test. + #[test] + fn the_stride_is_the_grid_times_the_block_dim() { + assert_eq!(RPX_BLOCK_DIM, 128, "the RPX arm's block dim"); + assert_eq!(BLOCK_DIM, 256, "the keccak arm's block dim"); + assert_eq!( + Knobs::DEFAULT.stride(RPX_BLOCK_DIM), + 131_072, + "1024 blocks * 128 threads" + ); + assert_eq!( + Knobs::DEFAULT.stride(BLOCK_DIM), + 262_144, + "1024 blocks * 256 threads" + ); + // Both directions of the sweep, as products. + for (grid, rpx) in [ + (256u32, 32_768u64), + (512, 65_536), + (1024, 131_072), + (2048, 262_144), + (4096, 524_288), + ] { + let knobs = Knobs { scan: 8, grid }; + assert_eq!( + knobs.stride(RPX_BLOCK_DIM), + rpx, + "grid {grid} on the RPX arm" + ); + } + } + + /// The defaults are the record posture: every measurement in the campaign + /// was taken at a 2^23 block over a 1024-block grid, and only an ABBA may + /// move either. + #[test] + fn the_defaults_are_the_record_posture() { + assert_eq!(SCAN_FACTOR_DEFAULT, 8, "the record posture's scan factor"); + assert_eq!(GRID_DEFAULT, 1024, "the record posture's grid"); + assert_eq!(Knobs::DEFAULT.scan, 8, "DEFAULT carries the scan factor"); + assert_eq!(Knobs::DEFAULT.grid, 1024, "DEFAULT carries the grid"); + assert_eq!( + Knobs::DEFAULT.block(20), + 1 << 23, + "the record posture's per-launch block at grinding factor 20" + ); + } + + /// Both ends of the clamp still bind, including the u64 overflow the + /// saturating multiply exists for. The grid does not enter the clamp at + /// all — it sizes the stride, not the block — and that separation is the + /// thing this asserts. + #[test] + fn the_clamp_binds_at_both_ends_and_the_grid_does_not_touch_it() { + assert_eq!( + at_scan(1).block(12), + MIN_BLOCK, + "the min-factor gate's own factor, at the narrowest scan, floors" + ); + assert_eq!(at_scan(8).block(26), MAX_BLOCK, "2^26 * 8 = 2^29 ceils"); + assert_eq!( + at_scan(8).block(64), + MAX_BLOCK, + "2^64 saturates before the multiply, then ceils" + ); + assert!( + (MIN_BLOCK..=MAX_BLOCK).contains(&Knobs::DEFAULT.block(20)), + "the record posture sits strictly inside the clamp, so neither end \ + is silently setting it" + ); + for grid in [256u32, 1024, 4096] { + assert_eq!( + Knobs { scan: 8, grid }.block(20), + 1 << 23, + "the grid must not move the per-launch block" + ); + } + } +} diff --git a/crypto/math-cuda/src/sumcheck.rs b/crypto/math-cuda/src/sumcheck.rs index 1f1dc0b42..538c75178 100644 --- a/crypto/math-cuda/src/sumcheck.rs +++ b/crypto/math-cuda/src/sumcheck.rs @@ -584,6 +584,7 @@ pub fn evaluate_many_base( // of the card: the caller's fallback is to evaluate the columns one at // a time, which needs almost nothing. let Some(_room) = crate::device::reserve(group_len as u64 * per_column) else { + crate::device::note_device_fallback(); return Err(cudarc::driver::DriverError( cudarc::driver::sys::CUresult::CUDA_ERROR_OUT_OF_MEMORY, )); @@ -753,6 +754,7 @@ impl DeviceFactors { let be = backend()?; let Some(room) = be.reserve(factors.len() as u64 * span as u64 * 8) else { + crate::device::note_device_fallback(); return Err(cudarc::driver::DriverError( cudarc::driver::sys::CUresult::CUDA_ERROR_OUT_OF_MEMORY, )); @@ -820,6 +822,7 @@ impl DeviceFactors { // What stays: the factors. The base columns they are gathered from are // a third of that and are freed as soon as the gather has read them. let Some(room) = be.reserve(width as u64 * rows as u64 * 24) else { + crate::device::note_device_fallback(); return Err(cudarc::driver::DriverError( cudarc::driver::sys::CUresult::CUDA_ERROR_OUT_OF_MEMORY, )); diff --git a/crypto/math-cuda/src/whir.rs b/crypto/math-cuda/src/whir.rs index dfb8b8993..40785c084 100644 --- a/crypto/math-cuda/src/whir.rs +++ b/crypto/math-cuda/src/whir.rs @@ -5,14 +5,14 @@ //! one NTT onto the blown-up domain, then the strided-coset leaf hash and the //! Merkle tree. Parity against that pipeline is checked by `tests/whir_commit.rs`. -use std::sync::Arc; +use std::sync::{Arc, Mutex, Once, Weak}; use cudarc::driver::{CudaSlice, CudaStream, LaunchConfig, PushKernelArg}; -use core::sync::atomic::{AtomicU64, Ordering}; +use core::sync::atomic::{AtomicBool, AtomicU64, Ordering}; use crate::Result; -use crate::device::{alloc_or_trim, backend}; +use crate::device::{DeviceReservation, alloc_or_trim, backend}; /// Leaf-hash passes over a codeword — one per tree actually built. /// @@ -41,6 +41,285 @@ pub fn reset_leaf_hash_calls() { /// it rather than to whoever ran alongside. type BuildCount = Arc; +/// Trees ASSEMBLED, process-wide — the twin of [`LEAF_HASH_CALLS`] and no +/// longer the same number as it. +/// +/// ★ Before the leaf layer was retained these two counted the same event, which +/// is exactly why the retention needs both: a tree is still assembled for every +/// opening (the inner levels are rebuilt), but its LEAF PASS is skipped when a +/// matching layer is in hand. `leaf_hash_calls` well below `tree_builds` is the +/// retention working; the two equal is the retention not taken. +static TREE_BUILDS: AtomicU64 = AtomicU64::new(0); + +pub fn tree_builds() -> u64 { + TREE_BUILDS.load(Ordering::Relaxed) +} + +/// Leaf layers this process asked to retain, and what happened. +/// +/// ⛔ REPORTED, ALWAYS. A run that retains nothing must say so on its own line +/// rather than reading as a lever that quietly did not fire: the refusal path +/// below is what makes the scheme safe at full size, so how often it fires is +/// the first thing any reading of the lever has to know. +static RETAIN_ADMITTED: AtomicU64 = AtomicU64::new(0); +static RETAIN_REFUSED: AtomicU64 = AtomicU64::new(0); +static RETAIN_BYTES_ASKED: AtomicU64 = AtomicU64::new(0); +static RETAIN_BYTES_ADMITTED: AtomicU64 = AtomicU64::new(0); +/// Bytes the budget still had when it first refused — zero if it never did. +static RETAIN_FIRST_REFUSAL_HEADROOM: AtomicU64 = AtomicU64::new(0); +/// Leaf passes SKIPPED because a matching layer was in hand. The saving, counted +/// where it happens rather than inferred from two other counters. +static LEAF_PASSES_SAVED: AtomicU64 = AtomicU64::new(0); +/// ★ The SIMULTANEOUS retained device bytes held RIGHT NOW — `fetch_add` on +/// admit, `fetch_sub` in [`RetainedLeaves`]'s `Drop`. Unlike +/// [`RETAIN_BYTES_ADMITTED`], which only ever rises (cumulative over the run), +/// this rises and falls with the live layers, so it is the TRUE footprint: the +/// quantity that contends with argue for the shared budget, and the one an +/// eviction gives back. wt16 reported the cumulative 39,057 MiB and had no name +/// for the ~2 GiB that actually bound. +static RETAIN_BYTES_LIVE: AtomicU64 = AtomicU64::new(0); +/// ★ The PEAK of [`RETAIN_BYTES_LIVE`] over the run — the number the print wants. +/// The instantaneous live count is ~0 by the time the base readback prints (the +/// codewords, and their layers, have dropped), so the reporting quantity is this +/// high-water: the largest simultaneous retained footprint the run ever held. +static RETAIN_BYTES_LIVE_PEAK: AtomicU64 = AtomicU64::new(0); + +/// What the leaf-layer retention did this process: admitted, refused, the bytes +/// on each side, the headroom at the first refusal, and the leaf passes skipped. +pub fn retention_report() -> (u64, u64, u64, u64, u64, u64) { + ( + RETAIN_ADMITTED.load(Ordering::Relaxed), + RETAIN_REFUSED.load(Ordering::Relaxed), + RETAIN_BYTES_ASKED.load(Ordering::Relaxed), + RETAIN_BYTES_ADMITTED.load(Ordering::Relaxed), + RETAIN_FIRST_REFUSAL_HEADROOM.load(Ordering::Relaxed), + LEAF_PASSES_SAVED.load(Ordering::Relaxed), + ) +} + +/// ★ The retained leaf-layer device bytes held RIGHT NOW — the live footprint, +/// as opposed to `retention_report`'s cumulative `admitted` bytes. This is the +/// quantity that contends with argue for the budget; a reading well below the +/// cumulative total is the layers being freed as their codewords drop (or an +/// eviction giving them back). +pub fn retained_bytes_live() -> u64 { + RETAIN_BYTES_LIVE.load(Ordering::Relaxed) +} + +/// ★ The PEAK simultaneous retained device footprint over the run — the largest +/// [`retained_bytes_live`] ever reached. This is the number that bound argue in +/// wt16 (~2 GiB); the instantaneous count is ~0 once the codewords drop, so this +/// is what a report prints. Process-wide and monotone (a high-water), so a fresh +/// process — which is how the launcher runs each prove — starts it at 0. +pub fn retained_bytes_peak() -> u64 { + RETAIN_BYTES_LIVE_PEAK.load(Ordering::Relaxed) +} + +/// A leaf layer kept past the call that built it — and NOTHING else. +/// +/// ⛔ NOT A TREE. H4 kept the whole node array, `2·num_leaves − 1` nodes, and +/// lost: at fold width `k` that is `C · 2^(6−k)` bytes against this layer's +/// `C · 2^(5−k)`, so at the production `k = 4` H4 held half a base codeword per +/// commitment where this holds a quarter. The inner levels are cheap to rebuild +/// — one permutation a node against two per leaf on a base codeword and six on +/// an extension one — so the expensive two thirds is what is kept. +/// +/// ⛔ THE KEY IS PART OF THE OBJECT. A leaf is the `2^log_folding` coset that +/// folds onto one position, so a layer is valid ONLY for the width it was built +/// at, and only for the hash that built it. `paths()` takes `log_folding` as a +/// PARAMETER, and `whir_tree_cache.rs`'s blocking test opens the same codeword +/// at two widths on purpose: served across widths, this would hand back paths +/// that are internally consistent and WRONG, which is the outcome that file +/// exists to forbid. An exact match or a rebuild; there is no near miss. +struct RetainedLeaves { + nodes: CudaSlice, + log_folding: usize, + hash: crate::DeviceHash, + num_leaves: usize, + bytes: u64, +} + +impl Drop for RetainedLeaves { + /// The other half of the live-footprint accounting: whatever admitted the + /// layer added to [`RETAIN_BYTES_LIVE`] is given back when the layer drops — + /// with its codeword, or when an eviction sets the slot to `None`. `nodes` + /// (a `CudaSlice`) frees its device bytes on the line after this returns. + fn drop(&mut self) { + RETAIN_BYTES_LIVE.fetch_sub(self.bytes, Ordering::Relaxed); + } +} + +/// Retention evictions this process has done, and the bytes they gave back — the +/// evictable scheme working under pressure. `evicted > 0` with `device fallbacks +/// 0` is argue getting its budget FROM the retention instead of from the host. +static RETAIN_EVICTED: AtomicU64 = AtomicU64::new(0); +static RETAIN_BYTES_EVICTED: AtomicU64 = AtomicU64::new(0); + +/// (evictions, bytes evicted) this process — see [`RETAIN_EVICTED`]. +pub fn retention_evictions() -> (u64, u64) { + ( + RETAIN_EVICTED.load(Ordering::Relaxed), + RETAIN_BYTES_EVICTED.load(Ordering::Relaxed), + ) +} + +/// Whether the leaf-layer retention is enabled this process. `LFM_WHIR_RETENTION=0` +/// turns it OFF (capture becomes a no-op), which is how the ABBA's control arm +/// runs off the SAME binary as the retention arm — the two differ only by this +/// flag. Read once and cached; default ON. +fn retention_enabled() -> bool { + static ENABLED: std::sync::OnceLock = std::sync::OnceLock::new(); + *ENABLED.get_or_init(|| { + std::env::var("LFM_WHIR_RETENTION") + .map(|v| v != "0") + .unwrap_or(true) + }) +} + +/// A handle the allocator can walk to reclaim a codeword's retained layer under +/// pressure. Both are `Weak`, so the registry never keeps a codeword alive: +/// `leaves` points at the codeword's `Arc>>` — the +/// MUTEX outlives any single layer, so this entry SURVIVES an eviction and a +/// later re-capture refills the same `Option` — and `room` at its reservation, so +/// the freed bytes can be `shrink`-ed back into the budget. A dead `leaves` +/// upgrade means the codeword is gone and the eviction walk prunes the entry. +struct RetentionHandle { + leaves: Weak>>, + room: Weak, +} + +/// Every retaining codeword's handle, registered once at its first capture and +/// pruned when the codeword drops. Walked only on the rare eviction path. +static RETENTION_REGISTRY: Mutex> = Mutex::new(Vec::new()); + +/// Install [`evict_retained_layers`] as the allocator's evictor, once. +fn ensure_evictor_installed() { + static INSTALLED: Once = Once::new(); + INSTALLED.call_once(|| crate::device::set_retention_evictor(evict_retained_layers)); +} + +/// Register a codeword's layer slot + reservation for eviction. Called OUTSIDE +/// the layer lock (it takes only the registry lock), so the one lock nesting in +/// the system is the evictor's registry→layer and there is no cycle. +fn register_retained(leaves: &Arc>>, room: &Arc) { + let handle = RetentionHandle { + leaves: Arc::downgrade(leaves), + room: Arc::downgrade(room), + }; + if let Ok(mut reg) = RETENTION_REGISTRY.lock() { + reg.push(handle); + } +} + +/// ⛔ THE EVICTOR — installed into `device::reserve`, called on its budget-miss +/// path on the RESERVING thread. Frees at least `target` bytes of BUDGET by +/// dropping retained leaf layers, FIFO (oldest first), pruning dead entries as it +/// goes; returns the bytes freed. +/// +/// SAFETY, the load-bearing facts: +/// - NO REENTRANCY. It calls only `room.shrink` (a lock-free `be.reserved` +/// subtract) and drops `RetainedLeaves` (a counter subtract + a stream-ordered +/// free). It NEVER calls `reserve`/`grow`, so it cannot re-enter the allocator +/// that called it. +/// - NO UAF. It takes the layer out UNDER the codeword's own `leaves` mutex, so a +/// concurrent [`serve_retained_leaves`] either ran first (its `memcpy_dtod` is +/// already enqueued on the codeword's stream) or sees `None` and rebuilds. The +/// evicted `nodes` was allocated on that SAME stream, and cudarc 0.19.4 +/// `CudaSlice::drop` (core.rs:776-795) first waits on the slice's own read/write +/// events, then frees on the slice's OWN stream — `free_async(ptr, +/// self.stream.cu_stream)` when `has_async_alloc`, else `synchronize()` + +/// `free_sync`. So the free is ordered after any serve copy on that stream (or +/// a full device sync). The only unsafe shape — an async free on a DIFFERENT +/// stream than the copy — cannot arise here: both are the codeword's stream. +/// - LOCK ORDER. Registry mutex, then ONE layer mutex at a time (released before +/// the next); `be.reserved` is lock-free. Capture registers OUTSIDE the layer +/// lock, so nothing ever holds layer→registry — acyclic. +/// - `shrink` frees the BUDGET, not the bytes: never-purge keeps the raw device +/// allocation pooled for reuse, and the budget is exactly what argue's +/// `reserve` is gated on ([[gpu-two-vramgates-overlap]]). +fn evict_retained_layers(target: u64) -> u64 { + let mut freed = 0u64; + let mut reg = match RETENTION_REGISTRY.lock() { + Ok(reg) => reg, + Err(poisoned) => poisoned.into_inner(), + }; + reg.retain(|handle| { + let Some(leaves) = handle.leaves.upgrade() else { + return false; // the codeword is gone; prune this dead entry + }; + if freed >= target { + return true; // enough freed; keep the rest for next time + } + // Take the layer out under its own mutex — a concurrent serve holds this + // same lock across its copy enqueue, so the free cannot race it. + let taken = match leaves.lock() { + Ok(mut slot) => slot.take(), + Err(poisoned) => poisoned.into_inner().take(), + }; + if let Some(layer) = taken { + let bytes = layer.bytes; + drop(layer); // RetainedLeaves::drop: live -= bytes; nodes freed (stream-ordered) + if let Some(room) = handle.room.upgrade() { + room.shrink(bytes); // give the BUDGET back to argue + } + freed += bytes; + RETAIN_EVICTED.fetch_add(1, Ordering::Relaxed); + RETAIN_BYTES_EVICTED.fetch_add(bytes, Ordering::Relaxed); + } + true // keep the entry: the mutex lives with the codeword and may refill + }); + freed +} + +/// May a layer built under `kept` be served for a tree asked for under `want`? +/// +/// ⛔ A FREE FUNCTION ON PURPOSE. Everything else on this path needs a device, +/// so this predicate would otherwise be checkable only on the box — and it is +/// the one piece of the retention whose failure is not slowness but WRONG +/// PATHS, internally consistent and verifying against nothing. Lifted out, it +/// takes unit cases on any machine. +/// +/// All three parts must agree. `num_leaves` is not redundant with +/// `log_folding`: the same width over a different codeword length is a +/// different tree, and a clone of a `DeviceCodeword` shares this layer. +fn leaf_key_matches( + kept: (usize, crate::DeviceHash, usize), + want: (usize, crate::DeviceHash, usize), +) -> bool { + kept.0 == want.0 && kept.1 == want.1 && kept.2 == want.2 +} + +#[cfg(test)] +mod leaf_key_tests { + use super::leaf_key_matches; + use crate::DeviceHash; + + /// The exact match is the only match, and each part is shown to matter on + /// its own — a predicate that only ever saw agreement would pass while + /// ignoring two of its three arguments. + #[test] + fn a_retained_leaf_layer_is_served_only_under_its_own_key() { + let kept = (4usize, DeviceHash::Rpx256, 1024usize); + assert!(leaf_key_matches(kept, kept), "an exact match must serve"); + assert!( + !leaf_key_matches(kept, (2, DeviceHash::Rpx256, 1024)), + "a different FOLD WIDTH describes a different tree: serving it would hand back paths of the wrong depth that verify against themselves" + ); + assert!( + !leaf_key_matches(kept, (4, DeviceHash::Keccak256, 1024)), + "a different HASH describes a different tree" + ); + assert!( + !leaf_key_matches(kept, (4, DeviceHash::Rpx256, 512)), + "the same width over a different codeword length is a different tree" + ); + assert!( + !leaf_key_matches(kept, (2, DeviceHash::Keccak256, 512)), + "and all three wrong is still not a match" + ); + } +} + use crate::merkle::{build_inner_tree_levels, keccak_launch_cfg}; /// A codeword the device holds, base-field or ext3. @@ -57,12 +336,25 @@ pub struct DeviceCodeword { /// Leaf-hash passes this codeword has paid for: one per tree built, so /// two for a commitment that is opened — the root's and the paths'. builds: BuildCount, + /// Leaf-hash passes this codeword actually paid for. Diverges from + /// [`builds`](Self::tree_builds) exactly when a retained layer was served. + leaf_passes: BuildCount, + /// ★ The leaf layer kept past the call that built it, with the key it is + /// valid under. `Arc` for the same reason `room` is one: `DeviceCodeword` + /// is `Clone` and the folds share the original's accounting, so a clone + /// must share the layer rather than silently rebuild beside it. + leaves: Arc>>, /// The room the chain promised itself: this codeword and the folds that /// halve it, shared with those folds because they live inside it. /// /// A tree is NOT in this number, because a tree is never held past the /// call that builds it — see [`with_tree`](Self::with_tree). room: Arc, + /// Whether this codeword has been registered with the eviction registry. + /// Set once, on the first capture, so a rebuild-after-eviction re-capture + /// does not push a duplicate handle. `Arc` so all clones share the one flag, + /// as they share `leaves` and `room`. + registered: Arc, } impl DeviceCodeword { @@ -104,40 +396,195 @@ impl DeviceCodeword { // kernel below, the inner nodes by the level loop after it. let mut nodes = unsafe { crate::device::alloc_or_trim::(&self.stream, total_nodes * 32) }?; - { - let leaves_offset = (num_leaves - 1) * 32; - let mut leaves = nodes.slice_mut(leaves_offset..leaves_offset + num_leaves * 32); - let num_leaves_u64 = num_leaves as u64; - let block = 1u64 << log_folding; - // ★ The hash is chosen HERE, not by the host backend that will - // label the result. `hash` is the key the caller's `WhirHash` - // supplied, so a tree labelled RPX was hashed by RPX's kernels or - // was not built here at all. - let kernel = match (hash, self.base) { - (crate::DeviceHash::Keccak256, true) => &be.keccak256_leaves_base_coset, - (crate::DeviceHash::Keccak256, false) => &be.keccak256_leaves_ext3_coset, - (crate::DeviceHash::Rpx256, true) => &be.rpx_leaves_base_coset, - (crate::DeviceHash::Rpx256, false) => &be.rpx_leaves_ext3_coset, - (other, _) => { - unimplemented!("no WHIR kernels for {} ({other:?})", other.name()) + let leaves_offset = (num_leaves - 1) * 32; + // ★ THE ONE BRANCH THIS CHANGE ADDS. A matching layer means the leaf + // pass is a device-to-device copy instead of a hash of every element; + // the inner levels below are built either way, so `build_tree` still + // returns the same tree it always did and `build_inner_tree_levels` is + // untouched. + let served = + self.serve_retained_leaves(&mut nodes, leaves_offset, num_leaves, log_folding, hash)?; + if !served { + { + let mut leaves = nodes.slice_mut(leaves_offset..leaves_offset + num_leaves * 32); + let num_leaves_u64 = num_leaves as u64; + let block = 1u64 << log_folding; + // ★ The hash is chosen HERE, not by the host backend that will + // label the result. `hash` is the key the caller's `WhirHash` + // supplied, so a tree labelled RPX was hashed by RPX's kernels or + // was not built here at all. + let kernel = match (hash, self.base) { + (crate::DeviceHash::Keccak256, true) => &be.keccak256_leaves_base_coset, + (crate::DeviceHash::Keccak256, false) => &be.keccak256_leaves_ext3_coset, + (crate::DeviceHash::Rpx256, true) => &be.rpx_leaves_base_coset, + (crate::DeviceHash::Rpx256, false) => &be.rpx_leaves_ext3_coset, + (other, _) => { + unimplemented!("no WHIR kernels for {} ({other:?})", other.name()) + } + }; + unsafe { + self.stream + .launch_builder(kernel) + .arg(self.buffer.as_ref()) + .arg(&num_leaves_u64) + .arg(&block) + .arg(&mut leaves) + .launch(keccak_launch_cfg(num_leaves_u64))?; } - }; - unsafe { - self.stream - .launch_builder(kernel) - .arg(self.buffer.as_ref()) - .arg(&num_leaves_u64) - .arg(&block) - .arg(&mut leaves) - .launch(keccak_launch_cfg(num_leaves_u64))?; } + // The borrow of `nodes` ends at the brace above, which is what lets + // the capture below read the region it just wrote. + // The pass was PAID here, so it is counted here — and the layer is + // offered for retention while it is in hand. + LEAF_HASH_CALLS.fetch_add(1, Ordering::Relaxed); + self.leaf_passes.fetch_add(1, Ordering::Relaxed); + self.capture_leaves(&nodes, leaves_offset, num_leaves, log_folding, hash); + } else { + LEAF_PASSES_SAVED.fetch_add(1, Ordering::Relaxed); } build_inner_tree_levels(self.stream.as_ref(), be, &mut nodes, num_leaves, hash)?; - LEAF_HASH_CALLS.fetch_add(1, Ordering::Relaxed); + TREE_BUILDS.fetch_add(1, Ordering::Relaxed); self.builds.fetch_add(1, Ordering::Relaxed); Ok((nodes, num_leaves)) } + /// Copy a retained layer into the node buffer's leaf region, if one matches. + /// + /// ⛔ THE MATCH IS EXACT ON BOTH KEY PARTS AND ON THE SHAPE. A layer built + /// at another fold width describes a different tree, and one built under + /// another hash describes a different tree again; either served here would + /// produce authentication paths that verify against themselves and against + /// nothing else. + fn serve_retained_leaves( + &self, + nodes: &mut CudaSlice, + leaves_offset: usize, + num_leaves: usize, + log_folding: usize, + hash: crate::DeviceHash, + ) -> Result { + let held = match self.leaves.lock() { + Ok(held) => held, + // A poisoned lock is not a reason to serve a layer nobody can + // vouch for: rebuild instead. + Err(_) => return Ok(false), + }; + let Some(kept) = held.as_ref() else { + return Ok(false); + }; + if !leaf_key_matches( + (kept.log_folding, kept.hash, kept.num_leaves), + (log_folding, hash, num_leaves), + ) { + return Ok(false); + } + let mut region = nodes.slice_mut(leaves_offset..leaves_offset + num_leaves * 32); + self.stream.memcpy_dtod(&kept.nodes, &mut region)?; + Ok(true) + } + + /// Offer the layer just hashed for retention, and take the answer. + /// + /// ⛔ ALLOCATE, THEN PROMISE, AND GIVE THE PROMISE BACK IF THE ALLOCATION + /// FAILED — in that order. Promising first and then failing to allocate + /// would leave the budget permanently short by bytes nothing holds, and + /// every later commitment would be refused because of it. `grow` returns + /// false and changes nothing when the budget will not take it, and + /// `shrink` gives back what an allocation could not use, so neither + /// direction can leak. + /// + /// ⚠ EVERY FAILURE PATH IS "NO RETENTION", NEVER AN ERROR. This is a cache: + /// the tree it would have saved is built anyway, the commit cannot fail + /// because of it, and `commit_stacked`'s device attempt cannot start + /// returning `None` — which is what would send a commitment to the host and + /// cost the 1.5 GiB per fallen-back chain that killed H4. + fn capture_leaves( + &self, + nodes: &CudaSlice, + leaves_offset: usize, + num_leaves: usize, + log_folding: usize, + hash: crate::DeviceHash, + ) { + // ⛔ LFM_WHIR_RETENTION=0 disables the retention entirely (the ABBA's + // control arm, off the SAME binary): capture is a no-op, nothing is + // retained, and the print reads admitted 0. Default ON. + if !retention_enabled() { + return; + } + // Install the evictor before any layer exists, so a later `reserve` that + // needs the budget can reclaim one. Idempotent (a `Once`). + ensure_evictor_installed(); + let Ok(mut held) = self.leaves.lock() else { + return; + }; + if held.is_some() { + return; + } + let bytes = (num_leaves as u64) * 32; + RETAIN_BYTES_ASKED.fetch_add(bytes, Ordering::Relaxed); + // SAFETY: every byte is written by the copy below before anything reads + // it, and the slice is dropped on every path that does not copy. + let Ok(mut copy) = (unsafe { alloc_or_trim::(&self.stream, num_leaves * 32) }) else { + Self::note_refusal(); + return; + }; + if !self.room.grow(bytes) { + Self::note_refusal(); + return; + } + let region = nodes.slice(leaves_offset..leaves_offset + num_leaves * 32); + if self.stream.memcpy_dtod(®ion, &mut copy).is_err() { + self.room.shrink(bytes); + Self::note_refusal(); + return; + } + RETAIN_ADMITTED.fetch_add(1, Ordering::Relaxed); + RETAIN_BYTES_ADMITTED.fetch_add(bytes, Ordering::Relaxed); + // The live footprint rises here and falls in `RetainedLeaves`'s Drop, so + // the two balance over the layer's lifetime; the peak is kept for the + // print, since the instantaneous count is ~0 by the time it is read. + let now_live = RETAIN_BYTES_LIVE.fetch_add(bytes, Ordering::Relaxed) + bytes; + RETAIN_BYTES_LIVE_PEAK.fetch_max(now_live, Ordering::Relaxed); + *held = Some(RetainedLeaves { + nodes: copy, + log_folding, + hash, + num_leaves, + bytes, + }); + // Release the layer lock BEFORE touching the registry, so the only lock + // nesting anywhere is the evictor's registry→layer — acyclic. Register + // ONCE: the handle is a Weak to this mutex, survives an eviction, and a + // re-capture refills the same slot, so a second push would only duplicate. + drop(held); + if !self.registered.swap(true, Ordering::Relaxed) { + register_retained(&self.leaves, &self.room); + } + } + + /// One refusal, with the headroom the budget had the FIRST time it happened + /// — the number that says whether the scheme was short by a little or by a + /// lot, and which a later refusal would overwrite with a smaller one. + fn note_refusal() { + RETAIN_REFUSED.fetch_add(1, Ordering::Relaxed); + // `max(1)` so "no headroom at all" is still distinguishable from "never + // refused", which is what a zero in this slot means. + let headroom = backend() + .map(|be| { + be.vram_budget_bytes() + .saturating_sub(be.reserved_bytes()) + .max(1) + }) + .unwrap_or(1); + let _ = RETAIN_FIRST_REFUSAL_HEADROOM.compare_exchange( + 0, + headroom, + Ordering::Relaxed, + Ordering::Relaxed, + ); + } + /// Run `f` against this codeword's tree, built here and freed on return. /// /// # Why the tree is not kept @@ -169,6 +616,31 @@ impl DeviceCodeword { /// and [`leaf_hash_calls`] make the two passes visible, and the group-scale /// test in `tests/whir_tree_cache.rs` fails if a tree is ever held past /// this call again. + /// + /// # ★ What DID work, and why it is a different object + /// + /// The tree is still not kept. Its LEAF LAYER is — see + /// [`capture_leaves`](Self::capture_leaves) — and that is not a softer + /// version of H4 but a different trade: + /// + /// - **Half the bytes.** A tree is `2·num_leaves − 1` nodes; the layer is + /// `num_leaves` of them. At fold width `k` that is `C · 2^(5−k)` bytes + /// against a tree's `C · 2^(6−k)` — at the production `k = 4`, a quarter + /// of a base codeword where H4 held half of one. + /// - **Most of the saving.** The leaf pass absorbs a whole `2^k` coset per + /// leaf: two permutations on a base codeword, six on an extension one, + /// against one per inner node. So the layer carries two thirds of a base + /// tree's work and six sevenths of an extension tree's, and rebuilding + /// the inner levels from it is the cheap third. + /// - **It can decline.** H4 could not: it allocated, and when the card said + /// no the commit fell back to the host at ~1.5 GiB a chain. The capture + /// asks [`DeviceReservation::grow`](crate::device::DeviceReservation::grow) + /// first, and a refusal costs exactly one leaf pass — the behaviour of + /// this file before the change. The cliff is unreachable rather than + /// unmeasured. + /// + /// The window is still the group's, because the window is the protocol's + /// and nothing here changes it. What changed is what sits in it. fn with_tree( &self, log_folding: usize, @@ -194,6 +666,26 @@ impl DeviceCodeword { self.builds.load(Ordering::Relaxed) } + /// ★ Leaf-hash passes over THIS codeword — the number the retention moves. + /// + /// Equal to [`tree_builds`](Self::tree_builds) when nothing is retained, and + /// 1 however many times the codeword is opened when the layer is kept. The + /// two together are what make the retention assertable in BOTH directions: + /// a cache that stopped working reads them equal, a tree kept past its call + /// reads `tree_builds` short. + pub fn leaf_passes(&self) -> u64 { + self.leaf_passes.load(Ordering::Relaxed) + } + + /// Bytes this codeword is holding as a retained leaf layer, or zero. + pub fn retained_leaf_bytes(&self) -> u64 { + self.leaves + .lock() + .ok() + .and_then(|h| h.as_ref().map(|k| k.bytes)) + .unwrap_or(0) + } + /// The root of that tree, which is the commitment. /// /// ★ The tree is KEPT (H4). The other thing anyone wants from it is a path @@ -475,7 +967,10 @@ fn commit_from( elements: n, base: true, builds: BuildCount::default(), + leaf_passes: BuildCount::default(), + leaves: Arc::new(Mutex::new(None)), room: Arc::new(room), + registered: Arc::new(AtomicBool::new(false)), }; let root = codeword.commit(log_folding, hash)?; Ok((codeword, root)) @@ -722,9 +1217,14 @@ pub fn fold_resident( // is committed and opened in its own right, and sharing the parent's // slot would make one of them evict the other every round. builds: BuildCount::default(), + leaf_passes: BuildCount::default(), + leaves: Arc::new(Mutex::new(None)), // The fold lives inside the room the codeword it came from promised: // it is half of it, and that one is still alive. room: codeword.room.clone(), + // A fold is its own codeword for retention too: its own slot, its own + // registry entry once it captures. + registered: Arc::new(AtomicBool::new(false)), }) } diff --git a/crypto/math-cuda/tests/grinding.rs b/crypto/math-cuda/tests/grinding.rs index 998c40090..f755ec9d1 100644 --- a/crypto/math-cuda/tests/grinding.rs +++ b/crypto/math-cuda/tests/grinding.rs @@ -76,3 +76,90 @@ fn gpu_grind_declines_below_min_factor() { "GPU grind should decline factor 1" ); } + +/// ★ THE BLOCK SIZE AND THE GRID CANNOT MOVE THE ANSWER, EXECUTED. +/// +/// `search` scans contiguous blocks from zero and returns the first hitting +/// block's minimum, so the nonce is a function of the inner hash and the +/// grinding factor alone. That is what makes a sweep of either knob move no +/// proof byte — and it is a property that can fail: a stride or bounds defect +/// would return a different valid nonce at a different stride, and only an +/// equality across settings can see it. +/// +/// Through `generate_nonce_gpu_at` rather than the environment because the +/// knobs cache in a `OnceLock`: one process cannot read two settings through +/// `LAMBDA_VM_GRIND_*`, and two processes would be two device contexts. +#[test] +fn the_nonce_is_the_same_at_every_scan_factor_and_grid() { + let seed = [20u8; 32]; + let factor = 20u8; + let lanes = inner_hash_lanes::(&seed, factor); + + let record = math_cuda::grinding::Knobs::DEFAULT; + let expected = math_cuda::grinding::generate_nonce_gpu_at(&lanes, factor, record) + .expect("GPU grind at the record posture (needs a GPU)"); + assert!( + is_valid_nonce::(&seed, expected, factor), + "the record posture's nonce {expected} fails is_valid_nonce" + ); + + // Both knobs, both directions, including the pair the ruling names. + for knobs in [ + math_cuda::grinding::Knobs { + scan: 1, + grid: 1024, + }, + math_cuda::grinding::Knobs { + scan: 2, + grid: 1024, + }, + math_cuda::grinding::Knobs { + scan: 8, + grid: 4096, + }, + math_cuda::grinding::Knobs { scan: 8, grid: 256 }, + math_cuda::grinding::Knobs { + scan: 1, + grid: 4096, + }, + ] { + let nonce = math_cuda::grinding::generate_nonce_gpu_at(&lanes, factor, knobs) + .expect("GPU grind (needs a GPU)"); + assert!( + is_valid_nonce::(&seed, nonce, factor), + "nonce {nonce} from {knobs:?} fails is_valid_nonce" + ); + assert_eq!( + nonce, expected, + "{knobs:?} returned {nonce}, the record posture returned {expected} — \ + the launch geometry moved the answer, so the search is not scanning \ + contiguously from zero" + ); + } +} + +/// And it is still the SMALLEST at every setting, not merely the same one. +/// +/// Factor 14 so the exhaustive host scan below the answer stays cheap. Equality +/// across settings (the test above) would be satisfied by a search that +/// consistently skipped the same range; minimality is what rules that out. +#[test] +fn the_nonce_is_still_the_smallest_at_a_narrow_grid() { + let seed = [14u8; 32]; + let factor = 14u8; + let lanes = inner_hash_lanes::(&seed, factor); + for knobs in [ + math_cuda::grinding::Knobs { scan: 1, grid: 256 }, + math_cuda::grinding::Knobs { + scan: 8, + grid: 4096, + }, + ] { + let nonce = math_cuda::grinding::generate_nonce_gpu_at(&lanes, factor, knobs) + .expect("GPU grind (needs a GPU)"); + assert!( + (0..nonce).all(|n| !is_valid_nonce::(&seed, n, factor)), + "nonce {nonce} from {knobs:?} is not the smallest valid nonce" + ); + } +} diff --git a/crypto/math-cuda/tests/host_kat/rpx_host_kat.cpp b/crypto/math-cuda/tests/host_kat/rpx_host_kat.cpp index 7cdf8cbd5..5f69ae403 100644 --- a/crypto/math-cuda/tests/host_kat/rpx_host_kat.cpp +++ b/crypto/math-cuda/tests/host_kat/rpx_host_kat.cpp @@ -1097,7 +1097,31 @@ void permute_probe_matches_the_oracle_table() { // --------------------------------------------------------------------------- uint64_t run_grind(const uint64_t inner[4], uint8_t factor, uint64_t base, uint64_t count) { const uint64_t limit = (uint64_t)1 << (64 - factor); - uint64_t result = UINT64_MAX; + // ⛔⛔ `unsigned long long`, NOT `uint64_t`, AND THAT DIFFERENCE WAS THE BUG. + // The kernel takes `volatile unsigned long long *` because that is the type + // CUDA's `atomicMin` overload wants, so this object's address is handed out + // under it. `uint64_t` is `unsigned long long` on Darwin/arm64 and + // `unsigned long` on LP64 glibc — a DIFFERENT type of the same width — so on + // Linux the cast type-punned, and with `#include "rpx.cu"` putting the whole + // kernel in this translation unit, GCC 13.3 at -O2 was free under TBAA to + // assume a write through `unsigned long long *` could not touch an + // `unsigned long`, and to keep `result` in a register across the inlined call. + // + // It did. `run_grind` returned `UINT64_MAX` for every input on the box, so + // every check below that expects the SENTINEL passed vacuously while every + // check that expects a FOUND nonce failed — six rows, at every sha back to + // the gated base `c00342c1f`, while the same source passed on a clang/arm64 + // laptop where the two types coincide. Measured on the box 2026-09-20: + // `-O2` → 6 failures; `-O2 -fno-strict-aliasing` → all pass; `-O0` → all pass. + // + // ⚠ NONE OF THAT WAS EVER A STATEMENT ABOUT THE DEVICE GRIND. This file is a + // HOST replay of the kernel source through `cuda_host_shim.h`; the defect + // was in the harness holding the result, not in the kernel it was testing. + // ⛔ Do not "tidy" this back to `uint64_t`: matching the pointer type the + // kernel is given is what makes the access well-defined, and the box gate + // carries a mutation that restores `uint64_t` and requires those six rows + // back — so the tidy-up would be caught, loudly, by a red nobody wants again. + unsigned long long result = UINT64_MAX; CUDA_HOST_SINGLE_THREAD(); rpx_grind_search(inner, limit, base, count, (volatile unsigned long long *)&result); return result; diff --git a/crypto/math-cuda/tests/whir_tree_cache.rs b/crypto/math-cuda/tests/whir_tree_cache.rs index e3e001c27..a6494ae49 100644 --- a/crypto/math-cuda/tests/whir_tree_cache.rs +++ b/crypto/math-cuda/tests/whir_tree_cache.rs @@ -1,5 +1,17 @@ -//! ★★ H4's result of record — a commitment does NOT keep its tree, and the -//! bytes it holds are only its codeword. +//! ★★ H4's result of record, and what replaced it — a commitment does NOT keep +//! its TREE; it keeps its LEAF LAYER, and the bytes it holds are its codeword +//! and that layer. +//! +//! ⛔ H4's finding stands and is not edited away below: keeping the whole node +//! array LOST, measured on the card at about +15 s, because the retention is +//! one object per commitment IN THE GROUP and ten of those put the device at +//! 96%. What changed is WHICH object. A tree is `2·num_leaves − 1` nodes; its +//! leaf layer is `num_leaves` of them — half the bytes — and the leaf pass it +//! saves is two thirds of a base tree's permutations and six sevenths of an +//! extension one, because a leaf absorbs a whole `2^k` coset while an inner +//! node absorbs two digests. Half the memory for most of the saving is a +//! different trade from the one H4 measured, and these tests now pin BOTH +//! sides of it: the layer must be held, and a whole tree must still never be. //! //! Needs a GPU: //! @@ -58,7 +70,7 @@ use math::field::element::FieldElement; use math::field::goldilocks::GoldilocksField as F; use math_cuda::DeviceHash; -use math_cuda::whir::{leaf_hash_calls, reset_leaf_hash_calls}; +use math_cuda::whir::{leaf_hash_calls, retention_report, tree_builds}; use multilinear::mle::Mle; use multilinear::whir::{self, Domain}; use multilinear::whir_commit::{CodewordCommitment, verify_opening}; @@ -140,12 +152,27 @@ fn a_commitment_hashes_its_leaves_once_per_tree_it_builds() { "{name}: the commit itself must hash the leaves exactly once" ); + assert_eq!( + codeword.leaf_passes(), + 1, + "{name}: the commit hashes the leaves once" + ); + let _ = codeword.paths(4, &[0, 1, 7], hash).expect("paths"); assert_eq!( codeword.tree_builds(), 2, "{name}: an opening builds its own tree — a 1 here means one is kept" ); + // ★ THE OTHER DIRECTION, and it is the whole point of the change: the + // tree was rebuilt, but its LEAF LAYER was not re-hashed. A 2 here is + // the retention not working. + assert_eq!( + codeword.leaf_passes(), + 1, + "{name}: the opening must serve the retained leaf layer — a 2 here \ + means the layer was not kept, or not matched" + ); // …and again, because a cache that served once and then evicted would // read 2 on the line above too. @@ -155,6 +182,17 @@ fn a_commitment_hashes_its_leaves_once_per_tree_it_builds() { 3, "{name}: and a second opening builds a third" ); + assert_eq!( + codeword.leaf_passes(), + 1, + "{name}: and still one leaf pass — a layer that served once and was \ + then evicted would read 2 here" + ); + assert!( + codeword.retained_leaf_bytes() > 0, + "{name}: the codeword reports no retained layer, so the counts above \ + are agreeing about the wrong thing" + ); } } @@ -355,29 +393,130 @@ fn a_group_holds_only_its_codewords_before_any_open() { // sample. math_cuda::device::drain_and_trim().expect("drain"); let free_before = be.free_vram_bytes().expect("cuMemGetInfo"); + // ★ THE SECOND INSTRUMENT, and it is the one that can tell the two stories + // apart. `free_vram_bytes` is the DRIVER's count and includes whatever the + // pool is sitting on; `reserved_bytes` is what the CODE promised and is + // blind to the pool by construction. Their difference is the pool's, and + // printing it turns "either the pool retained a block or the code holds one + // more" from a question into a read. + let reserved_before = be.reserved_bytes(); let held: Vec<_> = (0..4) .map(|_| commit_on_device(num_vars, log_folding, hash)) .collect(); + // ⛔ SYMMETRIC SAMPLING, AND THIS LINE IS THE FIX. `free_before` is taken + // AFTER a drain and this one was not, so the difference measured the code + // plus every transient the four commits made — and with the pool set to + // retain all freed blocks, that is all of them. A delta between two samples + // is about the code only if both are taken at the same pool state. + // + // ★ WHAT IT COST, and why a widened bound was the wrong repair. The run of + // 2026-09-20 read `taken` = 301,989,888 B — EXACTLY nine codewords, on a + // bound of nine, where the model says eight are held. The old one-sided + // bound was `8 × codeword` against four codewords held, so it carried 128 + // MiB of margin the pool had been living in unnoticed; the leaf-layer + // retention did not add pool retention, it CONSUMED that margin. And the + // slack was never sized against the transients anyway: the node buffer + // `build_tree` allocates is `(2L−1)·32` = 64 MiB, twice the bound's 32. + // + // The same run's mutation arm settles which it is. Holding a whole node + // array instead of a layer moved the measurement to 480 MiB where the code + // then holds 384 — an excess of 96 against the honest run's 32, tripling + // while the holding grew by half. No "the code holds one more object" form + // fits both (`4×(cw+tree)+tree` = 448, `+cw` = 416; `5×(cw+tree)` = 480 fits + // MUT C exactly and dies on the honest run, where `5×(cw+leaf)` = 320 ≠ 288). + // Every peak-demand model misses on BOTH sides (320 and 448 predicted), which + // is the signature of driver suballocation and not of anything this tree + // accounts for. ⇒ `free_vram_bytes()` cannot carry a bound this tight + // unless both samples are drained. + math_cuda::device::drain_and_trim().expect("drain"); let free_after = be.free_vram_bytes().expect("cuMemGetInfo"); let taken = free_before.saturating_sub(free_after); + let promised = be.reserved_bytes().saturating_sub(reserved_before); let codeword_bytes = ((1u64 << num_vars) << 2) * 8; let leaves = ((1u64 << num_vars) << 2) >> log_folding; + let leaf_bytes = leaves * 32; let tree_bytes = (2 * leaves - 1) * 32; - let bound = 8 * codeword_bytes; + // ⛔ TWO-SIDED, AND AGAINST THE FORM RATHER THAN A MULTIPLE. The layers must + // be HELD (so more than the codewords alone) and a whole TREE must still + // never be (so less than four of those). At this shape — `log_folding = 2`, + // chosen by the comment above because it makes a tree two codewords — a + // leaf layer is exactly ONE codeword, so the three cases are 128, 256 and + // 384 MiB and the bound sits between the last two with one codeword of + // slack. A one-sided bound passed either way and is what let the old + // arithmetic sit on its own edge. + let expect = 4 * (codeword_bytes + leaf_bytes); + let bound = expect + codeword_bytes; + let floor = 4 * codeword_bytes; let mib = |b: u64| b / (1 << 20); + // ★ THE TWO ACCOUNTINGS, SIDE BY SIDE, IN WHICHEVER MESSAGE FIRES. A failure + // here used to say only how many bytes the DRIVER lost, which cannot + // distinguish "the pool retained a block" from "the code holds one more" — + // and those call for opposite repairs. `promised` is the code's own number + // and is blind to the pool; the per-codeword pair says how many layers are + // actually in hand. `driver − promised` is the pool's share, and it should + // now be small: the drain above is what makes that true. + let ledger = { + let per: Vec = held + .iter() + .map(|(c, _)| { + format!( + "{}/{}", + mib(c.reserved_bytes()), + mib(c.retained_leaf_bytes()) + ) + }) + .collect(); + format!( + "driver {} MiB · promised {} MiB · pool share {} MiB · per codeword \ + reserved/retained MiB: [{}]", + mib(taken), + mib(promised), + mib(taken.saturating_sub(promised)), + per.join(", ") + ) + }; + // ⛔ UNCONDITIONAL, AND THAT IS THE WHOLE POINT. Built only inside the + // assertion messages, this line appears ONLY when the test fails — so on + // the honest path the numbers were inferred from a pass and never read. + // + // What a pass alone establishes is `driver < bound`, i.e. pool share under + // one codeword, and NOTHING narrower. The mutation run that keeps a whole + // node array reads `pool share 64 MiB` even after the symmetric drain — a + // fragmentation floor of one largest transient, because a best-effort + // `trim` cannot release a chunk still backing a live allocation. If the + // honest path sits anywhere near that, this guard has a margin of a few MiB + // and will flake, and nobody would learn it from a green run. + // + // ⇒ printed every time, under `--nocapture`, which is how the box gate runs + // this suite. The number becomes a READ, and a ceiling can be asserted + // against it once its honest value is known. + println!(" group guard: {ledger}"); assert!( taken < bound, "four unopened commitments took {} MiB from the device. Four codewords \ - are {} MiB and the bound is {} MiB; a tree is {} MiB, so four of those \ - kept would read {} MiB. Something is held per commitment.", + and their leaf layers are {} MiB and the bound is {} MiB; a TREE is {} \ + MiB, so four of those kept would read {} MiB. Something bigger than a \ + leaf layer is held per commitment.\n {}", mib(taken), - mib(4 * codeword_bytes), + mib(expect), mib(bound), mib(tree_bytes), mib(4 * (codeword_bytes + tree_bytes)), + ledger, + ); + assert!( + taken > floor, + "four unopened commitments took only {} MiB, which is at or under the {} \ + MiB their codewords alone need. The leaf layers ({} MiB for four) are \ + NOT being held — either the capture never ran or the budget refused it, \ + and in both cases every opening will re-hash its leaves.\n {}", + mib(taken), + mib(floor), + mib(4 * leaf_bytes), + ledger, ); // The commitments are alive up to here, which is the whole point: a `drop` @@ -406,11 +545,29 @@ fn a_tree_is_built_for_the_blocking_that_is_asked_for() { 2, "the opening must build a tree for the blocking it was given" ); + // ★ THE KEY, ASSERTED WHERE IT CAN FAIL. The retained layer was built at + // k=4; this opening is at k=2 and describes a DIFFERENT tree, so the layer + // must not be served and the leaves must be hashed again. A 1 here is a + // cache ignoring its key, which is the one way this change could hand back + // paths that are internally consistent and wrong. + assert_eq!( + codeword.leaf_passes(), + 2, + "a k=2 opening must NOT be served the k=4 leaf layer" + ); // And the rebuild answered the question that was asked: at k=2 the tree has // four times the leaves, so each path is two levels deeper. let at_four = codeword.paths(4, &[0, 1], hash).expect("paths at k=4"); assert_eq!(codeword.tree_builds(), 3, "and a third for the k=4 opening"); + // …and the k=4 layer IS still there and IS served, so the key rejects a + // mismatch without throwing away a match. Without this line the test above + // would also pass on a cache that had simply stopped working. + assert_eq!( + codeword.leaf_passes(), + 2, + "the k=4 opening matches the retained layer's key and must not re-hash" + ); assert_eq!( at_two.len(), at_four.len() + 2 * 2 * 32, @@ -444,20 +601,210 @@ fn the_process_wide_counter_tracks_the_same_passes() { let _exclusive = exclusive(); let hash = key::(); - reset_leaf_hash_calls(); + // ⛔ DELTAS, NOT ABSOLUTES. These three counters are process-wide and no + // longer resettable as a set, and the retention makes them diverge on + // purpose, so the assertions below are about what THIS codeword moved. + let at = || (tree_builds(), leaf_hash_calls(), retention_report().5); + + let (b0, p0, s0) = at(); let (codeword, _root) = commit_on_device(12, 4, hash); - let after_commit = leaf_hash_calls(); - assert_eq!(after_commit, 1, "one commit, one leaf-hash pass"); + let (b1, p1, s1) = at(); + assert_eq!(b1 - b0, 1, "one commit, one tree assembled"); + assert_eq!(p1 - p0, 1, "and it paid for its own leaf pass"); + assert_eq!(s1 - s0, 0, "with nothing yet in hand to reuse"); let _ = codeword.paths(4, &[0, 1], hash).expect("paths"); + let (b2, p2, s2) = at(); + assert_eq!(b2 - b1, 1, "the opening assembles its own tree"); + // ★ THE LINE THAT CHANGED WITH H4's REPLACEMENT. This used to assert the + // global counter moved by one too, because a tree and a leaf pass were the + // same event. They are not any more: the tree is assembled, the leaves are + // not re-hashed, and a 1 here is the retention failing to serve. + assert_eq!(p2 - p1, 0, "and does NOT re-hash the leaves"); + assert_eq!(s2 - s1, 1, "the saving is counted where it happens"); + + // ⭐ THE ACCOUNTING IDENTITY, which is what the old equality became and + // which fails in BOTH directions: a tree that skipped its pass without + // recording a saving breaks it, and so does a saving recorded for a tree + // that was never assembled. assert_eq!( - leaf_hash_calls(), - after_commit + 1, - "an opening builds a tree, so the global counter must move by one" + b2 - b0, + (p2 - p0) + (s2 - s0), + "every tree either paid for its leaf pass or reused one; trees {}, \ + passes {}, savings {}", + b2 - b0, + p2 - p0, + s2 - s0 ); +} + +/// ⛔ THE ARGUE-SURFACE DEVICE-FALLBACK COUNTER FIRES AT A REAL SITE. +/// +/// wt16 read as a win at the slot level because argue's per-table device work +/// fell back to the host UNCOUNTED — `multilinear::gpu::host_fallbacks()` +/// counts the COMMIT path only. `math_cuda::device::device_fallbacks()` is the +/// counter that closes that blind spot; this test proves it actually moves when +/// an argue-surface `reserve` is refused, using the cheapest of the five sites, +/// `DeviceColumns::upload`. +/// +/// It forces the refusal WITHOUT a budget setter and WITHOUT allocating any +/// device memory: `Backend::reserve` is a pure atomic bump on the reservation +/// total (no `cuMemAlloc`), so reserving the whole remaining budget makes every +/// later `reserve` return `None` at no memory cost. The reservation is dropped +/// at the end, giving the budget back. +/// +/// The gate mutation is deleting the `note_device_fallback()` call at +/// `columns.rs`'s `reserve`→`None` site: the count then reads 0 and the final +/// assertion reddens by name. The four undriven sites are covered card-free by +/// `note_device_fallback_is_called_at_exactly_the_five_argue_sites` in +/// `device.rs`. +#[test] +fn an_argue_reservation_refusal_bumps_the_device_fallback_counter() { + let _exclusive = exclusive(); + let be = math_cuda::device::backend().expect("device fallback test needs a GPU"); + + // Take the whole remaining budget as one reservation — an atomic bump, no + // device memory — so any further `reserve` must be refused. Held to the end + // of the test, then dropped. + let remaining = be.vram_budget_bytes().saturating_sub(be.reserved_bytes()); + let _hog = math_cuda::device::reserve(remaining) + .expect("reserving the remaining budget is an accounting move and cannot fail"); assert_eq!( - codeword.tree_builds(), - leaf_hash_calls(), - "with one codeword in flight the two counters must agree" + be.reserved_bytes(), + be.vram_budget_bytes(), + "the budget is now fully promised, so the next reserve must be refused" + ); + + math_cuda::device::reset_device_fallbacks(); + assert_eq!( + math_cuda::device::device_fallbacks(), + 0, + "the counter starts this measurement at zero" ); + + // The cheapest argue site: one column of one element wants 8 bytes the + // budget cannot promise, so `upload` returns `None` at its `reserve` and + // the site records the fallback. + let refused = math_cuda::columns::DeviceColumns::upload(&[&[0u64]]); + assert!( + refused.is_none(), + "with the budget fully promised, the device upload must decline" + ); + assert_eq!( + math_cuda::device::device_fallbacks(), + 1, + "an argue-surface reserve was refused, so the device-fallback counter \ + must read exactly one — a 0 here is the counter not wired to the site" + ); + + // `_hog` is dropped here at scope end, returning the reserved budget. +} + +/// ⛔ THE LIVE FOOTPRINT AND RESERVED HIGH-WATER TRACK THE RETENTION. +/// +/// The two instruments the evictable retention is built on: `retained_bytes_live` +/// (the SIMULTANEOUS footprint, not the cumulative `held`) and `reserved_high_water` +/// (the peak `be.reserved`, the quantity argue's `reserve` is checked against). +/// This proves both move with a real retention and, crucially, that the layer's +/// bytes are GIVEN BACK when its codeword drops — the balance the eviction relies +/// on. A broken `RetainedLeaves::drop` leaves the live count high and reddens the +/// final assertion; a missing `note_reserved` leaves the high-water below the +/// live reserved total and reddens the middle one. +#[test] +fn the_live_footprint_and_reserved_high_water_track_the_retention() { + let _exclusive = exclusive(); + let be = math_cuda::device::backend().expect("footprint test needs a GPU"); + let live_before = math_cuda::whir::retained_bytes_live(); + { + let (codeword, _root) = commit_on_device(14, 4, key::()); + let _ = codeword + .paths(4, &[0, 1, 7], key::()) + .expect("paths"); + // A leaf layer is captured and held on the codeword. + let live = math_cuda::whir::retained_bytes_live(); + assert!( + live > live_before, + "a held leaf layer must raise the live footprint (live {live}, before {live_before})" + ); + assert!( + math_cuda::whir::retained_bytes_peak() >= live, + "the peak footprint must be at least the current live count" + ); + assert!( + math_cuda::device::reserved_high_water() >= be.reserved_bytes(), + "the reservation high-water must be at least the current reserved total \ + — note_reserved must fire on every rise of be.reserved" + ); + } + // The codeword — and, synchronously in its Drop, the retained layer — is gone. + assert_eq!( + math_cuda::whir::retained_bytes_live(), + live_before, + "dropping the codeword must give the layer's bytes back: RetainedLeaves::drop \ + balances the admit, or the live footprint would only ever rise" + ); +} + +/// ⛔ A BUDGET MISS EVICTS A RETAINED LAYER AND THE RESERVE THEN SUCCEEDS. +/// +/// The whole point of the evictable retention: when a real caller (argue, here a +/// bare reserve) cannot get its bytes, the retention gives a layer back rather +/// than the caller falling to the host. Committing captures ONE base layer (no +/// folds), so exactly one layer is in the registry. Then the budget is filled to +/// leave LESS free than the layer's bytes, so a reserve of the layer's size must +/// miss — and it SUCCEEDS only because the evictor reclaims the layer. The +/// mutation that disables the evictor (skip the consult in reserve, or never +/// install it) makes this reserve return None and the test panic on the expect. +#[test] +fn a_budget_miss_evicts_a_retained_layer_and_the_reserve_succeeds() { + let _exclusive = exclusive(); + let be = math_cuda::device::backend().expect("eviction test needs a GPU"); + + // Commit captures the base leaf layer; hold the codeword so the layer stays. + let (codeword, _root) = commit_on_device(14, 4, key::()); + let layer_bytes = codeword.retained_leaf_bytes(); + assert!( + layer_bytes > 0, + "precondition: the commit must have retained a layer" + ); + let (evictions_before, _) = math_cuda::whir::retention_evictions(); + + // Fill the budget to leave a gap SMALLER than the layer, so a reserve of the + // layer's size cannot fit without eviction. reserve is a pure atomic bump, + // so the hog costs no device memory. + let gap = layer_bytes / 2; + let hog_bytes = be + .vram_budget_bytes() + .saturating_sub(be.reserved_bytes()) + .saturating_sub(gap); + let _hog = math_cuda::device::reserve(hog_bytes).expect("the hog reservation cannot fail"); + assert!( + be.vram_budget_bytes().saturating_sub(be.reserved_bytes()) < layer_bytes, + "the free budget must now be below the layer size, so the next reserve misses" + ); + + // This would return None without the evictor; with it, the layer is freed + // and the reserve succeeds. + let got = math_cuda::device::reserve(layer_bytes).expect( + "the reserve must SUCCEED by evicting the retained layer — a None here \ + is the evictor not consulted, or eviction freeing nothing", + ); + + let (evictions_after, bytes_evicted) = math_cuda::whir::retention_evictions(); + assert!( + evictions_after > evictions_before, + "an eviction must have been recorded ({evictions_before} -> {evictions_after})" + ); + assert!( + bytes_evicted >= layer_bytes, + "the eviction must have freed at least the layer's bytes" + ); + assert_eq!( + codeword.retained_leaf_bytes(), + 0, + "the evicted codeword's layer slot must read None (the sole registered layer)" + ); + + drop(got); + drop(_hog); } diff --git a/crypto/multilinear/src/stacked_eval.rs b/crypto/multilinear/src/stacked_eval.rs index 8ba8ddc7d..9f27e8b0a 100644 --- a/crypto/multilinear/src/stacked_eval.rs +++ b/crypto/multilinear/src/stacked_eval.rs @@ -107,6 +107,22 @@ where // sixteen, and what the difference buys is the widest tables getting a // device at all. If the card will not promise it, each commitment // promises its own, which is the conservative accounting. + // + // ★ AND THE LEAF LAYERS ARE NOW IN THIS NUMBER, as a named term rather + // than as a surprise. A commitment that is opened keeps the leaf layer + // its commit hashed — `num_leaves × 32` = `C · 2^(5−k)` bytes at fold + // width `k`, a QUARTER of a base codeword at the production `k = 4` — + // so a group of `n` commitments holds `n · C/4` more than the codewords + // alone. It is not reserved here, because the width is not known until + // the tree is built: each codeword grows THIS reservation when it + // captures (`DeviceCodeword::capture_leaves`), and `grow` refuses + // without changing anything when the budget will not take it. A refused + // retention costs the leaf pass again and nothing else. + // + // ⚠ A TREE is still not in this number and must never be. H4 kept the + // whole node array — twice these bytes — and the card reached 96%, after + // which commits fell back to the host at ~1.5 GiB each. The layer is + // half of what that held and two thirds of what it saved. let room = sources.first().and_then(|poly| { let codeword_bytes = (1u64 << (poly.num_vars() + config.log_blowup)) * 8; crate::gpu::reserve_room(codeword_bytes) diff --git a/crypto/multilinear/src/whir_chain.rs b/crypto/multilinear/src/whir_chain.rs index 8d49a1a23..d614b92e1 100644 --- a/crypto/multilinear/src/whir_chain.rs +++ b/crypto/multilinear/src/whir_chain.rs @@ -704,6 +704,9 @@ where T: IsTranscript, H: WhirHash, { + // One chain, counted — arm F's identity is per chain, and a group can open + // several (`stacked_eval::prove` runs one per commitment). + crate::whir_split::bump(&crate::whir_split::CHAIN_COUNT); let schedule = config.schedule(num_vars); // The codeword comes out of the commitment rather than being encoded // again: it is the same array, and the NTT is not cheap. @@ -721,6 +724,9 @@ where // after the last round — is reported as `setup_tail`, named rather // than left as a gap for a tolerance to swallow. let __wc_round = crate::whir_split::mark(); + // One per iteration, so arm F can derive the `open_many` calls this + // loop should have made (`2R − 1` per chain) from the run's own count. + crate::whir_split::bump(&crate::whir_split::ROUND_COUNT); // ── the round's six slots, under `LAMBDA_VM_BASE_SPLIT=1` ── // They partition the round, so `open_groups - Σ(six)` is loop overhead // and nothing else. All three grinds share one slot: they are the same @@ -929,9 +935,13 @@ where T: IsTranscript, H: WhirHash, { + let __wq_sample = crate::whir_split::mark(); let queries: Vec = (0..config.num_queries) .map(|_| transcript.sample_u64(current.num_leaves() as u64) as usize) .collect(); + crate::whir_split::add(&crate::whir_split::QUERY_SAMPLE, __wq_sample); + // ⛔ ONE `open_many` here, not two: the final round has no successor to + // open. That is the `− 1` in arm F's `2R − 1`. Ok(RoundProof { current: current.open_many(&queries)?, next: Vec::new(), diff --git a/crypto/multilinear/src/whir_commit.rs b/crypto/multilinear/src/whir_commit.rs index f8aed99aa..b1a72e7fa 100644 --- a/crypto/multilinear/src/whir_commit.rs +++ b/crypto/multilinear/src/whir_commit.rs @@ -308,9 +308,22 @@ where /// them. pub fn open_many(&self, indices: &[usize]) -> Result>, Error> { let num_leaves = self.num_leaves(); + // ★ ONE CALL, ONE DEVICE TREE REBUILD. Counted rather than inferred: + // `whir_round::prove` opens the current commitment AND its successor, + // so a non-final round passes here twice, and `check_closure`'s arm F + // asserts the total against the rounds that produced it. + crate::whir_split::bump(&crate::whir_split::REBUILD_CALLS); + // ⛔ THE SPLIT IS HERE AND NOT INSIDE `paths()`. On the device arm + // `paths()` is a range check, ONE device call and a `map` into `Proof`, + // so splitting inside it would weigh the rebuild against host + // bookkeeping and read ~100% every time. What competes with the rebuild + // is the COSET GATHER, which is the next statement, not a nested one. + let __wq_tree = crate::whir_split::mark(); let proofs = self.paths(indices)?; + crate::whir_split::add(&crate::whir_split::TREE_REBUILD, __wq_tree); let block = 1usize << self.log_folding; + let __wq_gather = crate::whir_split::mark(); let blocks: Vec>> = match &self.codeword { Codeword::Host(values) => indices .iter() @@ -331,11 +344,16 @@ where } }; - Ok(blocks + crate::whir_split::add(&crate::whir_split::COSET_GATHER, __wq_gather); + + let __wq_assemble = crate::whir_split::mark(); + let openings: Vec> = blocks .into_iter() .zip(proofs) .map(|(values, proof)| CosetOpening { values, proof }) - .collect()) + .collect(); + crate::whir_split::add(&crate::whir_split::OPEN_ASSEMBLE, __wq_assemble); + Ok(openings) } /// One authentication path per index, from wherever the tree is. diff --git a/crypto/multilinear/src/whir_round.rs b/crypto/multilinear/src/whir_round.rs index 3cffbcd93..66ea3c42a 100644 --- a/crypto/multilinear/src/whir_round.rs +++ b/crypto/multilinear/src/whir_round.rs @@ -93,12 +93,17 @@ where T: IsTranscript, H: WhirHash, { + // The transcript squeezes that choose the indices, and the map onto the + // successor's leaves. Host work, and the only part of a round's openings + // that is not `open_many`. + let __wq_sample = crate::whir_split::mark(); let queries = sample_queries(transcript, config.num_queries, current.num_leaves()); let leaves: Vec = queries .iter() .map(|q| leaf_and_slot(*q, next.num_leaves()).0) .collect(); + crate::whir_split::add(&crate::whir_split::QUERY_SAMPLE, __wq_sample); Ok(RoundProof { current: current.open_many(&queries)?, diff --git a/crypto/multilinear/src/whir_split.rs b/crypto/multilinear/src/whir_split.rs index 256c9aa49..17f290974 100644 --- a/crypto/multilinear/src/whir_split.rs +++ b/crypto/multilinear/src/whir_split.rs @@ -146,6 +146,35 @@ impl Slot { } } +/// An occurrence counter. Same discipline as [`Slot`] — process-global, and +/// cleared when read — but it counts CALLS, not nanoseconds. +/// +/// ★ It exists because a count survives a shape where a duration does not. On +/// the card-free fixture the query openings read `0.00` s, so no bound on a +/// TIME could ever see one of the four windows go missing there; the calls +/// still happen, so a count still can. That is what makes the gate's +/// real-wiring mutation able to fire at the shape the gate actually runs. +#[derive(Debug)] +pub struct Counter(AtomicU64); + +impl Counter { + const fn new() -> Self { + Self(AtomicU64::new(0)) + } + /// Read and CLEAR, for the reason [`Slot::take`] does. + fn take(&self) -> u64 { + self.0.swap(0, Ordering::Relaxed) + } +} + +/// Count one occurrence. Inert when the instrument is off, like [`mark`]. +#[inline] +pub fn bump(counter: &Counter) { + if enabled() { + counter.0.fetch_add(1, Ordering::Relaxed); + } +} + /// `absorb_roots_and_challenge`: the roots into the transcript and the /// challenge out. Host. pub static CHALLENGE: Slot = Slot::new(); @@ -186,6 +215,33 @@ pub static OOD: Slot = Slot::new(); /// the Merkle tree on device per query batch. pub static QUERIES: Slot = Slot::new(); +// ── inside the query openings: what `open_many` does, twice per round ─────── +// +// QUERIES is the largest slot in the chain and, like `open_groups` before it, +// one number. These four PARTITION it, and they are counted on EVERY +// `open_many` call — `whir_round::prove` opens the CURRENT commitment and the +// NEXT one (`whir_round.rs:104-105`), so a non-final round calls it twice and +// the final round once through `final_openings`. +// +// ⛔ THE BOUNDARY IS `open_many`, NOT `paths()`. On the device arm `paths()` is +// a range check, ONE device call, and a `map` that wraps each path in a +// `Proof`; splitting INSIDE it would put the rebuild against host bookkeeping +// over a hundred kilobyte-sized paths and read ~100% every time. The term that +// competes with the rebuild is the COSET GATHER, which sits beside `paths()` in +// `open_many` and would otherwise stay inside QUERIES as an unnamed remainder — +// the same shape as the setup gap the seventh slot was added to name. + +/// `sample_queries` / the `sample_u64` loop — the transcript squeezes that +/// choose the indices, and the `leaf_and_slot` map onto the successor. +pub static QUERY_SAMPLE: Slot = Slot::new(); +/// ★ `paths()` — where a device codeword's Merkle tree is REBUILT, because the +/// commitment kept only its root. This is round 3's whole question. +pub static TREE_REBUILD: Slot = Slot::new(); +/// `cosets()` — the query blocks gathered off the codeword where it lies. +pub static COSET_GATHER: Slot = Slot::new(); +/// The zip/map/collect that pairs each block with its path. +pub static OPEN_ASSEMBLE: Slot = Slot::new(); + /// The ROUND LOOP's own wall, summed over rounds — the six measured against /// their container rather than against `open_groups` directly. /// @@ -207,11 +263,56 @@ pub static QUERIES: Slot = Slot::new(); /// of the two owns the gap; neither is inferred. pub static ROUND: Slot = Slot::new(); -/// The six, as taken at the group-loop boundary, awaiting the record. -static CHAIN_AT_GROUPS: Mutex<([f64; 6], f64)> = Mutex::new(([0.0; 6], 0.0)); +/// Rounds the chain ran, summed over the group openings. +pub static ROUND_COUNT: Counter = Counter::new(); + +/// ★ CHAINS run in the group openings — NOT groups. +/// +/// ⛔ A group is not a chain. `stacked_eval::prove` runs one chain per +/// COMMITMENT in the stacked commitment (`stacked_eval.rs:367`, a loop over +/// `stacked.commitments`), so a single group can open several. Arm F's identity +/// is per chain, so it counts chains; deriving it from `groups` would have made +/// the arm red on the honest path the first time a group carried two +/// commitments — and a check that reddens honestly gets widened until it cannot +/// fail at all. +pub static CHAIN_COUNT: Counter = Counter::new(); + +/// ★ `open_many` calls — and therefore device tree REBUILDS, one per call. +/// +/// Pre-registered as an identity the record can check itself against: a chain +/// of `R` rounds calls `open_many` twice per non-final round and once in the +/// final one, so `2R − 1`. Summed over `g` chains that is +/// `2·Σ R − g = 2·round_count − groups`, which is [`check_closure`]'s arm F. +pub static REBUILD_CALLS: Counter = Counter::new(); + +/// Everything the chain parked at the group-loop boundary, awaiting the record. +#[derive(Clone, Copy, Debug, Default, PartialEq)] +pub struct ChainSlots { + /// grind, sumcheck, fold, commit_folded, ood, queries. + pub six: [f64; 6], + /// The round loop's own wall, summed over rounds. + pub round_wall: f64, + /// query_sample, tree_rebuild, coset_gather, open_assemble. + pub queries: [f64; 4], + /// Chains run, rounds run, and `open_many` calls made, in the group + /// openings. + pub chain_count: u64, + pub round_count: u64, + pub rebuild_calls: u64, +} -/// Park the six and the round wall for the record being built one layer up. -pub fn note_chain(chain: ([f64; 6], f64)) { +/// The chain's slots, as taken at the group-loop boundary. +static CHAIN_AT_GROUPS: Mutex = Mutex::new(ChainSlots { + six: [0.0; 6], + round_wall: 0.0, + queries: [0.0; 4], + chain_count: 0, + round_count: 0, + rebuild_calls: 0, +}); + +/// Park the chain's slots for the record being built one layer up. +pub fn note_chain(chain: ChainSlots) { if !enabled() { return; } @@ -220,10 +321,10 @@ pub fn note_chain(chain: ([f64; 6], f64)) { } } -/// Read and clear the six chain slots (in record order) and the round wall. -pub fn take_chain() -> ([f64; 6], f64) { - ( - [ +/// Read and clear every chain slot and counter, in record order. +pub fn take_chain() -> ChainSlots { + ChainSlots { + six: [ GRIND.take(), SUMCHECK.take(), FOLD.take(), @@ -231,8 +332,17 @@ pub fn take_chain() -> ([f64; 6], f64) { OOD.take(), QUERIES.take(), ], - ROUND.take(), - ) + round_wall: ROUND.take(), + queries: [ + QUERY_SAMPLE.take(), + TREE_REBUILD.take(), + COSET_GATHER.take(), + OPEN_ASSEMBLE.take(), + ], + chain_count: CHAIN_COUNT.take(), + round_count: ROUND_COUNT.take(), + rebuild_calls: REBUILD_CALLS.take(), + } } /// Close a region opened by [`mark`] into `slot`, returning its seconds. @@ -390,6 +500,14 @@ pub struct ProverSplit { /// The round loop's own wall, summed over rounds. The six live inside it; /// everything else in `open_groups` lives outside it. pub chain_round_wall: f64, + /// The four inside the QUERIES slot, in the order [`take_chain`] returns + /// them: query_sample, tree_rebuild, coset_gather, open_assemble. + pub queries: [f64; 4], + /// Chains and rounds the group openings ran, and `open_many` calls they + /// made — the three numbers arm F's identity is written over. + pub chain_count: u64, + pub round_count: u64, + pub rebuild_calls: u64, } /// The six chain slots' names, in record order — so a message can name the one @@ -403,6 +521,29 @@ pub const CHAIN_NAMES: [&str; 6] = [ "queries", ]; +/// The four query slots' names, in record order — same reason as +/// [`CHAIN_NAMES`]: a message names the slot that went missing. +pub const QUERY_NAMES: [&str; 4] = [ + "query_sample", + "tree_rebuild", + "coset_gather", + "open_assemble", +]; + +/// ★ The ABSOLUTE allowance arm E gives the four query slots, beside its +/// relative one. +/// +/// A purely relative bound cannot work at both shapes the gate runs. Card-free, +/// the fixture's query openings read `0.00` s — a couple of milliseconds — and +/// 3% of that is tens of microseconds, below the glue between the four windows +/// and below the clock's own resolution, so the bound would fire on the honest +/// path. At the block, QUERIES is about a second per epoch and this slack is +/// 0.2%, so any of the four going missing still trips it. +/// +/// ⇒ At the fixture the bound is inert BY DESIGN, and that is exactly why the +/// fixture's real-wiring mutation targets arm F's COUNT rather than a duration. +const QUERY_SLACK: f64 = 0.002; + /// The index the cross-epoch global stage records under. It is the last thing /// the base does and it is INSIDE the base's wall, so it belongs in the table — /// but it is not an epoch and must not be averaged with them. @@ -434,6 +575,22 @@ impl ProverSplit { pub fn chain_other(&self) -> f64 { self.open_groups - self.chain.iter().sum::() } + /// What the four leave over inside the QUERIES slot: the glue between the + /// windows `open_many` opens, plus `whir_round::prove`'s own frame. + pub fn queries_other(&self) -> f64 { + self.chain[5] - self.queries.iter().sum::() + } + /// ★ The `open_many` calls the group openings SHOULD have made, from the + /// chains and rounds they ran: `2R − 1` per chain, summed over the chains, + /// which is `2·ΣR − chains`. + /// + /// ⛔ CHAINS, not groups: `stacked_eval::prove` runs one chain per + /// commitment, so a group can open several. Both numbers are counted by the + /// run, so this is a claim about the call structure that the run can refute + /// — not a constant retyped from a reading of the source. + pub fn derived_rebuild_calls(&self) -> i128 { + 2 * self.round_count as i128 - self.chain_count as i128 + } pub fn is_global(&self) -> bool { self.index == GLOBAL_INDEX } @@ -472,12 +629,16 @@ pub fn push_prover(mut rec: ProverSplit) { // put DECODE's chain under `open_groups`'s name. So: drop what they hold // (which also keeps it out of the next epoch), then take the parked pair. let _ = take_chain(); - let (chain, round_wall) = CHAIN_AT_GROUPS + let parked = CHAIN_AT_GROUPS .lock() - .map(|mut h| std::mem::replace(&mut *h, ([0.0; 6], 0.0))) - .unwrap_or(([0.0; 6], 0.0)); - rec.chain = chain; - rec.chain_round_wall = round_wall; + .map(|mut h| std::mem::take(&mut *h)) + .unwrap_or_default(); + rec.chain = parked.six; + rec.chain_round_wall = parked.round_wall; + rec.queries = parked.queries; + rec.chain_count = parked.chain_count; + rec.round_count = parked.round_count; + rec.rebuild_calls = parked.rebuild_calls; let names = TABLE_NAMES .lock() .map(|mut h| std::mem::take(&mut *h)) @@ -513,7 +674,11 @@ pub fn push_prover(mut rec: ProverSplit) { open_prepared {open_prepared:.2} · other {prove_other:.2} || \ chain[Σ groups] grind {c0:.2} · sumcheck {c1:.2} · fold {c2:.2} · \ commit_folded {c3:.2} · ood {c4:.2} · queries {c5:.2} || \ - round_wall {cw:.2} · round_other {c6:.2} · setup_tail {cst:.2}", + round_wall {cw:.2} · round_other {c6:.2} · setup_tail {cst:.2} || \ + queries[Σ groups] query_sample {q0:.2} · tree_rebuild {q1:.2} · \ + coset_gather {q2:.2} · open_assemble {q3:.2} · queries_other {qo:.2} \ + || chains {chains} · rounds {rounds} · rebuild_calls {rebuilds} \ + (derived {rderiv})", tainted = if rec.overlapped { " ⛔OVERLAPPED" } else { "" }, airs = rec.airs, wall = rec.wall, @@ -537,6 +702,15 @@ pub fn push_prover(mut rec: ProverSplit) { cw = rec.chain_round_wall, c6 = rec.round_other(), cst = rec.setup_tail(), + q0 = rec.queries[0], + q1 = rec.queries[1], + q2 = rec.queries[2], + q3 = rec.queries[3], + qo = rec.queries_other(), + chains = rec.chain_count, + rounds = rec.round_count, + rebuilds = rec.rebuild_calls, + rderiv = rec.derived_rebuild_calls(), ); if let Ok(mut held) = PROVER.lock() { @@ -695,6 +869,74 @@ pub fn check_closure( named.join(" · "), )); } + // Arm E, the query half: the four inside QUERIES, containment first. + let queries_secs = r.chain[5]; + let query_bound = tol * queries_secs.max(0.0) + QUERY_SLACK; + let query_slots = || -> String { + QUERY_NAMES + .iter() + .zip(r.queries.iter()) + .map(|(n, v)| format!("{n} {v:.4}")) + .collect::>() + .join(" · ") + }; + if r.queries_other() < -query_bound { + return Err(format!( + "arm E: {who}'s four query slots OVERRUN the openings that \ + contain them — QUERIES is {queries_secs:.3}s but the four sum \ + to {:.3}s, a NEGATIVE remainder of {:.3}s. A sum of parts \ + cannot exceed the whole that contains it: either a window \ + reaches outside `open_many` or two of them nest. Slots: {}.", + r.queries.iter().sum::(), + r.queries_other(), + query_slots(), + )); + } + if r.queries_other() > query_bound { + return Err(format!( + "arm E: {who}'s four do not account for the query openings — \ + QUERIES is {queries_secs:.3}s but the four inside it sum to \ + only {:.3}s, leaving {:.3}s ({:.1}%) unattributed. The glue \ + between the windows is ~0 on a correct instrument, so a gap \ + this size is A SLOT THAT IS NOT BEING ADDED. Slots: {}.", + r.queries.iter().sum::(), + r.queries_other(), + 100.0 * r.queries_other() / queries_secs.max(1e-9), + query_slots(), + )); + } + // Arm F: the rebuild count, against the rounds that produced it. + // + // `whir_round::prove` opens the CURRENT commitment and the NEXT one, + // and the last round opens only the current through `final_openings`, + // so a chain of R rounds makes `2R − 1` `open_many` calls — one device + // tree rebuild each — and g chains make `2·ΣR − g`. BOTH SIDES ARE + // COUNTED BY THE RUN; neither is a constant read off the source, so + // this is a claim about the call structure that the run can refute. + // + // ⭐ It is also the only arm here that can fire on the card-free + // fixture, where every duration in the chain's query half reads 0.00. + // ⛔ THE GUARD IS "EITHER SIDE IS NONZERO", NOT "THE ROUNDS ARE". + // Guarding on the rounds alone would make the arm blind to exactly one + // of the two omissions it exists to catch: drop the ROUND count and + // `round_count` is 0, the guard skips, and the missing counter is + // invisible. A record that genuinely ran no chain has BOTH at zero, and + // that is the only state this arm may pass over. + if (r.chain_count > 0 || r.round_count > 0 || r.rebuild_calls > 0) + && r.rebuild_calls as i128 != r.derived_rebuild_calls() + { + return Err(format!( + "arm F: {who}'s query openings made {} `open_many` calls, but \ + {} round(s) over {} chain(s) derive {} (2R − 1 per chain). \ + Either a call is not being counted, or the chain no longer \ + opens the current commitment and its successor once each per \ + round.", + r.rebuild_calls, + r.round_count, + r.chain_count, + r.derived_rebuild_calls(), + )); + } } // Arm C: execute + collect + build + handoff partition the producer's wall. for r in producer { @@ -756,6 +998,17 @@ mod tests { "a disabled add must not move a slot" ); assert_eq!(ARGUE.take(), 0.0, "a disabled add must not move a slot"); + // ★ And the COUNTERS, which are the one thing here that is not a clock + // read: `bump` has to check the knob itself, because unlike `add` it + // takes no `Option` that a disabled `mark` could have emptied. + bump(&REBUILD_CALLS); + bump(&ROUND_COUNT); + bump(&CHAIN_COUNT); + assert_eq!( + (REBUILD_CALLS.take(), ROUND_COUNT.take(), CHAIN_COUNT.take()), + (0, 0, 0), + "a disabled bump must not move a counter" + ); assert_eq!(stage_done(0, "execute", mark()), 0.0); note_table(3, 9.0); set_table_names(vec!["KECCAK".to_string()]); @@ -895,6 +1148,22 @@ mod tests { // `open_groups` 0.90 leaves setup_tail 0.09, positive, which is // where the un-slotted setup legitimately lives. chain_round_wall: 0.81, + // The four partition the QUERIES slot (chain[5] = 0.05): + // 0.005 + 0.030 + 0.010 + 0.004 = 0.049, leaving 0.001 of + // glue — the same "~0 on a correct instrument" the round + // loop's own remainder reads. Same lesson as the round wall + // above: a fixture that is not a correct instrument makes + // every arm written against it meaningless. + queries: [0.005, 0.030, 0.010, 0.004], + // Two chains of six rounds: 2·12 − 2 = 22 `open_many` calls, + // which is arm F's identity satisfied rather than asserted. + // `groups` is carried too, and deliberately DIFFERENT from the + // chain count: a group can open several chains, and an identity + // written over groups would pass here by coincidence. + groups: 1, + chain_count: 2, + round_count: 12, + rebuild_calls: 22, ..Default::default() }) .collect(); @@ -1008,6 +1277,13 @@ mod tests { let (producer, mut prover, base) = honest(); prover[0].chain[5] = 0.0; prover[0].chain_round_wall = 0.76; // the wall loses it too + // ⛔ AND THE FOUR INSIDE IT GO WITH IT. This line was added when the + // query half of arm E reddened here on its first run: leaving the four + // at their honest values while zeroing the slot that CONTAINS them + // models four parts summing to more than their whole — an impossible + // instrument, and the very shape arm E exists to reject. The fixture + // was wrong, not the bound. Same lesson as `honest()`'s round wall. + prover[0].queries = [0.0; 4]; assert_eq!( check_closure(&producer, &prover, base, 0.03), Ok(()), @@ -1058,4 +1334,175 @@ mod tests { "a zeroed tolerance must refuse a run carrying real timer cost", ); } + + /// ⛔ EVERY ONE OF THE FOUR, OMITTED IN TURN, IS SEEN — and named. + /// + /// The same discipline as the six: `check_closure` is a pure function over + /// records, so an omission can be FED to it rather than waited for. That + /// matters more here than it did for the six, because the shape the gate + /// runs card-free has every one of these durations at 0.00 and no bound on + /// a time could see anything there (see `QUERY_SLACK`). + #[test] + fn closure_refuses_every_omitted_query_slot() { + for (i, name) in QUERY_NAMES.iter().enumerate() { + let (producer, mut prover, base) = honest(); + let dropped = prover[1].queries[i]; + assert!( + dropped > 0.0, + "{name} must be nonzero in the fixture or \ + this case cannot fail" + ); + prover[1].queries[i] = 0.0; + let err = check_closure(&producer, &prover, base, 0.03).unwrap_err(); + assert!( + err.starts_with("arm E:"), + "expected arm E for {name}, got: {err}" + ); + assert!(err.contains("epoch 1"), "arm E must name the epoch: {err}"); + assert!( + err.contains("A SLOT THAT IS NOT BEING ADDED"), + "arm E must say what the gap means: {err}" + ); + assert!( + err.contains(&format!("{name} 0.0000")), + "arm E must print the four BY NAME so the culprit is visible: {err}" + ); + } + } + + /// ⛔ AND ITS TWIN, WHICH MUST NOT REDDEN: a genuinely tiny QUERIES with + /// the four tiny alongside it — the card-free fixture's own shape. + /// + /// Without this case the next lane meets a gate that reds on every + /// card-free run and widens the bound until it cannot fail at all. That is + /// the failure this pair exists to make impossible. + #[test] + fn closure_accepts_query_slots_that_are_small_because_the_work_was() { + let (producer, mut prover, base) = honest(); + for rec in prover.iter_mut() { + // QUERIES shrinks from 0.05 to 0.002, and the four shrink with it. + rec.chain[5] = 0.002; + rec.queries = [0.0002, 0.0012, 0.0004, 0.0001]; + // The six and the walls move together so the outer arms still hold. + rec.chain_round_wall = rec.chain.iter().sum::() + 0.001; + rec.open_groups = rec.chain_round_wall + 0.09; + rec.prove = rec.challenge + rec.argue + rec.open_groups + rec.open_prepared; + rec.wall = rec.prep + rec.absorb + rec.commit + rec.prove; + } + assert_eq!( + check_closure(&producer, &prover, base, 0.03), + Ok(()), + "a small QUERIES with small parts is an honest run, not a defect" + ); + } + + /// ⛔ A NEGATIVE remainder in the query half: the four cannot exceed the + /// slot that contains them, and it means nesting or a window reaching + /// outside `open_many` — the two defects arm E caught in the six. + #[test] + fn closure_refuses_query_slots_that_overrun_their_opening() { + let (producer, mut prover, base) = honest(); + // tree_rebuild alone made larger than the whole QUERIES slot. + prover[2].queries[1] = 0.09; + let err = check_closure(&producer, &prover, base, 0.03).unwrap_err(); + assert!(err.starts_with("arm E:"), "expected arm E, got: {err}"); + assert!(err.contains("epoch 2"), "arm E must name the epoch: {err}"); + assert!( + err.contains("NEGATIVE"), + "arm E must say what a negative remainder means: {err}" + ); + assert!( + err.contains("tree_rebuild 0.0900"), + "arm E must print the culprit by name: {err}" + ); + } + + /// ★ ARM F — the rebuild count against the rounds that produced it, and + /// the ONE arm that can fire where every duration reads 0.00. + #[test] + fn closure_refuses_a_rebuild_call_that_is_not_counted() { + let (producer, mut prover, base) = honest(); + prover[0].rebuild_calls -= 1; + let err = check_closure(&producer, &prover, base, 0.03).unwrap_err(); + assert!(err.starts_with("arm F:"), "expected arm F, got: {err}"); + assert!(err.contains("epoch 0"), "arm F must name the epoch: {err}"); + assert!( + err.contains("21") && err.contains("22"), + "arm F must print BOTH the counted and the derived number: {err}" + ); + assert!( + err.contains("2R − 1 per chain"), + "arm F must state the identity it is checking: {err}" + ); + } + + /// And arm F fires on the other side too — a chain that stopped opening + /// its successor would make FEWER calls per round, not more. + #[test] + fn closure_refuses_a_round_that_stopped_opening_its_successor() { + let (producer, mut prover, base) = honest(); + // 12 rounds over 2 groups that made only one call per round. + prover[1].rebuild_calls = 12; + let err = check_closure(&producer, &prover, base, 0.03).unwrap_err(); + assert!(err.starts_with("arm F:"), "expected arm F, got: {err}"); + assert!( + err.contains("opens the current commitment and its successor"), + "arm F must name the structure it assumes: {err}" + ); + } + + /// ★ AND THE OTHER OMISSION: the ROUND counter dropped while the calls + /// are still counted. Guarding arm F on `round_count > 0` alone would skip + /// this record entirely and the missing counter would be invisible — the + /// same blind spot the seventh slot opened in arm E, in a new place. + #[test] + fn closure_refuses_a_round_that_was_not_counted() { + let (producer, mut prover, base) = honest(); + prover[2].round_count = 0; + let err = check_closure(&producer, &prover, base, 0.03).unwrap_err(); + assert!(err.starts_with("arm F:"), "expected arm F, got: {err}"); + assert!(err.contains("epoch 2"), "arm F must name the epoch: {err}"); + assert!( + err.contains("22") && err.contains("0 round(s)"), + "arm F must print both counted numbers: {err}" + ); + } + + /// ★ AND THE THIRD OMISSION: the CHAIN counter dropped. The identity is + /// written over all three numbers, so each of them going missing is a + /// different wrong answer and each must be seen. + #[test] + fn closure_refuses_a_chain_that_was_not_counted() { + let (producer, mut prover, base) = honest(); + prover[1].chain_count = 0; + let err = check_closure(&producer, &prover, base, 0.03).unwrap_err(); + assert!(err.starts_with("arm F:"), "expected arm F, got: {err}"); + assert!( + err.contains("0 chain(s)") && err.contains("24"), + "arm F must print the chains it counted and what they derive: {err}" + ); + } + + /// ⛔ AND ARM F MUST NOT FIRE ON A RECORD THAT RAN NO CHAIN. An epoch that + /// opened no groups has no rounds and no calls, and `2·0 − 0 = 0` would + /// hold anyway — but a record with `groups` set and no chain at all would + /// read a derived `-groups`, which is not a defect in the run. + #[test] + fn closure_accepts_a_record_that_ran_no_chain() { + let (producer, mut prover, base) = honest(); + for rec in prover.iter_mut() { + rec.chain_count = 0; + rec.round_count = 0; + rec.rebuild_calls = 0; + rec.chain = [0.0; 6]; + rec.queries = [0.0; 4]; + rec.chain_round_wall = 0.0; + rec.open_groups = 0.9; + } + assert_eq!( + check_closure(&producer, &prover, base, 0.03), + Ok(()), + "no chain is not a broken chain" + ); + } } diff --git a/prover/src/lfm/per_table_aggregator_tests.rs b/prover/src/lfm/per_table_aggregator_tests.rs index 86f763309..8b523b951 100644 --- a/prover/src/lfm/per_table_aggregator_tests.rs +++ b/prover/src/lfm/per_table_aggregator_tests.rs @@ -7420,6 +7420,75 @@ open_groups {v_groups:.1}s ({:.1}%) · open_prepared {v_prepared:.1}s ({:.1}%)", pct(v_groups), pct(v_prepared) ); + // ★ THE CHAIN, AND THE QUERY OPENINGS INSIDE IT — summed over the EPOCH + // records only. The global stage carries its own and is reported on its own + // line below; folding the two would make a ranking nobody could attribute. + let c = |i: usize| sum(&|r| r.chain[i]); + let q = |i: usize| sum(&|r| r.queries[i]); + let queries_secs = c(5); + let share = |x: f64| 100.0 * x / queries_secs.max(1e-9); + println!( + " chain[Σ epochs] grind {:.2}s · sumcheck {:.2}s · fold {:.2}s · \ +commit_folded {:.2}s · ood {:.2}s · queries {:.2}s", + c(0), + c(1), + c(2), + c(3), + c(4), + queries_secs + ); + // ⭐ `tree_rebuild`'s SHARE is round 3's kill condition, computed here + // rather than by whoever reads the log: retention removes the rebuilds and + // nothing else, so if they are not the bulk of the query openings the + // lever is dead before any lifetime code is written. + println!( + " queries[Σ epochs] query_sample {:.2}s ({:.0}%) · tree_rebuild {:.2}s ({:.0}%) · \ +coset_gather {:.2}s ({:.0}%) · open_assemble {:.2}s ({:.0}%) · rebuild_calls {} over {} rounds in {} chains", + q(0), + share(q(0)), + q(1), + share(q(1)), + q(2), + share(q(2)), + q(3), + share(q(3)), + epochs.iter().map(|r| r.rebuild_calls).sum::(), + epochs.iter().map(|r| r.round_count).sum::(), + epochs.iter().map(|r| r.chain_count).sum::() + ); + // ⛔ THE RETENTION LINE IS REQUIRED, and it prints on every run including + // the ones that retain nothing. + // + // The leaf-layer retention has a REFUSAL PATH — `DeviceReservation::grow` + // declines without changing anything when the budget will not take the + // bytes, and the commitment then pays its leaf pass again. That path is + // what makes the scheme safe to run at full size, and it is also what makes + // a silent zero indistinguishable from a lever that never fired. So the + // counts are printed rather than inferred, and a run with nothing to report + // says so in words: the launcher refuses to report a block number without + // this line, exactly as it refuses one without the grind-knobs banner. + let (admitted, refused, asked, got, headroom, saved) = math_cuda::whir::retention_report(); + let (evicted, evicted_bytes) = math_cuda::whir::retention_evictions(); + let mib = |b: u64| b as f64 / (1024.0 * 1024.0); + println!( + " retention[leaf layers] admitted {admitted} · refused {refused} · asked {:.0} MiB · held {:.0} MiB · leaf passes saved {saved} · leaf_passes {} of tree_builds {} · peak simultaneous footprint {:.0} MiB · evicted {evicted} ({:.0} MiB){}", + mib(asked), + mib(got), + math_cuda::whir::leaf_hash_calls(), + math_cuda::whir::tree_builds(), + mib(math_cuda::whir::retained_bytes_peak()), + mib(evicted_bytes), + if refused > 0 { + format!( + " · FIRST REFUSAL at {:.0} MiB of budget headroom", + mib(headroom) + ) + } else if admitted == 0 { + " · NOT TAKEN: no leaf layer was retained in this run".to_string() + } else { + String::new() + }, + ); println!( " global (in base) wall {g_wall:.1}s ({:.1}%)", pct(g_wall) @@ -7447,7 +7516,7 @@ open_groups {v_groups:.1}s ({:.1}%) · open_prepared {v_prepared:.1}s ({:.1}%)", panic!("WHIR BASE SPLIT does not close: {why}"); } println!( - " WHIR BASE SPLIT: closure GREEN (arms A-E at {:.0}% tolerance)", + " WHIR BASE SPLIT: closure GREEN (arms A-F at {:.0}% tolerance)", 100.0 * TOL ); @@ -8082,6 +8151,32 @@ fn the_whir_production_tree_composes_to_a_root() { ), Err(why) => println!(" ⚠ NO ceiling read, so NO percentage: {why}"), } + + // ⛔ THE FALLBACK COUNTS, WHOLE-RUN SCOPE, ALWAYS PRINTED — two DIFFERENT + // device surfaces, each of which silently moves work to the host and leaves + // only host memory and a utilisation dip as its symptoms: + // · commit fallbacks — a WHIR commitment declined the device + // (`multilinear::gpu::host_fallbacks`, one call site, the commit path); + // · device fallbacks — an argue-surface reservation was refused in + // math-cuda (sumcheck/gkr/columns; `math_cuda::device::device_fallbacks`). + // wt16 read as a slot-level win because THIS second number had no name: the + // leaf-layer retention took the shared budget and argue fell to the host + // uncounted. Printed here, whole-run, so the launcher can refuse a block + // number unless BOTH read zero. + println!(" commit fallbacks {}", multilinear::gpu::host_fallbacks()); + println!( + " device fallbacks {}", + math_cuda::device::device_fallbacks() + ); + // The PEAK simultaneous device reservation the run reached — the quantity + // argue's `reserve` is checked against (not the raw device peak, which the + // never-purge pool inflates above the budget). A control run reads argue's + // reservation demand here; while the evictable retention holds only spare + // bytes, this stays below the budget by construction. + println!( + " reserved high-water {:.0} MiB", + math_cuda::device::reserved_high_water() as f64 / (1024.0 * 1024.0) + ); } /// The WHIR tree at FIXTURE scale — the same driver, card-free, on a guest small diff --git a/prover/tests/rpx_grind_bench.rs b/prover/tests/rpx_grind_bench.rs new file mode 100644 index 000000000..f83b2633e --- /dev/null +++ b/prover/tests/rpx_grind_bench.rs @@ -0,0 +1,605 @@ +//! ★ What one device grind costs, and what the two launch knobs do to it. +//! +//! # v3 — the order control, and the one split that can falsify a mechanism +//! +//! v2 ran its arms in one fixed order inside one process and read a monotone +//! fall down the SCAN column — ratios of 1.000, 0.953, 0.872 and 0.800 at scan +//! factors 8, 4, 2 and 1. That contradicts the kernels, which both carry +//! `if (nonce >= *result) break;` against a `volatile` result the `atomicMin` +//! writes through L2, so once the first valid nonce `h` is found every thread +//! stops within one stride round and the executed permutations are `h + stride` +//! WHATEVER the block size. For the median seed, whose `h` is below even the +//! narrowest block here, the two launches are the same kernel doing the same +//! rounds — there is nothing for the knob to change. +//! +//! Pairing on seeds cancels the seed spread. It does not cancel DRIFT, and a +//! monotone fall in run order is what a boosting clock looks like. So: +//! +//! 1. **Every seed runs EVERY arm, in a rotating order** (seed `s` starts at arm +//! `s % N`). Drift now pairs out too: each arm sits in every position of the +//! rotation equally often. +//! 2. **The control is repeated as the LAST arm.** Identical knobs to arm 0, so +//! its ratio must read ≈ 1.00. It is the noise floor of the whole procedure, +//! measured rather than assumed; anything the other arms claim must clear it. +//! 3. **Per-seed PAIRED statistics** — the median of the per-seed ratios and the +//! COUNT of seeds on which the arm beat the control. A real 20% shows on most +//! of 256 seeds; drift shows as a trend that a rotation destroys. +//! 4. ⭐ **THE SPLIT THAT CAN FALSIFY A MECHANISM.** The returned nonce IS `h`, +//! so each seed can be labelled by whether `h` fell inside the arm's block. +//! Seeds with `h < block` take ONE launch at every arm and run the identical +//! kernel — no mechanism can touch them. Seeds with `h >= block` are the only +//! ones that miss and relaunch. So if an arm's gain is the same on both +//! groups, it is NOT the knob; if it lives entirely in the `h >= block` +//! group, there is a real miss-path effect to explain. +//! 5. The COMBINED arms, in case the two effects are real and additive. +//! +//! # v5 — the scan factor swept UPWARD, because `count` is the only thing left +//! +//! v3 and v4 settled the shape and killed two of the three candidates. The +//! typical seed does not see the knob (the per-seed median ratios read 1.000 / +//! 0.999 / 0.988 and the `h`-split is equal in both columns), yet the MEANS +//! fall to 0.804 at scan 1, so a minority of seeds carries a large absolute +//! saving. v4's top-20 named that minority and it is not what anyone predicted: +//! they are SMALL-`h` seeds (258k-512k, well inside 2^20), ONE launch at both +//! arms, and it is the CONTROL that is slow by 5-11 ms while scan 1 costs what +//! `h + stride` predicts. The power limiter is out (these are not the long +//! seeds), the miss-and-relaunch path is out (one launch at both arms), and the +//! volatile load's per-iteration cost is out (the same iterations either way). +//! +//! ⛔ So read what is left. For a one-launch seed the host does NOTHING that +//! scales with `count`: `search` sends one 8-byte sentinel, launches at a grid +//! fixed by the knob, copies 8 bytes back and synchronises. Inside the kernel +//! `count` reaches exactly one thing — the loop bound `i < count`. Work is +//! `h + stride` ONLY IF the early exit stops every thread; a thread that does +//! not observe the `atomicMin` runs to `count`, and its waste is proportional +//! to `count`. That is the one hypothesis the data has not ruled out, and it +//! predicts something the knob can test without touching the kernel. +//! +//! ⭐ **Sweep the scan factor UP.** If the excess is bounded by `count` it must +//! be roughly proportional to it: `excess(k) ∝ (k − 1)·2^20`, so normalised to +//! the control the slope reads (k − 1)/7 — 0.14, 0.43, 1.00, 2.14, 4.43, 9.00 +//! at k = 2, 4, 8, 16, 32, 64. If the excess saturates instead, every ratio +//! above k = 8 reads ~1.00 and this hypothesis dies with the other three. The +//! two branches are a factor of EIGHT apart at k = 64; no clock ramp, thermal +//! drift, ordering or seed spread produces that. +//! +//! The `COUNT SLOPE` section prints that ratio against its prediction, by name, +//! so the verdict is not recomputed downstream — v3's lesson, which cost this +//! file a noise floor that could not pass. +//! +//! ⚠ **This measures wasted work, never a wrong answer.** The nonce lists are +//! identical across every arm and that is asserted before any timing is read: +//! the search returns the globally smallest valid nonce whatever the block +//! size. A straggler burns permutations it did not need to burn. Nothing here +//! is a soundness finding and nothing here moves a proof byte. +//! +//! The grid arms are kept and extended to the same question from the other +//! side: if more resident blocks mean more readers contending the `atomicMin`'s +//! cache line, a wide grid should make stragglers WORSE and a tight `count` +//! should mask it — so `grid 4096` is run at scan 8 AND at scan 1. That is one +//! prediction, not an assumption, and the run is free to refuse it. +//! +//! ```text +//! cargo test -p lambda-vm-prover --release --features cuda \ +//! --test rpx_grind_bench -- --ignored --nocapture +//! ``` +//! +//! Lives in the prover crate rather than `math-cuda` for the same reason +//! `rpx_device_parity.rs` does: the host side (`RpxStarkHash`) lives here, and +//! `math-cuda` is a dev-dependency of this crate, not the reverse. RPX because +//! that is the hash the record posture grinds under. Needs a GPU. +#![cfg(feature = "cuda")] + +use std::time::Instant; + +use lambda_vm_prover::lfm::algebraic_commit::RpxStarkHash; +use math_cuda::grinding::Knobs; +use stark::config::GrindingDigest; +use stark::grinding::{inner_hash_felts, is_valid_nonce}; + +/// The digest the RPX configuration grinds over — its transcript's hash. +type RpxGrind = GrindingDigest; + +/// The production grinding factor. Not a knob here: the bit count is the +/// security parameter, and this measurement is about the launch, not the bits. +const GRINDING_FACTOR: u8 = 20; + +/// Grinds per arm, every arm on the same seeds. +const RUNS: usize = 256; + +/// ★ THE CALIBRATION WINDOW, from the block and not from this bench. +/// +/// The base performs 3,428 device grinds (lb19/lb20: `rpx grinds 3428`, and +/// `states 3428 = grinds`). wt14's grind wall over the base is 14.84 s in the +/// 15 epochs' group openings plus 0.89 s in the global, plus the prepared +/// openings' share which is not separately measured and is bounded by +/// `open_prepared` = 2.3 s. So 15.73-18.03 s over 3,428 grinds. +const BASE_GRINDS: f64 = 3428.0; +const CALIBRATION_LOW_MS: f64 = 15.73 * 1000.0 / BASE_GRINDS; +const CALIBRATION_HIGH_MS: f64 = 18.03 * 1000.0 / BASE_GRINDS; + +/// The arms, in DEFINITION order. The rotation decides RUN order. +/// +/// Arm 0 is the control and the record posture. The last arm repeats it +/// byte-for-byte: same knobs, different position in every rotation, so its +/// ratio is the procedure's own noise floor. +fn arms() -> Vec<(&'static str, Knobs)> { + vec![ + ( + "control 8/1024", + Knobs { + scan: 8, + grid: 1024, + }, + ), + ( + "scan 1", + Knobs { + scan: 1, + grid: 1024, + }, + ), + ( + "scan 2", + Knobs { + scan: 2, + grid: 1024, + }, + ), + ( + "scan 4", + Knobs { + scan: 4, + grid: 1024, + }, + ), + ( + "scan 16", + Knobs { + scan: 16, + grid: 1024, + }, + ), + ( + "scan 32", + Knobs { + scan: 32, + grid: 1024, + }, + ), + ( + "scan 64", + Knobs { + scan: 64, + grid: 1024, + }, + ), + ("grid 512", Knobs { scan: 8, grid: 512 }), + ( + "grid 4096", + Knobs { + scan: 8, + grid: 4096, + }, + ), + ("scan 1 grid 512", Knobs { scan: 1, grid: 512 }), + ( + "scan 1 grid 4096", + Knobs { + scan: 1, + grid: 4096, + }, + ), + ( + "control AGAIN", + Knobs { + scan: 8, + grid: 1024, + }, + ), + ] +} + +/// ★ The scan column, in the order the slope is read down. +/// +/// Every name here must appear in [`arms`] — asserted, not assumed, because a +/// renamed arm would otherwise drop silently out of the slope and the verdict +/// would be computed over fewer points than it claims. Every entry shares the +/// record grid, so `count` is the only thing that moves down this column. +const SCAN_COLUMN: [(&str, u32); 7] = [ + ("scan 1", 1), + ("scan 2", 2), + ("scan 4", 4), + ("control 8/1024", 8), + ("scan 16", 16), + ("scan 32", 32), + ("scan 64", 64), +]; + +/// The arm the slope's excess is measured against: the tightest cap in the +/// sweep, where a straggler can waste least. +const SLOPE_BASE: &str = "scan 1"; + +#[test] +#[ignore = "device benchmark; run with --ignored --nocapture on the GPU box"] +fn what_one_grind_costs_at_each_launch_geometry() { + let arms = arms(); + let n = arms.len(); + let seeds: Vec<[u8; 32]> = (0..RUNS).map(seed_for).collect(); + + print_device_fill(); + + // Warm-up, excluded by name: the first launch in a process pays context + // creation and the cubin load. + let warm = one_grind(&seeds[0], Knobs::DEFAULT); + println!( + "GRIND BENCH: warm-up nonce {} in {:.3} ms (EXCLUDED)", + warm.0, warm.1 + ); + + // The environment path, exercised once and tied to the explicit path, so + // the knob is shown READ on the same path the tree runs use. + let env_knobs = math_cuda::grinding::knobs_in_effect(); + let env_nonce = math_cuda::grinding::generate_nonce_rpx_gpu( + &inner_hash_felts::(&seeds[0], GRINDING_FACTOR), + GRINDING_FACTOR, + ) + .expect("GPU RPX grind through the environment path (needs a GPU)"); + assert_eq!( + env_nonce, + one_grind(&seeds[0], env_knobs).0, + "the environment path and the explicit path disagree at {env_knobs:?}" + ); + println!("GRIND BENCH: environment path agrees with the explicit path at {env_knobs:?}"); + + // ── THE RUN: seeds outer, arms inner, ROTATED ────────────────────────── + let mut times = vec![vec![0.0f64; RUNS]; n]; + let mut nonces = vec![vec![0u64; RUNS]; n]; + for (s, seed) in seeds.iter().enumerate() { + for j in 0..n { + let a = (s + j) % n; + let (nonce, ms) = one_grind(seed, arms[a].1); + times[a][s] = ms; + nonces[a][s] = nonce; + } + } + println!( + "GRIND BENCH: {} arms x {RUNS} seeds, rotated (seed s starts at arm s % {})", + n, n + ); + + // ★ THE NONCE CONTROL, first: if the geometry moved the answer, nothing + // else here means anything. + for (a, (name, knobs)) in arms.iter().enumerate() { + assert_eq!( + nonces[a], nonces[0], + "{name} ({knobs:?}) returned a different nonce list from the control \ + — the launch geometry moved the answer, so the search is not \ + scanning contiguously from zero" + ); + } + println!("NONCE CONTROL: all {n} arms returned identical nonce lists over {RUNS} seeds"); + + // ── THE TABLE ────────────────────────────────────────────────────────── + println!("\n=== THE ARMS (paired per seed, rotated order) ==="); + println!( + "{:<18} {:>9} {:>9} {:>9} {:>11} {:>10} {:>9} {:>9}", + "arm", "mean ms", "median", "ns/perm", "median rat", "beat/256", "rat h=blk" + ); + for (a, (name, knobs)) in arms.iter().enumerate() { + print_arm(name, *knobs, ×[a], ×[0], &nonces[0]); + } + + // ⛔ THE NOISE FLOOR, ON A LINE OF ITS OWN AND NAMED. + // + // v3's launcher read this by COLUMN POSITION — `awk '{print $5}'` over the + // repeated control's row — and `control AGAIN` is two words, so it read the + // `ns/perm` column instead of the ratio and declared the procedure 327% + // unstable on a run whose ratio was 1.000. A verdict read by position + // breaks the first time a label gains a space. This line exists so nothing + // downstream has to count spaces to find the truth. + let floor_ratios: Vec = times[n - 1] + .iter() + .zip(×[0]) + .map(|(a, b)| a / b) + .collect(); + let floor = median(&floor_ratios); + println!( + "\nNOISE FLOOR: repeated control median per-seed ratio {floor:.4} \ + (|1 - r| = {:.4}; the procedure's own movement — no arm may claim less)", + (1.0 - floor).abs() + ); + + // ── THE DISTRIBUTION, because a flat median with a falling mean is a + // distribution statement and deciles are the cheapest way to make it one. + println!("\n=== THE PER-SEED RATIO, BY DECILE ==="); + print!("{:<18}", "arm"); + for d in 1..10 { + print!("{:>7}", format!("p{}0", d)); + } + println!(); + for (a, (name, _)) in arms.iter().enumerate() { + let mut rs: Vec = times[a].iter().zip(×[0]).map(|(x, y)| x / y).collect(); + rs.sort_by(|x, y| x.partial_cmp(y).expect("no NaN in a ratio")); + print!("{name:<18}"); + for d in 1..10 { + print!("{:>7.3}", rs[(d * rs.len()) / 10]); + } + println!(); + } + + // ★ THE TWENTY SEEDS THAT MOVE THE MEAN, for the arm that moves it most. + // + // The mean falls while every median stays flat, so a minority of seeds + // carries the saving. These are that minority, with the one quantity that + // actually differs between the arms for a given seed: the LAUNCH COUNT. + // It is DERIVED, not instrumented — the loop advances `base` by the block + // each miss and returns on the block containing the hit, so the count is + // `h / block + 1` exactly. + let worst = (1..n) + .min_by(|a, b| mean(×[*a]).total_cmp(&mean(×[*b]))) + .expect("at least one arm beside the control"); + let (worst_name, worst_knobs) = arms[worst]; + let ctl_block = arms[0].1.block(GRINDING_FACTOR); + let arm_block = worst_knobs.block(GRINDING_FACTOR); + println!("\n=== THE 20 SEEDS THAT MOVE THE MEAN MOST — {worst_name} vs the control ==="); + println!( + "{:>5} {:>12} {:>9} {:>9} {:>9} {:>9} {:>9}", + "seed", "h", "ctl ms", "arm ms", "delta ms", "ctl lch", "arm lch" + ); + let mut order: Vec = (0..RUNS).collect(); + order.sort_by(|x, y| { + (times[worst][*y] - times[0][*y]) + .abs() + .total_cmp(&(times[worst][*x] - times[0][*x]).abs()) + }); + for s in order.into_iter().take(20) { + let h = nonces[0][s]; + println!( + "{s:>5} {h:>12} {:>9.3} {:>9.3} {:>9.3} {:>9} {:>9}", + times[0][s], + times[worst][s], + times[worst][s] - times[0][s], + h / ctl_block + 1, + h / arm_block + 1, + ); + } + println!( + " ⭐ v4 ALREADY READ THIS TABLE and it killed three candidates: the movers are \ + SMALL-h seeds taking ONE launch at BOTH arms, and the CONTROL is the slow one. \ + So not the power limiter (these are not the long seeds), not the miss-and-relaunch \ + path (one launch either way), not the volatile load's per-iteration cost (the same \ + iterations either way). What is left is a thread that did not stop, and the COUNT \ + SLOPE below is what tests it." + ); + + // ── ★ THE COUNT SLOPE — v5's verdict, computed here and named ────────── + // + // ⛔ Printed BESIDE its prediction and read by NAME, so no launcher has to + // recompute it or count columns to find it. That is v3's lesson: its noise + // floor was read by column position, `control AGAIN` is two words, and the + // verdict it produced could not pass on any run. + // + // The excess is measured against the TIGHTEST cap in the sweep rather than + // against a model: at scan 1 a straggler can waste at most 2^20 nonces, so + // whatever sits above that arm is what a larger `count` bought. No fitted + // constant enters, which is what keeps this from being a model checking + // itself. + let at = |name: &str| -> usize { + arms.iter() + .position(|(n, _)| *n == name) + .unwrap_or_else(|| panic!("the scan column names `{name}`, which is not an arm")) + }; + let base_mean = mean(×[at(SLOPE_BASE)]); + let ctl_excess = mean(×[at("control 8/1024")]) - base_mean; + println!("\n=== THE COUNT SLOPE (pre-registered: excess proportional to count - 2^20) ==="); + println!( + "{:<18} {:>11} {:>10} {:>11} {:>12} {:>11}", + "arm", "count", "mean ms", "excess ms", "vs control", "PREDICTED" + ); + for (name, k) in SCAN_COLUMN { + let a = at(name); + let excess = mean(×[a]) - base_mean; + let block = Knobs { + scan: k, + grid: 1024, + } + .block(GRINDING_FACTOR); + println!( + "{name:<18} {block:>11} {:>10.3} {:>11.3} {:>12.3} {:>11.3}", + mean(×[a]), + excess, + excess / ctl_excess, + (f64::from(k) - 1.0) / 7.0, + ); + } + if ctl_excess.abs() < 1.0e-3 { + println!( + " ⛔ THE CONTROL SHOWS NO EXCESS over {SLOPE_BASE} ({ctl_excess:.6} ms), so the \ + `vs control` column is a ratio to zero and says NOTHING. That is itself the \ + answer: without an excess at the record posture there is no tail to explain." + ); + } + println!( + " ⭐ COUNT-BOUND if `vs control` tracks `PREDICTED` up the sweep (2.14 / 4.43 / 9.00 \ + at k = 16 / 32 / 64): a subset of threads runs to `count` because it never observed \ + the early exit, and the scan factor is a CAP on that waste, not a performance knob." + ); + println!( + " ⭐ SATURATED if every ratio at k >= 8 reads ~1.00: the excess is a fixed per-launch \ + cost that merely correlates with the scan factor, the straggler hypothesis dies with \ + the other three, and that cost owes a name." + ); + println!( + " ⚠ WASTED WORK, NEVER A WRONG ANSWER: the nonce control above already asserted that \ + every arm returned the identical nonce list. Whatever this column reads, no proof \ + byte moves and no verifier check is touched." + ); + + // ── THE CALIBRATION CONTROL ──────────────────────────────────────────── + let control_mean = mean(×[0]); + let projected = control_mean * BASE_GRINDS / 1000.0; + let inside = (CALIBRATION_LOW_MS..=CALIBRATION_HIGH_MS).contains(&control_mean); + println!("\n=== THE CALIBRATION CONTROL ==="); + println!( + "this bench at the record posture: {control_mean:.3} ms/grind; the block's window \ + {CALIBRATION_LOW_MS:.3}-{CALIBRATION_HIGH_MS:.3} ms/grind (15.73-18.03 s over \ + {BASE_GRINDS:.0} grinds, wt14)" + ); + println!( + "projected over the base's grinds: {projected:.2} s against 15.73-18.03 s => {}", + if inside { + "IN — this bench measures the block's grind" + } else { + "OUT — this bench is NOT the block's grind and nothing above is quotable" + } + ); + + println!("\n=== HOW TO READ IT (pre-registered) ==="); + println!( + " `control AGAIN` is the NOISE FLOOR: its median ratio must read ~1.00. Any arm \ + claiming less than that floor is claiming noise." + ); + println!( + " ⭐ `rat h=blk`: seeds whose hit fell INSIDE the arm's block take \ + ONE launch and run the identical kernel at every arm, so no knob can touch them. \ + An arm whose gain is the SAME on both groups is not the knob — it is the procedure. \ + A gain living only in `h>=blk` is a real miss-path effect and owes a mechanism." + ); + println!( + " ⭐ v5's VERDICT IS THE COUNT SLOPE, and it was written down before the run: \ + COUNT-BOUND means `vs control` follows `PREDICTED` to 9.00 at scan 64; SATURATED \ + means it flattens at ~1.00 from scan 8 up. Nothing between those two readings is \ + claimed here, and the arms that decide it are the three ABOVE the record posture — \ + which no earlier version of this bench ever ran." + ); + println!( + " ⭐ THE GRID PAIR TESTS THE SAME DEFECT FROM THE BLOCK-COUNT SIDE: `grid 4096` at \ + scan 8 against `grid 4096` at scan 1. If wide grids make stragglers worse by \ + contending the atomicMin's line, the wide arm should hurt at scan 8 and be MASKED \ + at scan 1, where `count` caps the waste. Equal damage at both caps refuses that \ + unification and leaves the grid column its own explanation." + ); + println!( + " ⚠ EVERY ARM HERE IS A MEASUREMENT, NOT A PROPOSAL. The record posture is scan 8 / \ + grid 1024 and this run changes no default. Scan 16, 32 and 64 exist to make the \ + waste visible by exaggerating it; they are not candidates for anything." + ); +} + +/// One arm's row, paired against the control seed by seed. +fn print_arm(name: &str, knobs: Knobs, t: &[f64], ctl: &[f64], nonces: &[u64]) { + let block = knobs.block(GRINDING_FACTOR); + let ratios: Vec = t.iter().zip(ctl).map(|(a, b)| a / b).collect(); + let beat = ratios.iter().filter(|r| **r < 1.0).count(); + let inside: Vec = ratios + .iter() + .zip(nonces) + .filter(|(_, h)| **h < block) + .map(|(r, _)| *r) + .collect(); + let outside: Vec = ratios + .iter() + .zip(nonces) + .filter(|(_, h)| **h >= block) + .map(|(r, _)| *r) + .collect(); + let stride = knobs.stride(math_cuda::grinding::RPX_BLOCK_DIM) as f64; + let perms: f64 = nonces.iter().map(|h| *h as f64 + stride).sum(); + println!( + "{name:<18} {:>9.3} {:>9.3} {:>9.2} {:>11.3} {:>6}/{:<3} {:>9} {:>10}", + mean(t), + median(t), + t.iter().sum::() * 1.0e6 / perms, + median(&ratios), + beat, + ratios.len(), + fmt_group(&inside), + fmt_group(&outside), + ); +} + +/// A group's median ratio and its size, or a dash when the group is empty. +fn fmt_group(rs: &[f64]) -> String { + if rs.is_empty() { + "—".to_string() + } else { + format!("{:.3}/{}", median(rs), rs.len()) + } +} + +fn mean(xs: &[f64]) -> f64 { + xs.iter().sum::() / xs.len() as f64 +} + +fn median(xs: &[f64]) -> f64 { + let mut v = xs.to_vec(); + v.sort_by(|a, b| a.partial_cmp(b).expect("no NaN in a timing")); + v[v.len() / 2] +} + +/// One grind at the production factor, timed around the device call alone. +/// +/// The nonce is re-validated on the host exactly as the prover's dispatch does, +/// so a timing arm cannot quietly become a measurement of a kernel that returns +/// garbage quickly. +fn one_grind(seed: &[u8; 32], knobs: Knobs) -> (u64, f64) { + let felts = inner_hash_felts::(seed, GRINDING_FACTOR); + let started = Instant::now(); + let nonce = math_cuda::grinding::generate_nonce_rpx_gpu_at(&felts, GRINDING_FACTOR, knobs) + .expect("GPU RPX grind (needs a GPU)"); + let ms = started.elapsed().as_secs_f64() * 1000.0; + assert!( + is_valid_nonce::(seed, nonce, GRINDING_FACTOR), + "GPU nonce {nonce} from {knobs:?} fails is_valid_nonce — a timing number \ + from an invalid nonce is not a measurement" + ); + (nonce, ms) +} + +/// Distinct seeds, so each arm pays its own spread of hit distances. +fn seed_for(i: usize) -> [u8; 32] { + let mut seed = [0u8; 32]; + seed[..8].copy_from_slice(&(i as u64).to_le_bytes()); + seed[8] = 0xA5; + seed +} + +/// ★ What the driver says about filling this card — every field READ, so the +/// residency question the grid knob turns on is answered rather than estimated. +fn print_device_fill() { + println!("\n=== DEVICE FILL (read from the driver, not estimated) ==="); + match math_cuda::grinding::device_fill() { + None => println!("DEVICE FILL: unavailable — no device, or the driver refused the query"), + Some(fill) => { + println!( + "SMs {} · max threads/SM {} · rpx grind kernel: {} regs/thread, block dim {}, \ + {} resident blocks/SM", + fill.sm_count, + fill.max_threads_per_sm, + fill.rpx_regs_per_thread, + fill.rpx_block_dim, + fill.rpx_blocks_per_sm + ); + println!( + "⇒ resident ceiling: {} blocks = {} threads", + fill.resident_blocks(), + fill.resident_threads() + ); + println!("{:<8} {:>16} {:>10}", "grid", "fill vs ceiling", "stride"); + for grid in [256u32, 512, 1024, 2048, 4096] { + println!( + "{:<8} {:>15.2}x {:>10}", + grid, + fill.fill(grid), + Knobs { scan: 8, grid }.stride(fill.rpx_block_dim) + ); + } + println!( + "⇒ a grid above {} queues: the surplus blocks buy no parallelism and their \ + stride is pure overshoot past the first hit", + fill.resident_blocks() + ); + } + } +} diff --git a/prover/tests/rpx_grind_counted.rs b/prover/tests/rpx_grind_counted.rs new file mode 100644 index 000000000..2948dc3a9 --- /dev/null +++ b/prover/tests/rpx_grind_counted.rs @@ -0,0 +1,363 @@ +//! ★ STAGE B: do the slow grind launches do MORE WORK, or the same work slower? +//! +//! # What the timings could not answer +//! +//! `rpx_grind_bench` reads a mean about 20% above the model at the record +//! posture while the MEDIAN seed sits exactly on it, so a minority of launches +//! carries a large absolute cost. Three candidates died to that bench: the +//! power limiter (the movers are SMALL-`h` seeds, not long sustained launches), +//! the miss-and-relaunch path (the movers take ONE launch at both arms) and the +//! volatile load's per-iteration cost (the same iterations either way). +//! +//! v5 swept the scan factor UPWARD to test the last one standing — threads +//! running to the loop bound. They do not: the excess SATURATES above `count` +//! = 2^23 instead of growing with it (`vs control` 1.05 at scan 16, 32 and 64 +//! against a prediction of 2.14, 4.43 and 9.00). +//! +//! But it does not saturate immediately either. Converted to iterations per +//! thread, `N = count / stride`, the excess fits a bounded quantity approached +//! geometrically: +//! +//! ```text +//! N 8 16 32 64 128 256 512 +//! ms 0 0.354 0.763 1.004 1.055 1.057 1.060 +//! excess = T·(1 − e^(−(N−8)/τ)), T ≈ 1.06 ms, τ ≈ 19 iterations +//! ``` +//! +//! Three independent points agree on τ within 4%. That is the shape of a +//! thread that keeps scanning for a bounded TIME after the answer is known — +//! about 19 further iterations, a per-iteration stopping probability near 5% — +//! and NOT the shape of one running to the loop bound. One iteration at grid +//! 1024 is 131,072 permutations, ≈ 0.56 ms at the bench's own 4.27 ns/perm, so +//! 19 iterations is ≈ 10.6 ms, which is where v4's top movers sat (5-11 ms). +//! +//! ⚠ That is a three-point fit with two parameters, measured against a baseline +//! (scan 1, eight iterations) that may itself be carrying capped excess. It is +//! a hypothesis, not a reading. This file replaces it with a reading. +//! +//! # What this measures +//! +//! `rpx_grind_search_counted` counts, on the device, the permutations its +//! threads actually ran. So `executed − (h + stride)` is the overrun, per +//! search, as a number rather than a model — with `max_iters` saying how deep +//! the deepest thread went and `ran_to_end` how many threads left by the loop +//! bound rather than by the early exit. +//! +//! **PRE-REGISTERED, before the run:** +//! +//! - STALE-POLL: the overrun is ≈ 0 on most searches and ≈ 19 × stride on a +//! minority — and on those, *the same at scan 8 and at scan 64*, because a +//! stop bounded by time does not care about the bound. `max_iters` reads +//! `ceil(h/stride) + ~20`, never `count/stride`. ⭐ `ran_to_end` is the +//! discriminator: near ZERO at scan 64, where 512 iterations are available +//! and a thread stops after ~19; NONZERO at scan 1, where the cap of 8 bites +//! first. +//! - NO OVERRUN: `executed − (h + stride)` is ≈ 0 everywhere, including on the +//! slow searches. Then no thread over-scans, the slow launches run the SAME +//! permutations more slowly, the whole straggler family is dead, and the cost +//! is outside this loop and owes a name. +//! +//! # ⛔ The control that decides whether this file may be read at all +//! +//! A counted kernel with different register pressure has different occupancy +//! and therefore measures a different kernel. Every arm therefore runs the +//! SHIPPED kernel on the same seed immediately beside the counted one, and the +//! two must agree on the nonce (always) and on the milliseconds (within the +//! procedure's own noise floor). If the milliseconds disagree, this run reports +//! that the instrument changed the phenomenon and draws no conclusion. +//! +//! ```text +//! cargo test -p lambda-vm-prover --release --features cuda \ +//! --test rpx_grind_counted -- --ignored --nocapture +//! ``` +//! +//! Needs a GPU. Changes no default and no shipped kernel. +#![cfg(feature = "cuda")] + +use std::time::Instant; + +use lambda_vm_prover::lfm::algebraic_commit::RpxStarkHash; +use math_cuda::grinding::{GrindCounts, Knobs}; +use stark::config::GrindingDigest; +use stark::grinding::{inner_hash_felts, is_valid_nonce}; + +type RpxGrind = GrindingDigest; + +/// The production grinding factor: this is about the launch, not the bits. +const GRINDING_FACTOR: u8 = 20; + +/// The same 256 seeds v5 used, so the two runs describe the same population. +const RUNS: usize = 256; + +/// The noise floor v5 measured on this box: the repeated control's median +/// per-seed ratio came back 0.9992. The admissibility control below allows +/// five times that, because it compares two DIFFERENT kernels and a tie is not +/// what is being claimed — only that the twin did not change the phenomenon. +const ADMISSIBLE_MS_RATIO: f64 = 0.05; + +/// The arms: the two scan factors that bracket the saturation, the one above +/// it, and the grid that moves the typical seed. +fn arms() -> Vec<(&'static str, Knobs)> { + vec![ + ( + "scan 1", + Knobs { + scan: 1, + grid: 1024, + }, + ), + ( + "scan 8 (record)", + Knobs { + scan: 8, + grid: 1024, + }, + ), + ( + "scan 64", + Knobs { + scan: 64, + grid: 1024, + }, + ), + ( + "scan 8 grid 4096", + Knobs { + scan: 8, + grid: 4096, + }, + ), + ( + "scan 1 grid 4096", + Knobs { + scan: 1, + grid: 4096, + }, + ), + ] +} + +fn seed_for(i: usize) -> [u8; 32] { + let mut seed = [0u8; 32]; + seed[..8].copy_from_slice(&(i as u64).to_le_bytes()); + seed[8] = 0xA5; + seed +} + +/// The shipped kernel, timed — the control arm of every pair. +fn shipped(seed: &[u8; 32], knobs: Knobs) -> (u64, f64) { + let felts = inner_hash_felts::(seed, GRINDING_FACTOR); + let started = Instant::now(); + let nonce = math_cuda::grinding::generate_nonce_rpx_gpu_at(&felts, GRINDING_FACTOR, knobs) + .expect("GPU RPX grind (needs a GPU)"); + (nonce, started.elapsed().as_secs_f64() * 1000.0) +} + +/// The counted twin, timed the same way. +fn counted(seed: &[u8; 32], knobs: Knobs) -> (GrindCounts, f64) { + let felts = inner_hash_felts::(seed, GRINDING_FACTOR); + let started = Instant::now(); + let counts = math_cuda::grinding::search_counted(&felts, GRINDING_FACTOR, knobs) + .expect("counted RPX grind (needs a GPU)"); + (counts, started.elapsed().as_secs_f64() * 1000.0) +} + +fn mean(xs: &[f64]) -> f64 { + xs.iter().sum::() / xs.len() as f64 +} + +fn median(xs: &[f64]) -> f64 { + let mut v = xs.to_vec(); + v.sort_by(|a, b| a.partial_cmp(b).expect("no NaN in a timing")); + v[v.len() / 2] +} + +#[test] +#[ignore = "device diagnostic; run with --ignored --nocapture on the GPU box"] +fn what_the_slow_grind_launches_actually_execute() { + let arms = arms(); + let seeds: Vec<[u8; 32]> = (0..RUNS).map(seed_for).collect(); + + // Warm-up, excluded by name, on BOTH kernels: the first launch of each + // pays its own cubin load. + let w1 = shipped(&seeds[0], Knobs::DEFAULT); + let w2 = counted(&seeds[0], Knobs::DEFAULT); + println!( + "WARM-UP (EXCLUDED): shipped {:.3} ms, counted {:.3} ms, nonces {} / {}", + w1.1, w2.1, w1.0, w2.0.nonce + ); + + // Per arm: the mean overrun in strides, the summed `ran_to_end`, and + // whether the twin was admissible. The cross-arm verdict is computed from + // these HERE rather than by whoever reads the log — the same reason the + // count slope is printed beside its prediction in `rpx_grind_bench`. + let mut summary: Vec<(&'static str, f64, u64, bool)> = Vec::new(); + + for (name, knobs) in arms.iter() { + let stride = knobs.stride(math_cuda::grinding::RPX_BLOCK_DIM); + let block = knobs.block(GRINDING_FACTOR); + let mut ship_ms = Vec::with_capacity(RUNS); + let mut cnt_ms = Vec::with_capacity(RUNS); + let mut rows: Vec<(usize, GrindCounts, f64, f64, i128)> = Vec::with_capacity(RUNS); + + for (s, seed) in seeds.iter().enumerate() { + // Paired and adjacent, so a clock that drifts drifts through both. + let (nonce, t_ship) = shipped(seed, *knobs); + let (counts, t_cnt) = counted(seed, *knobs); + + // ⛔ THE ANSWER IS PINNED FIRST. A diagnostic that returns a + // different nonce is measuring a different search. + assert_eq!( + counts.nonce, nonce, + "{name}, seed {s}: the counted kernel returned {} and the \ + shipped kernel {nonce} — they are not running the same search", + counts.nonce + ); + assert!( + is_valid_nonce::(seed, counts.nonce, GRINDING_FACTOR), + "{name}, seed {s}: nonce {} fails is_valid_nonce", + counts.nonce + ); + + // The model: every nonce below the hit, plus one stride round for + // the threads that were mid-permutation when it landed. + // + // ⛔ IT IS `nonce + stride`, FULL STOP, AND THE MISSED BLOCKS ARE + // ALREADY IN IT. The first draft added `(launches − 1) · block` on + // top, reasoning that a missed block costs its whole `count`. It + // does — but `nonce` is ABSOLUTE, so those nonces are counted once + // already and the term double-counted them. The symptom was a + // NEGATIVE overrun on exactly the arms where searches miss (scan 1 + // read −4.17 strides, which is not a quantity that can be negative), + // while scan 8 and scan 64 were untouched because at those block + // sizes every seed here hits on its first launch. + let ideal = counts.nonce + stride; + let _ = block; + let overrun = counts.executed as i128 - ideal as i128; + ship_ms.push(t_ship); + cnt_ms.push(t_cnt); + rows.push((s, counts, t_ship, t_cnt, overrun)); + } + + // ── the admissibility control, before any reading ────────────────── + let (m_ship, m_cnt) = (mean(&ship_ms), mean(&cnt_ms)); + let ratio = m_cnt / m_ship; + let admissible = (ratio - 1.0).abs() <= ADMISSIBLE_MS_RATIO; + println!( + "\n=== {name}: count {block}, stride {stride}, {} iterations available ===", + block / stride + ); + println!( + "ADMISSIBILITY: shipped {m_ship:.3} ms vs counted {m_cnt:.3} ms, ratio {ratio:.4} \ + (allowed |1 - r| <= {ADMISSIBLE_MS_RATIO}) => {}", + if admissible { + "ADMISSIBLE — the twin did not change the phenomenon" + } else { + "NOT ADMISSIBLE — the counters changed the kernel; read no overrun from this arm" + } + ); + + let overruns: Vec = rows.iter().map(|r| r.4 as f64).collect(); + let affected = rows.iter().filter(|r| r.4 > stride as i128).count(); + let to_end: u64 = rows.iter().map(|r| r.1.ran_to_end).sum(); + println!( + "OVERRUN (executed - ideal), permutations: mean {:.0} · median {:.0} · \ + in strides mean {:.2} · searches overrunning by > 1 stride {affected}/{RUNS}", + mean(&overruns), + median(&overruns), + mean(&overruns) / stride as f64, + ); + println!( + "RAN_TO_END (threads leaving by the loop bound, summed over {RUNS} searches): \ + {to_end}" + ); + summary.push((name, mean(&overruns) / stride as f64, to_end, admissible)); + + // The ten searches with the largest overrun, with everything that could + // explain them. + let mut order: Vec = (0..RUNS).collect(); + order.sort_by(|a, b| rows[*b].4.cmp(&rows[*a].4)); + println!( + "{:>5} {:>12} {:>8} {:>14} {:>11} {:>10} {:>9} {:>9}", + "seed", + "h", + "launches", + "overrun perms", + "in strides", + "max_iters", + "ship ms", + "cnt ms" + ); + for &i in order.iter().take(10) { + let (s, c, t_ship, t_cnt, over) = &rows[i]; + println!( + "{s:>5} {:>12} {:>8} {over:>14} {:>11.2} {:>10} {t_ship:>9.3} {t_cnt:>9.3}", + c.nonce, + c.launches, + *over as f64 / stride as f64, + c.max_iters, + ); + } + } + + // ── ★ THE VERDICT, by name, beside what each branch predicted ────────── + let at = + |n: &str| -> Option<&(&'static str, f64, u64, bool)> { summary.iter().find(|r| r.0 == n) }; + println!("\n=== ★ THE OVERRUN VERDICT ==="); + println!( + "{:<20} {:>16} {:>14} {:>14}", + "arm", "overrun/stride", "ran_to_end", "admissible" + ); + for (name, over, ends, ok) in &summary { + println!("{name:<20} {over:>16.2} {ends:>14} {:>14}", ok); + } + match (at("scan 8 (record)"), at("scan 64")) { + (Some(s8), Some(s64)) if s8.3 && s64.3 => { + // Time-bounded means the SAME overrun however much room the loop + // bound leaves; count-bounded would have grown eightfold here. + let grew = s64.1 / s8.1.max(1e-9); + if s8.1 > 5.0 && (0.5..=2.0).contains(&grew) { + println!( + " ⇒ STALE-POLL: the overrun is {:.1} strides at scan 8 and {:.1} at scan \ + 64, a ratio of {grew:.2} where a loop-bound cause would read about 8. \ + Threads DO execute extra permutations and the excess is bounded by TIME, \ + not by `count` — a poll of `*result` served stale. The fix is reader-side.", + s8.1, s64.1 + ); + } else if s8.1 <= 1.0 && s64.1 <= 1.0 { + println!( + " ⇒ NO OVERRUN: {:.2} and {:.2} strides. Nothing over-scans; the slow \ + launches run the SAME permutations more slowly, the straggler family is \ + dead, and the cost is outside this loop and owes a name.", + s8.1, s64.1 + ); + } else { + println!( + " ⇒ NEITHER BRANCH: {:.2} strides at scan 8 and {:.2} at scan 64 (ratio \ + {grew:.2}) match neither pre-registered reading. Report it unresolved \ + rather than rounding it to a verdict.", + s8.1, s64.1 + ); + } + } + (Some(_), Some(_)) => println!( + " ⛔ NO VERDICT: one of the two arms was NOT ADMISSIBLE, so its counters describe \ + a kernel with different occupancy from the one that ships." + ), + _ => { + println!(" ⛔ NO VERDICT: the two arms the verdict is written over did not both run.") + } + } + + println!( + "\n=== HOW TO READ IT (pre-registered, not re-derived) ===\n \ + STALE-POLL: the overrun is ~0 on most searches and ~19 strides on a minority, the \ + SAME on those at scan 8 and scan 64; max_iters ~ ceil(h/stride) + 20, never \ + count/stride; ran_to_end near ZERO at scan 64 and NONZERO at scan 1.\n \ + NO OVERRUN: the overrun is ~0 everywhere including the slow searches — nothing \ + over-scans, the slow launches run the same permutations more slowly, and the cost \ + is outside this loop.\n \ + ⚠ An arm reported NOT ADMISSIBLE says nothing either way: its counters describe a \ + kernel with different occupancy from the one that ships." + ); +}