From 2ba8eed7118eb1c9ba033c980729bb553368edbb Mon Sep 17 00:00:00 2001 From: MauroFab Date: Sat, 19 Sep 2026 16:03:34 -0300 Subject: [PATCH 01/17] feat(math-cuda): the grind's scan factor as a read-once knob, and what it is not MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit `LAMBDA_VM_GRIND_SCAN_FACTOR` (default 8 — the record posture, unmoved) replaces the literal 8 at the one site that sizes a device grind's launch block. Read once per process through a `OnceLock`, refused outside 1..=64 with the offending value named, and printed as `★ GRIND SCAN FACTOR: n` on the first device grind, so a log that quotes the factor can be shown to have read it rather than assumed it. ⛔ The knob is NOT the lever it was ruled to be, and the doc comment now says why. Both grind kernels carry `if (nonce >= *result) break;` against a `volatile` result the `atomicMin` writes through L2, and the stride walk gives every nonce in `[base, base+count)` exactly one owner — so the scan stops at the first hit. The permutations executed are `h + stride` whatever the block size, and the launches before the hitting one cover exactly the part of `[0, h)` below it. The factor buys only the probability that one launch suffices, `1 - e^-k`. Lowering it removes no permutations (they were never executed) and adds `1/(1 - e^-k)` expected round trips. The same reading says the returned nonce is the globally smallest valid one at any factor — which `tests/grinding.rs::gpu_grind_returns_smallest_valid_nonce` already pins — so sweeping the knob moves no proof byte. `prover/tests/rpx_grind_bench.rs` is the arm that settles this on the card in seconds rather than in four tree runs: ms/grind and the full nonce list at one scan factor per process. Flat means the block is a ceiling; halving means the scan dominates; an identical nonce list across the arms is the byte control. --- crypto/math-cuda/src/grinding.rs | 210 ++++++++++++++++++++++++++++++- prover/tests/rpx_grind_bench.rs | 129 +++++++++++++++++++ 2 files changed, 333 insertions(+), 6 deletions(-) create mode 100644 prover/tests/rpx_grind_bench.rs diff --git a/crypto/math-cuda/src/grinding.rs b/crypto/math-cuda/src/grinding.rs index 1a83a859a..42672e172 100644 --- a/crypto/math-cuda/src/grinding.rs +++ b/crypto/math-cuda/src/grinding.rs @@ -42,6 +42,113 @@ 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`]; read once per process and printed. +/// +/// # ⛔ 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; + +/// ★ The scan factor for this process, read once and cached. +/// +/// Prints on every setting including the default, so a log that quotes the +/// factor can be shown to have read it rather than assumed it. Aborts on a +/// value outside [`SCAN_FACTOR_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. +fn scan_factor() -> u32 { + static FACTOR: std::sync::OnceLock = std::sync::OnceLock::new(); + *FACTOR.get_or_init(|| { + let factor = match std::env::var(SCAN_FACTOR_ENV) { + Err(_) => SCAN_FACTOR_DEFAULT, + Ok(raw) => parse_scan_factor(raw.trim()).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!( + "{SCAN_FACTOR_ENV}={raw:?} is not a scan factor this path accepts. \ + Accepted: an integer in {}..={}.", + SCAN_FACTOR_RANGE.start(), + SCAN_FACTOR_RANGE.end() + ); + std::process::abort() + }), + }; + println!("★ GRIND SCAN FACTOR: {factor}"); + factor + }) +} + +/// ★ The scan factor this process is searching under — the SAME cached value +/// [`search`] uses, not a second read of the environment. +/// +/// Exists so a measurement can print the knob it actually ran under rather +/// than the one its launcher believed it exported (a quoted knob has to be +/// shown read on the path). Calling it also forces the banner, so a log with no +/// `★ GRIND SCAN FACTOR:` line is a run that never reached the device search. +pub fn scan_factor_in_effect() -> u32 { + scan_factor() +} + +/// ★ The nonces one launch would cover at `grinding_factor` under the scan +/// factor in effect — what [`search`] passes the kernel as `count`. +pub fn per_launch_block(grinding_factor: u8) -> u64 { + block_size(grinding_factor, scan_factor()) +} + +/// The accepted spellings: a plain decimal integer inside the range. +fn parse_scan_factor(raw: &str) -> Option { + let factor: u32 = raw.parse().ok()?; + SCAN_FACTOR_RANGE.contains(&factor).then_some(factor) +} + +/// 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. +fn block_size(grinding_factor: u8, scan_factor: u32) -> u64 { + let expected = 1u64.checked_shl(grinding_factor as u32).unwrap_or(u64::MAX); + expected + .saturating_mul(scan_factor as u64) + .clamp(MIN_BLOCK, MAX_BLOCK) +} + /// Which outer hash the search runs. #[derive(Clone, Copy, Debug, PartialEq, Eq)] enum Arm { @@ -89,6 +196,12 @@ 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). +/// +/// ★ **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) -> Option { if !(GRIND_MIN_FACTOR..=64).contains(&grinding_factor) { return None; @@ -103,12 +216,7 @@ 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 = block_size(grinding_factor, scan_factor()); let cfg = LaunchConfig { grid_dim: (GRID_DIM, 1, 1), @@ -146,3 +254,93 @@ fn search(arm: Arm, inner: &[u64; 4], grinding_factor: u8) -> Option { base = base.checked_add(count)?; } } + +#[cfg(test)] +mod tests { + //! The knob's parser and its block arithmetic, card-free: nothing here + //! touches `backend()`, so these run wherever the crate compiles. + //! + //! `scan_factor()` itself is deliberately NOT exercised — it caches in a + //! `OnceLock`, so a test that set the variable would fix the value for + //! every other test in the binary and the second assertion would pass on + //! the first one's cache. The sweep sets it per PROCESS, which is what the + //! bench arms do. + + use super::{MAX_BLOCK, MIN_BLOCK, SCAN_FACTOR_DEFAULT, block_size, parse_scan_factor}; + + /// Every arm of the sweep parses, and an untouched environment is the + /// default. Spelled as the four values the sweep actually launches, so a + /// parser that stopped accepting one of them reddens here. + #[test] + fn the_parser_accepts_every_arm_of_the_sweep() { + for k in [8u32, 4, 2, 1] { + assert_eq!( + parse_scan_factor(&k.to_string()), + Some(k), + "scan factor {k} is a sweep arm and must parse" + ); + } + assert_eq!(parse_scan_factor("64"), Some(64), "the top of the range"); + assert_eq!(parse_scan_factor("16"), Some(16), "inside the range"); + } + + /// The refusing half, executed. Each of these would otherwise reach the + /// clamp 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", "65", "abc", "", "-1", "8.0", "0x8", " ", "4294967296"] { + assert_eq!( + parse_scan_factor(raw), + None, + "{raw:?} must be refused, not defaulted" + ); + } + } + + /// 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!(block_size(20, 8), 8 * 1_048_576, "2^20 * 8 = 2^23"); + assert_eq!(block_size(20, 4), 4 * 1_048_576, "2^20 * 4 = 2^22"); + assert_eq!(block_size(20, 2), 2 * 1_048_576, "2^20 * 2 = 2^21"); + assert_eq!(block_size(20, 1), 1_048_576, "2^20 * 1 = 2^20"); + } + + /// The default is the record posture: every measurement in the campaign + /// was taken at a 2^23 block, and the ABBA is what may move it. + #[test] + fn the_default_is_the_record_posture() { + assert_eq!(SCAN_FACTOR_DEFAULT, 8, "the record posture's scan factor"); + assert_eq!( + block_size(20, SCAN_FACTOR_DEFAULT), + 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. + #[test] + fn the_clamp_binds_at_both_ends() { + assert_eq!( + block_size(12, 1), + MIN_BLOCK, + "the min-factor gate's own factor, at the narrowest scan, floors" + ); + assert_eq!(block_size(26, 8), MAX_BLOCK, "2^26 * 8 = 2^29 ceils"); + assert_eq!( + block_size(64, 8), + MAX_BLOCK, + "2^64 saturates before the multiply, then ceils" + ); + assert!( + (MIN_BLOCK..=MAX_BLOCK).contains(&block_size(20, 8)), + "the record posture sits strictly inside the clamp, so neither \ + end is silently setting it" + ); + } +} diff --git a/prover/tests/rpx_grind_bench.rs b/prover/tests/rpx_grind_bench.rs new file mode 100644 index 000000000..692cffa7c --- /dev/null +++ b/prover/tests/rpx_grind_bench.rs @@ -0,0 +1,129 @@ +//! ★ What one device grind costs, and what the scan factor does to it. +//! +//! Round 2's lever was ruled from `math-cuda/src/grinding.rs` — one launch +//! sizes its block at 8x the expected hit distance — on the reading that the +//! 8x is 8x the permutations. The kernels say otherwise: both carry +//! `if (nonce >= *result) break;` against a `volatile` result the `atomicMin` +//! writes through L2, and the stride walk gives every nonce in the block +//! exactly one owner, so the scan STOPS at the first hit and the factor is a +//! ceiling the launch never reaches. +//! +//! This is the arm that settles it on the card rather than from the source, in +//! seconds rather than in four tree runs: +//! +//! ```text +//! LAMBDA_VM_GRIND_SCAN_FACTOR=8 cargo test -p lambda-vm-prover --release \ +//! --features cuda --test rpx_grind_bench -- --ignored --nocapture +//! ``` +//! +//! run once per factor in 8 / 4 / 2 / 1. Read ACROSS the arms: +//! +//! * **ms/grind** — flat across the four ⇒ the block size is a ceiling and the +//! ruled lever is dead; halving with the factor ⇒ the scan dominates and it +//! is real. The k = 1 arm additionally prices the extra round trips (1.58 +//! expected launches per grind against 1.00 at k = 8). +//! * **the nonces** — IDENTICAL across the four, or the search is not scanning +//! contiguously from zero and the block size is moving the answer. The nonce +//! is a function of the inner hash and the grinding factor alone; this is the +//! control that says so, and it is also why the sweep moves no proof byte. +//! +//! 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 +//! (`LAMBDA_VM_WHIR_HASH=rpx`). Needs a GPU. +#![cfg(feature = "cuda")] + +use std::time::Instant; + +use lambda_vm_prover::lfm::algebraic_commit::RpxStarkHash; +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 block, not the bits. +const GRINDING_FACTOR: u8 = 20; + +/// Grinds timed per arm. Each uses a distinct seed, so each pays its own hit +/// distance and the mean is taken over the geometric spread rather than over +/// one lucky nonce repeated. +const RUNS: usize = 32; + +#[test] +#[ignore = "device benchmark; run with --ignored --nocapture on the GPU box"] +fn what_one_grind_costs_at_this_scan_factor() { + // Read through the production accessor, so the line below reports the knob + // the search is actually using and not a second reading of the + // environment. This also forces the `★ GRIND SCAN FACTOR:` banner. + let scan = math_cuda::grinding::scan_factor_in_effect(); + let block = math_cuda::grinding::per_launch_block(GRINDING_FACTOR); + let expected: u64 = 1 << GRINDING_FACTOR; + + println!("GRIND BENCH: rpx256 · grinding factor {GRINDING_FACTOR} · scan factor {scan}"); + println!( + "GRIND BENCH: expected hit distance {expected} · per-launch block {block} ({:.2}x) \ + · expected launches per grind {:.3}", + block as f64 / expected as f64, + 1.0 / (1.0 - (-(scan as f64)).exp()) + ); + + // A warm-up grind, excluded from the statistics: the first launch in the + // process pays context creation and the cubin load, which is not what a + // grind costs in a proof that has already done three thousand of them. + let (warm, warm_ms) = grind(0xFF); + println!("GRIND BENCH: warm-up nonce {warm} {warm_ms:.3} ms (EXCLUDED)"); + + let mut total = 0.0f64; + let mut worst = 0.0f64; + let mut best = f64::INFINITY; + for i in 0..RUNS { + let (nonce, ms) = grind(i as u8); + total += ms; + worst = worst.max(ms); + best = best.min(ms); + // Every nonce printed, because the cross-arm control is an equality + // between two runs' nonce LISTS and a mean cannot carry it. + println!("GRIND BENCH: seed {i:02} nonce {nonce} {ms:.3} ms"); + } + + let mean = total / RUNS as f64; + println!( + "GRIND BENCH RESULT: scan {scan} · {RUNS} grinds · mean {mean:.3} ms · min {best:.3} \ + · max {worst:.3}" + ); + // Both competing models, printed, so the arm reading this log does not have + // to recompute them — and so the comparison across arms stays a comparison + // of MEASURED means rather than of a mean against a model. + println!( + "GRIND BENCH MODEL: if the block is scanned in full, {:.1} M perm/s; if the scan stops \ + at the first hit, {:.1} M perm/s", + block as f64 / (mean / 1000.0) / 1e6, + expected as f64 / (mean / 1000.0) / 1e6 + ); +} + +/// One grind at the production factor over a seed derived from `tag`, 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. A device search that failed would return `None` +/// here and take the `expect`, so every timed call is a real device grind. +fn grind(tag: u8) -> (u64, f64) { + let seed = [tag; 32]; + let felts = inner_hash_felts::(&seed, GRINDING_FACTOR); + let started = Instant::now(); + let nonce = math_cuda::grinding::generate_nonce_rpx_gpu(&felts, GRINDING_FACTOR) + .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} fails is_valid_nonce (factor {GRINDING_FACTOR}) — a timing number \ + from an invalid nonce is not a measurement" + ); + (nonce, ms) +} From 45388b143b5a6994c8b5fe9e5b3233b0bca4c4a5 Mon Sep 17 00:00:00 2001 From: MauroFab Date: Sat, 19 Sep 2026 16:23:37 -0300 Subject: [PATCH 02/17] feat(math-cuda): the grind's launch geometry as knobs, and the fill read off the driver MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit `LAMBDA_VM_GRIND_GRID` joins `LAMBDA_VM_GRIND_SCAN_FACTOR` in one module, both read once, both defaulting to today's exact values (8 and 1024) so the record posture is byte-unchanged. ONE line prints both AND the stride each arm gets — `★ GRIND KNOBS: scan 8 · grid 1024 · stride rpx 131072 / keccak 262144` — because the stride is the mechanism and a reader should not have to multiply it back out. The two block dims stay constants: they are tuned per kernel against register pressure, which belongs to the kernel body, and moving them would change what an occupancy reading means. Why the GRID is the candidate lever now that the scan factor is not. The kernels stop at the first hit, so a search executes `h + stride` permutations and `stride = grid × block_dim` is the term left behind — the overshoot is `stride/h`, 12.5% at the default. That gives the knob two opposite edges: while the card is not filled a wider grid raises throughput faster than overshoot, and once it is filled the surplus blocks only queue and the wider stride is pure added work. The sweep therefore has to run BOTH ways. `device_fill()` answers which edge the default sits on by READING the driver — SM count, max threads per SM, the kernel's registers per thread, and the occupancy the driver will actually grant — instead of estimating residency from a block dim. A grid above the resident-block ceiling buys no parallelism. `search` now takes its knobs as a parameter, so `generate_nonce_{gpu,rpx_gpu}_at` can sweep them inside ONE process. That is not a convenience: the knobs cache in a `OnceLock`, so comparing settings through the environment would need a process per arm, and four processes are four device contexts, four cubin loads and four clock domains compared across an exponential spread of hit distances. Paired arms on identical seeds make the ratios exact instead. `prover/tests/rpx_grind_bench.rs` runs the nine arms that way — scan 8/4/2/1 at grid 1024 and grid 256/512/1024/2048/4096 at scan 8, the 8/1024 arm shared — over 256 seeds at the production factor, reporting mean, median and `ns/perm`. The nonce IS the hit distance, so the permutations a launch executed are known exactly and `ns/perm` is the seed-independent throughput the grid question turns on. Three controls travel with it: the environment path is exercised and asserted to agree with the explicit one, every arm's nonce list must be identical, and the measured ms/grind is projected over the base's 3,428 grinds against the window wt14 read (15.73-18.03 s) and reported in or out. `crypto/math-cuda/tests/grinding.rs` gains the card-side twin: the nonce is the same, and still the smallest, at every scan factor and every grid. --- crypto/math-cuda/src/grinding.rs | 433 ++++++++++++++++++++++------- crypto/math-cuda/tests/grinding.rs | 87 ++++++ prover/tests/rpx_grind_bench.rs | 337 ++++++++++++++++------ 3 files changed, 681 insertions(+), 176 deletions(-) diff --git a/crypto/math-cuda/src/grinding.rs b/crypto/math-cuda/src/grinding.rs index 42672e172..57272623b 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 @@ -50,7 +54,7 @@ const MIN_BLOCK: u64 = 1 << 18; const MAX_BLOCK: u64 = 1 << 28; /// ★ `LAMBDA_VM_GRIND_SCAN_FACTOR` — how many expected hit distances one launch -/// covers. Default [`SCAN_FACTOR_DEFAULT`]; read once per process and printed. +/// covers. Default [`SCAN_FACTOR_DEFAULT`]. /// /// # ⛔ This is a CEILING on the block, not a multiplier on the work /// @@ -85,68 +89,196 @@ pub const SCAN_FACTOR_DEFAULT: u32 = 8; /// one, and silently so. pub const SCAN_FACTOR_RANGE: std::ops::RangeInclusive = 1..=64; -/// ★ The scan factor for this process, read once and cached. +/// ★ `LAMBDA_VM_GRIND_GRID` — blocks per launch. Default [`GRID_DEFAULT`]. /// -/// Prints on every setting including the default, so a log that quotes the -/// factor can be shown to have read it rather than assumed it. Aborts on a -/// value outside [`SCAN_FACTOR_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. -fn scan_factor() -> u32 { - static FACTOR: std::sync::OnceLock = std::sync::OnceLock::new(); - *FACTOR.get_or_init(|| { - let factor = match std::env::var(SCAN_FACTOR_ENV) { - Err(_) => SCAN_FACTOR_DEFAULT, - Ok(raw) => parse_scan_factor(raw.trim()).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!( - "{SCAN_FACTOR_ENV}={raw:?} is not a scan factor this path accepts. \ - Accepted: an integer in {}..={}.", - SCAN_FACTOR_RANGE.start(), - SCAN_FACTOR_RANGE.end() - ); - std::process::abort() - }), - }; - println!("★ GRIND SCAN FACTOR: {factor}"); - factor - }) +/// # 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, } -/// ★ The scan factor this process is searching under — the SAME cached value -/// [`search`] uses, not a second read of the environment. +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. /// -/// Exists so a measurement can print the knob it actually ran under rather -/// than the one its launcher believed it exported (a quoted knob has to be -/// shown read on the path). Calling it also forces the banner, so a log with no -/// `★ GRIND SCAN FACTOR:` line is a run that never reached the device search. -pub fn scan_factor_in_effect() -> u32 { - scan_factor() +/// 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 scan -/// factor in effect — what [`search`] passes the kernel as `count`. +/// ★ 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 { - block_size(grinding_factor, scan_factor()) + knobs_in_effect().block(grinding_factor) } -/// The accepted spellings: a plain decimal integer inside the range. -fn parse_scan_factor(raw: &str) -> Option { - let factor: u32 = raw.parse().ok()?; - SCAN_FACTOR_RANGE.contains(&factor).then_some(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() + }), + } } -/// Nonces one launch covers: the expected hit distance `2^grinding_factor` -/// times the scan factor, clamped to `[MIN_BLOCK, MAX_BLOCK]`. +/// 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. /// -/// `2^grinding_factor` can overflow u64 (factor 64), so saturate. -fn block_size(grinding_factor: u8, scan_factor: u32) -> u64 { - let expected = 1u64.checked_shl(grinding_factor as u32).unwrap_or(u64::MAX); - expected - .saturating_mul(scan_factor as u64) - .clamp(MIN_BLOCK, MAX_BLOCK) +/// 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. @@ -162,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 @@ -186,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. @@ -202,7 +367,7 @@ pub fn generate_nonce_rpx_gpu(inner_felts: &[u64; 4], grinding_factor: u8) -> Op /// 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) -> Option { +fn search(arm: Arm, inner: &[u64; 4], grinding_factor: u8, knobs: Knobs) -> Option { if !(GRIND_MIN_FACTOR..=64).contains(&grinding_factor) { return None; } @@ -216,10 +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()?; - let count = block_size(grinding_factor, scan_factor()); + 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, }; @@ -257,45 +422,77 @@ fn search(arm: Arm, inner: &[u64; 4], grinding_factor: u8) -> Option { #[cfg(test)] mod tests { - //! The knob's parser and its block arithmetic, card-free: nothing here + //! The knobs' parser and their launch arithmetic, card-free: nothing here //! touches `backend()`, so these run wherever the crate compiles. //! - //! `scan_factor()` itself is deliberately NOT exercised — it caches in a - //! `OnceLock`, so a test that set the variable would fix the value for + //! `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. The sweep sets it per PROCESS, which is what the - //! bench arms do. + //! 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::{MAX_BLOCK, MIN_BLOCK, SCAN_FACTOR_DEFAULT, block_size, parse_scan_factor}; + 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 sweep parses, and an untouched environment is the - /// default. Spelled as the four values the sweep actually launches, so a - /// parser that stopped accepting one of them reddens here. + /// 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_the_sweep() { + fn the_parser_accepts_every_arm_of_both_sweeps() { for k in [8u32, 4, 2, 1] { assert_eq!( - parse_scan_factor(&k.to_string()), + parse_knob(&k.to_string(), &SCAN_FACTOR_RANGE), Some(k), "scan factor {k} is a sweep arm and must parse" ); } - assert_eq!(parse_scan_factor("64"), Some(64), "the top of the range"); - assert_eq!(parse_scan_factor("16"), Some(16), "inside the range"); + 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. Each of these would otherwise reach the - /// clamp and become a DIFFERENT search that no log could distinguish from - /// the one the operator asked for. + /// 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", "65", "abc", "", "-1", "8.0", "0x8", " ", "4294967296"] { + 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_scan_factor(raw), + parse_knob(raw, &GRID_RANGE), None, - "{raw:?} must be refused, not defaulted" + "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 @@ -304,43 +501,91 @@ mod tests { /// ⇒ 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!(block_size(20, 8), 8 * 1_048_576, "2^20 * 8 = 2^23"); - assert_eq!(block_size(20, 4), 4 * 1_048_576, "2^20 * 4 = 2^22"); - assert_eq!(block_size(20, 2), 2 * 1_048_576, "2^20 * 2 = 2^21"); - assert_eq!(block_size(20, 1), 1_048_576, "2^20 * 1 = 2^20"); + 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 default is the record posture: every measurement in the campaign - /// was taken at a 2^23 block, and the ABBA is what may move it. + /// ★ 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_default_is_the_record_posture() { + 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!( - block_size(20, SCAN_FACTOR_DEFAULT), + 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. + /// 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() { + fn the_clamp_binds_at_both_ends_and_the_grid_does_not_touch_it() { assert_eq!( - block_size(12, 1), + at_scan(1).block(12), MIN_BLOCK, "the min-factor gate's own factor, at the narrowest scan, floors" ); - assert_eq!(block_size(26, 8), MAX_BLOCK, "2^26 * 8 = 2^29 ceils"); + assert_eq!(at_scan(8).block(26), MAX_BLOCK, "2^26 * 8 = 2^29 ceils"); assert_eq!( - block_size(64, 8), + at_scan(8).block(64), MAX_BLOCK, "2^64 saturates before the multiply, then ceils" ); assert!( - (MIN_BLOCK..=MAX_BLOCK).contains(&block_size(20, 8)), - "the record posture sits strictly inside the clamp, so neither \ - end is silently setting it" + (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/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/prover/tests/rpx_grind_bench.rs b/prover/tests/rpx_grind_bench.rs index 692cffa7c..dd3414ca5 100644 --- a/prover/tests/rpx_grind_bench.rs +++ b/prover/tests/rpx_grind_bench.rs @@ -1,43 +1,58 @@ -//! ★ What one device grind costs, and what the scan factor does to it. +//! ★ What one device grind costs, and what the two launch knobs do to it. //! -//! Round 2's lever was ruled from `math-cuda/src/grinding.rs` — one launch -//! sizes its block at 8x the expected hit distance — on the reading that the -//! 8x is 8x the permutations. The kernels say otherwise: both carry +//! Round 2's lever was ruled from `math-cuda/src/grinding.rs`: one launch sizes +//! its block at 8x the expected hit distance, so "drop the 8 and drop 8x the +//! permutations". The kernels say otherwise — both carry //! `if (nonce >= *result) break;` against a `volatile` result the `atomicMin` //! writes through L2, and the stride walk gives every nonce in the block -//! exactly one owner, so the scan STOPS at the first hit and the factor is a -//! ceiling the launch never reaches. -//! -//! This is the arm that settles it on the card rather than from the source, in -//! seconds rather than in four tree runs: +//! exactly one owner, so the scan STOPS at the first hit. The executed +//! permutations are `h + stride` at ANY scan factor; the block size is a +//! ceiling the launch never reaches, and `stride = grid x block_dim` is the +//! term that is left. //! //! ```text -//! LAMBDA_VM_GRIND_SCAN_FACTOR=8 cargo test -p lambda-vm-prover --release \ -//! --features cuda --test rpx_grind_bench -- --ignored --nocapture +//! cargo test -p lambda-vm-prover --release --features cuda \ +//! --test rpx_grind_bench -- --ignored --nocapture //! ``` //! -//! run once per factor in 8 / 4 / 2 / 1. Read ACROSS the arms: +//! ONE process, every arm, and that is deliberate: the knobs cache in a +//! `OnceLock`, so reading two settings through the environment would need two +//! processes — two device contexts, two cubin loads, two clock domains. Every +//! arm here runs on the SAME seeds through `generate_nonce_rpx_gpu_at`, so the +//! comparison is PAIRED on identical hit distances and the seed-to-seed spread +//! (which is exponential, and wide) cancels between arms instead of being +//! averaged away. The environment path is exercised too, and asserted to agree. //! -//! * **ms/grind** — flat across the four ⇒ the block size is a ceiling and the -//! ruled lever is dead; halving with the factor ⇒ the scan dominates and it -//! is real. The k = 1 arm additionally prices the extra round trips (1.58 -//! expected launches per grind against 1.00 at k = 8). -//! * **the nonces** — IDENTICAL across the four, or the search is not scanning -//! contiguously from zero and the block size is moving the answer. The nonce -//! is a function of the inner hash and the grinding factor alone; this is the -//! control that says so, and it is also why the sweep moves no proof byte. +//! WHAT TO READ, pre-registered: //! -//! 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. +//! * **the scan arms** (8/4/2/1 at grid 1024) — FLAT means the block size is a +//! ceiling and the ruled lever is dead; halving with the factor means the +//! scan dominates and it is real. The k=1 arm prices the extra round trips +//! (1.58 expected launches per grind against 1.00). +//! * **the grid arms** (256/512/1024/2048/4096 at scan 8) — the sweep runs BOTH +//! ways because the knob has two edges: under-filled, a wider grid raises +//! throughput faster than overshoot and the wall falls; filled, extra blocks +//! only queue and the wider stride is pure added work. `DEVICE FILL` below +//! says which edge the default sits on, READ from the driver. +//! * **`ns/perm`** — the invariant. Time is about `(h/stride + 1)` rounds, so +//! `ms/grind` carries the seed's hit distance while `ns/perm` does not; the +//! grid question is a throughput question and this is the throughput. +//! * **the nonce lists** — identical across every arm, or the search is not +//! scanning contiguously from zero and the geometry is moving the answer. +//! * **the CALIBRATION CONTROL** — this bench's ms/grind at the record posture +//! times the base's 3,428 grinds, against the base's grind wall read from +//! wt14. Out of that window and nothing here describes the block's grind. //! -//! RPX because that is the hash the record posture grinds under -//! (`LAMBDA_VM_WHIR_HASH=rpx`). Needs a GPU. +//! 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}; @@ -45,85 +60,243 @@ use stark::grinding::{inner_hash_felts, is_valid_nonce}; type RpxGrind = GrindingDigest; /// The production grinding factor. Not a knob here: the bit count is the -/// security parameter, and this measurement is about the block, not the bits. +/// security parameter, and this measurement is about the launch, not the bits. const GRINDING_FACTOR: u8 = 20; -/// Grinds timed per arm. Each uses a distinct seed, so each pays its own hit -/// distance and the mean is taken over the geometric spread rather than over -/// one lucky nonce repeated. -const RUNS: usize = 32; +/// Grinds per arm, all arms on the same seeds. The hit distance is exponential +/// with mean `2^20`, so a single seed says nothing; pairing across arms is what +/// makes the ratios precise, and this many keeps the unpaired mean's standard +/// error near 6%. +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; + +/// One arm's reading. +struct Arm { + knobs: Knobs, + mean_ms: f64, + median_ms: f64, + min_ms: f64, + max_ms: f64, + /// Time divided by the permutations actually executed, `h + stride`. The + /// returned nonce IS `h`, so this is measured rather than modelled. + ns_per_perm: f64, + nonces: Vec, +} #[test] #[ignore = "device benchmark; run with --ignored --nocapture on the GPU box"] -fn what_one_grind_costs_at_this_scan_factor() { - // Read through the production accessor, so the line below reports the knob - // the search is actually using and not a second reading of the - // environment. This also forces the `★ GRIND SCAN FACTOR:` banner. - let scan = math_cuda::grinding::scan_factor_in_effect(); - let block = math_cuda::grinding::per_launch_block(GRINDING_FACTOR); - let expected: u64 = 1 << GRINDING_FACTOR; - - println!("GRIND BENCH: rpx256 · grinding factor {GRINDING_FACTOR} · scan factor {scan}"); +fn what_one_grind_costs_at_each_launch_geometry() { + 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, which is not what a grind costs in a proof + // that has already done three thousand of them. + let warm = one_grind(&seeds[0], Knobs::DEFAULT); println!( - "GRIND BENCH: expected hit distance {expected} · per-launch block {block} ({:.2}x) \ - · expected launches per grind {:.3}", - block as f64 / expected as f64, - 1.0 / (1.0 - (-(scan as f64)).exp()) + "GRIND BENCH: warm-up nonce {} in {:.3} ms (EXCLUDED)", + warm.0, warm.1 ); - // A warm-up grind, excluded from the statistics: the first launch in the - // process pays context creation and the cubin load, which is not what a - // grind costs in a proof that has already done three thousand of them. - let (warm, warm_ms) = grind(0xFF); - println!("GRIND BENCH: warm-up nonce {warm} {warm_ms:.3} ms (EXCLUDED)"); - - let mut total = 0.0f64; - let mut worst = 0.0f64; - let mut best = f64::INFINITY; - for i in 0..RUNS { - let (nonce, ms) = grind(i as u8); - total += ms; - worst = worst.max(ms); - best = best.min(ms); - // Every nonce printed, because the cross-arm control is an equality - // between two runs' nonce LISTS and a mean cannot carry it. - println!("GRIND BENCH: seed {i:02} nonce {nonce} {ms:.3} ms"); + // The environment path, exercised once so the knob is shown READ on the + // same path the tree runs use — and its answer tied to the explicit path. + 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)"); + let explicit_nonce = one_grind(&seeds[0], env_knobs).0; + assert_eq!( + env_nonce, explicit_nonce, + "the environment path and the explicit path disagree at {env_knobs:?} — \ + one of them is not the search the other is" + ); + println!("GRIND BENCH: environment path agrees with the explicit path at {env_knobs:?}"); + + let scan_arms: Vec = [8u32, 4, 2, 1] + .into_iter() + .map(|scan| run_arm(Knobs { scan, grid: 1024 }, &seeds)) + .collect(); + let grid_arms: Vec = [256u32, 512, 2048, 4096] + .into_iter() + .map(|grid| run_arm(Knobs { scan: 8, grid }, &seeds)) + .collect(); + + // The record posture is the shared arm: the head of the scan sweep and the + // 1024 point of the grid sweep are the same measurement, not two. + let record = &scan_arms[0]; + + println!("\n=== THE ARMS (paired: every arm ran the same {RUNS} seeds) ==="); + println!( + "{:<14} {:>10} {:>10} {:>10} {:>10} {:>10} {:>9}", + "knobs", "mean ms", "median", "min", "max", "ns/perm", "vs 8/1024" + ); + print_arm(record, record); + println!("-- the scan sweep, grid 1024 (pre-registered FLAT if the block is a ceiling) --"); + for arm in &scan_arms[1..] { + print_arm(arm, record); + } + println!("-- the grid sweep, scan 8 (pre-registered to fall on ONE side of 1024, not both) --"); + for arm in &grid_arms { + print_arm(arm, record); } - let mean = total / RUNS as f64; + // ★ THE NONCE CONTROL. Every arm must have returned the same answer for the + // same seed; the launch geometry cannot move it, and that is why sweeping + // either knob moves no proof byte. + for arm in scan_arms.iter().chain(&grid_arms) { + assert_eq!( + arm.nonces, record.nonces, + "{:?} returned a different nonce list from the record posture — the \ + launch geometry moved the answer, so the search is not scanning \ + contiguously from zero and nothing else here is quotable", + arm.knobs + ); + } println!( - "GRIND BENCH RESULT: scan {scan} · {RUNS} grinds · mean {mean:.3} ms · min {best:.3} \ - · max {worst:.3}" + "\nNONCE CONTROL: all {} arms returned identical nonce lists over {RUNS} seeds", + scan_arms.len() + grid_arms.len() + ); + + // ★ THE CALIBRATION CONTROL, said in or out. + let projected = record.mean_ms * BASE_GRINDS / 1000.0; + let inside = (CALIBRATION_LOW_MS..=CALIBRATION_HIGH_MS).contains(&record.mean_ms); + println!("\n=== THE CALIBRATION CONTROL ==="); + println!( + "this bench at the record posture: {:.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)", + record.mean_ms ); - // Both competing models, printed, so the arm reading this log does not have - // to recompute them — and so the comparison across arms stays a comparison - // of MEASURED means rather than of a mean against a model. println!( - "GRIND BENCH MODEL: if the block is scanned in full, {:.1} M perm/s; if the scan stops \ - at the first hit, {:.1} M perm/s", - block as f64 / (mean / 1000.0) / 1e6, - expected as f64 / (mean / 1000.0) / 1e6 + "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" + } ); } -/// One grind at the production factor over a seed derived from `tag`, timed -/// around the device call alone. +/// One arm: every seed, one process, explicit knobs. +fn run_arm(knobs: Knobs, seeds: &[[u8; 32]]) -> Arm { + let stride = knobs.stride(math_cuda::grinding::RPX_BLOCK_DIM) as f64; + let mut times = Vec::with_capacity(seeds.len()); + let mut nonces = Vec::with_capacity(seeds.len()); + let mut total_ns = 0.0f64; + let mut total_perms = 0.0f64; + for seed in seeds { + let (nonce, ms) = one_grind(seed, knobs); + total_ns += ms * 1.0e6; + // The returned nonce IS the first valid one, so the permutations the + // launch executed are `nonce + stride` — measured, not modelled. + total_perms += nonce as f64 + stride; + times.push(ms); + nonces.push(nonce); + } + let mean_ms = times.iter().sum::() / times.len() as f64; + let mut sorted = times.clone(); + sorted.sort_by(|a, b| a.partial_cmp(b).expect("no NaN in a timing")); + Arm { + knobs, + mean_ms, + median_ms: sorted[sorted.len() / 2], + min_ms: sorted[0], + max_ms: sorted[sorted.len() - 1], + ns_per_perm: total_ns / total_perms, + nonces, + } +} + +/// 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. A device search that failed would return `None` -/// here and take the `expect`, so every timed call is a real device grind. -fn grind(tag: u8) -> (u64, f64) { - let seed = [tag; 32]; - let felts = inner_hash_felts::(&seed, GRINDING_FACTOR); +/// 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. A device search that failed returns `None` and takes the +/// `expect`, so every timed call is a real device grind. +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(&felts, GRINDING_FACTOR) + 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} fails is_valid_nonce (factor {GRINDING_FACTOR}) — a timing number \ + 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} {:>12} {:>10}", "grid", "fill vs ceiling", "stride"); + for grid in [256u32, 512, 1024, 2048, 4096] { + println!( + "{:<8} {:>11.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() + ); + } + } +} + +fn print_arm(arm: &Arm, record: &Arm) { + println!( + "scan {:<3} grid {:<5} {:>10.3} {:>10.3} {:>10.3} {:>10.3} {:>10.2} {:>8.3}", + arm.knobs.scan, + arm.knobs.grid, + arm.mean_ms, + arm.median_ms, + arm.min_ms, + arm.max_ms, + arm.ns_per_perm, + arm.mean_ms / record.mean_ms + ); +} From 41a2457c6d26ef0f50e0e8a4ceb4e82128565461 Mon Sep 17 00:00:00 2001 From: MauroFab Date: Sat, 19 Sep 2026 16:53:18 -0300 Subject: [PATCH 03/17] test(bench): the order control, and the split that can falsify a mechanism MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The first run 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 — which the kernels say cannot exist. Both stop at the first hit, so the executed permutations are `h + stride` whatever the block size, and for the median seed, whose hit falls inside even the narrowest block here, the two launches are the same kernel doing the same rounds. There is nothing for the knob to change. The arms ran in one fixed order in one process, so that fall is confounded with drift. Pairing on seeds cancels the seed spread; it does not cancel a boosting clock. Four changes to the procedure, none to the measurement: Every seed now runs every arm in a rotating order, so each arm sits in every position of the rotation equally often and drift pairs out too. The control is repeated as a final arm with identical knobs: its ratio is the noise floor, measured rather than assumed, and no arm may claim less than it. Statistics are per-seed and paired — the median of the per-seed ratios and the count of seeds the arm actually beat, because a real twenty percent shows on most of 256 seeds while drift shows as a trend a rotation destroys. And the combined arms run, in case the two effects are real and additive. ★ The split that can falsify a mechanism. The returned nonce IS the hit distance, so every seed can be labelled by whether its hit fell inside the arm's block. Seeds inside take one launch and run the identical kernel at every arm, so no knob can touch them; seeds outside are the only ones that miss and relaunch. An arm whose gain is the same on both groups is not the knob, it is the procedure. A gain living only in the outside group is a real miss-path effect and owes a mechanism from the kernel before it is priced. --- prover/tests/rpx_grind_bench.rs | 353 ++++++++++++++++++-------------- 1 file changed, 203 insertions(+), 150 deletions(-) diff --git a/prover/tests/rpx_grind_bench.rs b/prover/tests/rpx_grind_bench.rs index dd3414ca5..ca4121603 100644 --- a/prover/tests/rpx_grind_bench.rs +++ b/prover/tests/rpx_grind_bench.rs @@ -1,48 +1,43 @@ //! ★ What one device grind costs, and what the two launch knobs do to it. //! -//! Round 2's lever was ruled from `math-cuda/src/grinding.rs`: one launch sizes -//! its block at 8x the expected hit distance, so "drop the 8 and drop 8x the -//! permutations". The kernels say otherwise — both carry +//! # 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, and the stride walk gives every nonce in the block -//! exactly one owner, so the scan STOPS at the first hit. The executed -//! permutations are `h + stride` at ANY scan factor; the block size is a -//! ceiling the launch never reaches, and `stride = grid x block_dim` is the -//! term that is left. +//! 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. //! //! ```text //! cargo test -p lambda-vm-prover --release --features cuda \ //! --test rpx_grind_bench -- --ignored --nocapture //! ``` //! -//! ONE process, every arm, and that is deliberate: the knobs cache in a -//! `OnceLock`, so reading two settings through the environment would need two -//! processes — two device contexts, two cubin loads, two clock domains. Every -//! arm here runs on the SAME seeds through `generate_nonce_rpx_gpu_at`, so the -//! comparison is PAIRED on identical hit distances and the seed-to-seed spread -//! (which is exponential, and wide) cancels between arms instead of being -//! averaged away. The environment path is exercised too, and asserted to agree. -//! -//! WHAT TO READ, pre-registered: -//! -//! * **the scan arms** (8/4/2/1 at grid 1024) — FLAT means the block size is a -//! ceiling and the ruled lever is dead; halving with the factor means the -//! scan dominates and it is real. The k=1 arm prices the extra round trips -//! (1.58 expected launches per grind against 1.00). -//! * **the grid arms** (256/512/1024/2048/4096 at scan 8) — the sweep runs BOTH -//! ways because the knob has two edges: under-filled, a wider grid raises -//! throughput faster than overshoot and the wall falls; filled, extra blocks -//! only queue and the wider stride is pure added work. `DEVICE FILL` below -//! says which edge the default sits on, READ from the driver. -//! * **`ns/perm`** — the invariant. Time is about `(h/stride + 1)` rounds, so -//! `ms/grind` carries the seed's hit distance while `ns/perm` does not; the -//! grid question is a throughput question and this is the throughput. -//! * **the nonce lists** — identical across every arm, or the search is not -//! scanning contiguously from zero and the geometry is moving the answer. -//! * **the CALIBRATION CONTROL** — this bench's ms/grind at the record posture -//! times the base's 3,428 grinds, against the base's grind wall read from -//! wt14. Out of that window and nothing here describes the block's grind. -//! //! 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 @@ -63,10 +58,7 @@ type RpxGrind = GrindingDigest; /// security parameter, and this measurement is about the launch, not the bits. const GRINDING_FACTOR: u8 = 20; -/// Grinds per arm, all arms on the same seeds. The hit distance is exponential -/// with mean `2^20`, so a single seed says nothing; pairing across arms is what -/// makes the ratios precise, and this many keeps the unpaired mean's standard -/// error near 6%. +/// Grinds per arm, every arm on the same seeds. const RUNS: usize = 256; /// ★ THE CALIBRATION WINDOW, from the block and not from this bench. @@ -80,105 +72,148 @@ 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; -/// One arm's reading. -struct Arm { - knobs: Knobs, - mean_ms: f64, - median_ms: f64, - min_ms: f64, - max_ms: f64, - /// Time divided by the permutations actually executed, `h + stride`. The - /// returned nonce IS `h`, so this is measured rather than modelled. - ns_per_perm: f64, - nonces: Vec, +/// 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 4", + Knobs { + scan: 4, + grid: 1024, + }, + ), + ( + "scan 2", + Knobs { + scan: 2, + grid: 1024, + }, + ), + ( + "scan 1", + Knobs { + scan: 1, + grid: 1024, + }, + ), + ("grid 256", Knobs { scan: 8, grid: 256 }), + ("grid 512", Knobs { scan: 8, grid: 512 }), + ( + "grid 2048", + Knobs { + scan: 8, + grid: 2048, + }, + ), + ( + "grid 4096", + Knobs { + scan: 8, + grid: 4096, + }, + ), + ("scan 1 grid 512", Knobs { scan: 1, grid: 512 }), + ("scan 1 grid 256", Knobs { scan: 1, grid: 256 }), + ( + "control AGAIN", + Knobs { + scan: 8, + grid: 1024, + }, + ), + ] } #[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, which is not what a grind costs in a proof - // that has already done three thousand of them. + // 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 so the knob is shown READ on the - // same path the tree runs use — and its answer tied to the explicit path. + // 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)"); - let explicit_nonce = one_grind(&seeds[0], env_knobs).0; assert_eq!( - env_nonce, explicit_nonce, - "the environment path and the explicit path disagree at {env_knobs:?} — \ - one of them is not the search the other is" + 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:?}"); - let scan_arms: Vec = [8u32, 4, 2, 1] - .into_iter() - .map(|scan| run_arm(Knobs { scan, grid: 1024 }, &seeds)) - .collect(); - let grid_arms: Vec = [256u32, 512, 2048, 4096] - .into_iter() - .map(|grid| run_arm(Knobs { scan: 8, grid }, &seeds)) - .collect(); - - // The record posture is the shared arm: the head of the scan sweep and the - // 1024 point of the grid sweep are the same measurement, not two. - let record = &scan_arms[0]; - - println!("\n=== THE ARMS (paired: every arm ran the same {RUNS} seeds) ==="); + // ── 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!( - "{:<14} {:>10} {:>10} {:>10} {:>10} {:>10} {:>9}", - "knobs", "mean ms", "median", "min", "max", "ns/perm", "vs 8/1024" + "GRIND BENCH: {} arms x {RUNS} seeds, rotated (seed s starts at arm s % {})", + n, n ); - print_arm(record, record); - println!("-- the scan sweep, grid 1024 (pre-registered FLAT if the block is a ceiling) --"); - for arm in &scan_arms[1..] { - print_arm(arm, record); - } - println!("-- the grid sweep, scan 8 (pre-registered to fall on ONE side of 1024, not both) --"); - for arm in &grid_arms { - print_arm(arm, record); - } - // ★ THE NONCE CONTROL. Every arm must have returned the same answer for the - // same seed; the launch geometry cannot move it, and that is why sweeping - // either knob moves no proof byte. - for arm in scan_arms.iter().chain(&grid_arms) { + // ★ 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!( - arm.nonces, record.nonces, - "{:?} returned a different nonce list from the record posture — the \ - launch geometry moved the answer, so the search is not scanning \ - contiguously from zero and nothing else here is quotable", - arm.knobs + 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!( - "\nNONCE CONTROL: all {} arms returned identical nonce lists over {RUNS} seeds", - scan_arms.len() + grid_arms.len() + "{:<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 CALIBRATION CONTROL, said in or out. - let projected = record.mean_ms * BASE_GRINDS / 1000.0; - let inside = (CALIBRATION_LOW_MS..=CALIBRATION_HIGH_MS).contains(&record.mean_ms); + // ── 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: {:.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)", - record.mean_ms + "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 => {}", @@ -188,44 +223,76 @@ fn what_one_grind_costs_at_each_launch_geometry() { "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." + ); } -/// One arm: every seed, one process, explicit knobs. -fn run_arm(knobs: Knobs, seeds: &[[u8; 32]]) -> Arm { +/// 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 mut times = Vec::with_capacity(seeds.len()); - let mut nonces = Vec::with_capacity(seeds.len()); - let mut total_ns = 0.0f64; - let mut total_perms = 0.0f64; - for seed in seeds { - let (nonce, ms) = one_grind(seed, knobs); - total_ns += ms * 1.0e6; - // The returned nonce IS the first valid one, so the permutations the - // launch executed are `nonce + stride` — measured, not modelled. - total_perms += nonce as f64 + stride; - times.push(ms); - nonces.push(nonce); - } - let mean_ms = times.iter().sum::() / times.len() as f64; - let mut sorted = times.clone(); - sorted.sort_by(|a, b| a.partial_cmp(b).expect("no NaN in a timing")); - Arm { - knobs, - mean_ms, - median_ms: sorted[sorted.len() / 2], - min_ms: sorted[0], - max_ms: sorted[sorted.len() - 1], - ns_per_perm: total_ns / total_perms, - nonces, + 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. A device search that failed returns `None` and takes the -/// `expect`, so every timed call is a real device grind. +/// garbage quickly. fn one_grind(seed: &[u8; 32], knobs: Knobs) -> (u64, f64) { let felts = inner_hash_felts::(seed, GRINDING_FACTOR); let started = Instant::now(); @@ -256,8 +323,8 @@ fn print_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", + "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, @@ -269,34 +336,20 @@ fn print_device_fill() { fill.resident_blocks(), fill.resident_threads() ); - println!("{:<8} {:>12} {:>10}", "grid", "fill vs ceiling", "stride"); + println!("{:<8} {:>16} {:>10}", "grid", "fill vs ceiling", "stride"); for grid in [256u32, 512, 1024, 2048, 4096] { println!( - "{:<8} {:>11.2}x {:>10}", + "{:<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", + "⇒ a grid above {} queues: the surplus blocks buy no parallelism and their \ + stride is pure overshoot past the first hit", fill.resident_blocks() ); } } } - -fn print_arm(arm: &Arm, record: &Arm) { - println!( - "scan {:<3} grid {:<5} {:>10.3} {:>10.3} {:>10.3} {:>10.3} {:>10.2} {:>8.3}", - arm.knobs.scan, - arm.knobs.grid, - arm.mean_ms, - arm.median_ms, - arm.min_ms, - arm.max_ms, - arm.ns_per_perm, - arm.mean_ms / record.mean_ms - ); -} From 5125a6fc60ea85f922d0250f4083df0d3cc52aed Mon Sep 17 00:00:00 2001 From: MauroFab Date: Sat, 19 Sep 2026 19:03:43 -0300 Subject: [PATCH 04/17] feat(multilinear): the query openings in four slots, and a count that survives a zero MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit QUERIES is the largest slot in the WHIR chain and, like `open_groups` before it, one number. Four slots partition it — `QUERY_SAMPLE`, `TREE_REBUILD`, `COSET_GATHER`, `OPEN_ASSEMBLE` — counted on EVERY `open_many` call, which is twice per non-final round because `whir_round::prove` opens the current commitment and its successor, and once in the final round. ⛔ The boundary is `open_many`, not `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 over a hundred kilobyte-sized paths and read ~100% every time. What competes with the rebuild is the coset gather, which sits beside `paths()` rather than inside it and would otherwise stay in QUERIES as an unnamed remainder — the same shape as the setup gap the seventh slot was added to name. `queries_other` is bounded above as well as below, and the bound carries an ABSOLUTE allowance beside the relative one. Card-free the fixture's query openings read 0.00 s, and three percent of two milliseconds is below the glue between the windows and below the clock itself, so a purely relative bound would fire on the honest path at the shape the gate actually runs. ★ Arm F is the arm that survives that shape: `rebuild_calls` must equal `2·round_count − chain_count`, every term counted by the run rather than read off the source. The durations vanish card-free; the calls do not. ⛔ And the term is CHAINS, not groups. `stacked_eval::prove` runs one chain per COMMITMENT in the stacked commitment, so a group can open several, and an identity written over groups would have been red on the honest path the first time one did. Its guard is "any of the three counters is nonzero" rather than "the rounds are", because guarding on the rounds alone makes a dropped ROUND counter invisible — the same blind spot the seventh slot opened in arm E. The harness sums the four over the epoch records and prints `tree_rebuild`'s SHARE of the query openings, which is round 3's kill condition: retention removes the rebuilds and nothing else, so if they are not the bulk of QUERIES the lever is dead before any lifetime code is written. The closure line now names arms A-F, because a line that under-names what it checked reads exactly like a check that never ran. ★ The new arm found a defect in the existing fixture on its first run: the "must NOT redden" twin zeroed the QUERIES slot while leaving the four inside it at their honest values, which models four parts summing to more than their whole. The fixture was wrong, not the bound. --- crypto/multilinear/src/whir_chain.rs | 10 + crypto/multilinear/src/whir_commit.rs | 22 +- crypto/multilinear/src/whir_round.rs | 5 + crypto/multilinear/src/whir_split.rs | 479 ++++++++++++++++++- prover/src/lfm/per_table_aggregator_tests.rs | 38 +- 5 files changed, 535 insertions(+), 19 deletions(-) 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..54591fb13 100644 --- a/prover/src/lfm/per_table_aggregator_tests.rs +++ b/prover/src/lfm/per_table_aggregator_tests.rs @@ -7420,6 +7420,42 @@ 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::() + ); println!( " global (in base) wall {g_wall:.1}s ({:.1}%)", pct(g_wall) @@ -7447,7 +7483,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 ); From c766c232796a7675e2b878093444e245d8f38bc5 Mon Sep 17 00:00:00 2001 From: MauroFab Date: Sat, 19 Sep 2026 19:04:02 -0300 Subject: [PATCH 05/17] test(bench): read the noise floor by name, and show the tail that moves the mean MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit ⛔ The launcher read the repeated control's ratio by COLUMN POSITION, and that arm's label is two words, so it read the throughput column instead. It reported the procedure as 327% unstable on a run whose ratio column read 1.000 — a check that could not pass, in a script that reads every other verdict by name. The fix is not a better column index. The bench now prints the floor on its own named line, so nothing downstream has to count spaces to find the truth. And two readings the means still owe. The per-seed median ratios are flat while the means fall, which is a statement about a distribution, so the distribution is now printed: the deciles of the per-seed ratio per arm, and the twenty seeds that move the mean most against the control. Each of those twenty carries its hit distance and its launch count at both arms. The launch count is DERIVED rather than instrumented — the search advances its base by one block per miss and returns on the block containing the hit, so the count is the hit distance over the block plus one, exactly. It is the only quantity that differs between two arms for one seed, which makes it the discriminator: if the twenty are the largest-hit-distance seeds and their launch counts exceed one, the effect lives in the miss-and-relaunch path or in what a long sustained launch costs under the board power limiter, and the card drew its full power on that run. If they are ordinary seeds, neither survives. --- prover/tests/rpx_grind_bench.rs | 81 +++++++++++++++++++++++++++++++++ 1 file changed, 81 insertions(+) diff --git a/prover/tests/rpx_grind_bench.rs b/prover/tests/rpx_grind_bench.rs index ca4121603..362463312 100644 --- a/prover/tests/rpx_grind_bench.rs +++ b/prover/tests/rpx_grind_bench.rs @@ -205,6 +205,87 @@ fn what_one_grind_costs_at_each_launch_geometry() { 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!( + " ⭐ if these are the LARGEST-h seeds and their `arm lch` is > 1, the effect \ + lives in the miss-and-relaunch path or in what a long sustained launch costs \ + under the power limiter — and the card drew its full board power on this run. \ + If they are ordinary-h seeds, neither candidate survives." + ); + // ── THE CALIBRATION CONTROL ──────────────────────────────────────────── let control_mean = mean(×[0]); let projected = control_mean * BASE_GRINDS / 1000.0; From 65327b87f1a8771ae70ade8935bf6de2bb590213 Mon Sep 17 00:00:00 2001 From: MauroFab Date: Sat, 19 Sep 2026 20:05:40 -0300 Subject: [PATCH 06/17] test(bench): sweep the grind's scan factor UPWARD, where the loop bound is the only thing left MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit v4's top-20 killed three of the four candidates for the tail effect. The seeds that move the mean are SMALL-h (258k-512k, inside 2^20), take 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. That rules out the power limiter (these are not the long sustained launches), the miss-and-relaunch path (one launch either way) and the volatile load's per-iteration cost (the same iterations either way). What is left is readable from the code. For a one-launch seed `search` does nothing that scales with `count` — one 8-byte sentinel, a launch at a grid the knob fixes, 8 bytes back, a synchronize — and 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 never observes the atomicMin runs to `count` and wastes in proportion to it. So sweep the knob UP instead of down. Scan 16, 32 and 64 join the arms, and the new COUNT SLOPE section prints each arm's excess over the tightest cap in the sweep, normalised to the control, BESIDE its prediction (k-1)/7. Count-bound reads 2.14 / 4.43 / 9.00 at k = 16 / 32 / 64; saturated reads ~1.00 from k = 8 up. The two branches are a factor of eight apart at k = 64, which no clock ramp, thermal drift, ordering or seed spread produces. The grid pair is run at BOTH caps (grid 4096 at scan 8 and at scan 1) so the same defect can be tested from the block-count side: if wide grids make stragglers worse by contending the atomicMin's line, a tight `count` should mask it. Equal damage at both caps refuses that unification. The verdict is printed on a named line and the section is read by its name, not by column position or section order: `scan 64` also begins a row in THE ARMS and in the DECILE tables, so the launcher anchors the read to the COUNT SLOPE section itself. That is v3's lesson, which cost this file a noise floor that could not pass. No default moves. Scan 16/32/64 exist to make waste visible by exaggerating it and are candidates for nothing; the record posture stays scan 8 / grid 1024. This measures wasted work, never a wrong answer: the nonce control asserts all twelve arms return identical nonce lists before any timing is read. --- prover/tests/rpx_grind_bench.rs | 197 +++++++++++++++++++++++++++++--- 1 file changed, 183 insertions(+), 14 deletions(-) diff --git a/prover/tests/rpx_grind_bench.rs b/prover/tests/rpx_grind_bench.rs index 362463312..f83b2633e 100644 --- a/prover/tests/rpx_grind_bench.rs +++ b/prover/tests/rpx_grind_bench.rs @@ -33,6 +33,52 @@ //! 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 @@ -87,9 +133,9 @@ fn arms() -> Vec<(&'static str, Knobs)> { }, ), ( - "scan 4", + "scan 1", Knobs { - scan: 4, + scan: 1, grid: 1024, }, ), @@ -101,21 +147,34 @@ fn arms() -> Vec<(&'static str, Knobs)> { }, ), ( - "scan 1", + "scan 4", Knobs { - scan: 1, + scan: 4, grid: 1024, }, ), - ("grid 256", Knobs { scan: 8, grid: 256 }), - ("grid 512", Knobs { scan: 8, grid: 512 }), ( - "grid 2048", + "scan 16", Knobs { - scan: 8, - grid: 2048, + 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 { @@ -124,7 +183,13 @@ fn arms() -> Vec<(&'static str, Knobs)> { }, ), ("scan 1 grid 512", Knobs { scan: 1, grid: 512 }), - ("scan 1 grid 256", Knobs { scan: 1, grid: 256 }), + ( + "scan 1 grid 4096", + Knobs { + scan: 1, + grid: 4096, + }, + ), ( "control AGAIN", Knobs { @@ -135,6 +200,26 @@ fn arms() -> Vec<(&'static str, Knobs)> { ] } +/// ★ 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() { @@ -280,10 +365,75 @@ fn what_one_grind_costs_at_each_launch_geometry() { ); } println!( - " ⭐ if these are the LARGEST-h seeds and their `arm lch` is > 1, the effect \ - lives in the miss-and-relaunch path or in what a long sustained launch costs \ - under the power limiter — and the card drew its full board power on this run. \ - If they are ordinary-h seeds, neither candidate survives." + " ⭐ 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 ──────────────────────────────────────────── @@ -316,6 +466,25 @@ fn what_one_grind_costs_at_each_launch_geometry() { 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. From d8b7ca62f72e4be0aefbaaf807f8a6c0f9bdbf2c Mon Sep 17 00:00:00 2001 From: MauroFab Date: Sat, 19 Sep 2026 20:25:02 -0300 Subject: [PATCH 07/17] test(math-cuda): count the permutations a grind launch actually executes MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Stage A swept the scan factor upward and killed the last timing-only hypothesis: threads do not run to the loop bound. The excess saturates above count = 2^23 instead of growing with it, reading 1.05 at scan 16, 32 and 64 against a prediction of 2.14, 4.43 and 9.00. It does not saturate immediately either. In iterations per thread the excess fits T·(1 - e^(-(N-8)/tau)) with T = 1.06 ms and tau = 19 on three independent points, which is the shape of a thread scanning for a bounded TIME after the answer is known rather than to a bound. One iteration at grid 1024 is 131,072 permutations, about 0.56 ms, so 19 iterations is about 10.6 ms -- where the earlier top movers sat. That is a fit, not a reading, and this commit replaces it with a reading. `rpx_grind_search_counted` is the shipped search with three device counters: the permutations its threads ran, the deepest thread's iteration count, and how many threads left by the loop bound rather than by the early exit. So executed - (h + stride) becomes a number per search. The counters reduce inside the warp and hit memory four times a warp rather than once a thread -- 131,072 serialised updates of one L2 line would be the same order as the effect under measurement, and an instrument that manufactures its own signal answers a different question. The twin is device-only. The host KAT compiles this file through a shim that supplies gridDim, blockDim and atomicMin but not the shuffle, and a KAT has no answer to check for a kernel that produces no digest. Its agreement with the shipped kernel is asserted instead where it can be: on the device, seed by seed, on the nonce. Nothing on a proving path launches it, no default moves, and the shipped kernel is not touched. The bench pairs every counted search with a shipped one on the same seed and refuses to draw a conclusion from any arm where the two disagree on milliseconds beyond the noise floor -- a counted kernel with different register pressure has different occupancy and measures a different kernel. Both branches are written into the file before the run: a stale poll means real extra permutations, bounded by time and therefore the same at scan 8 and scan 64, with ran_to_end near zero at scan 64 and nonzero at scan 1; no overrun means the slow launches run the same permutations more slowly and the cost is outside this loop. The cross-arm verdict is computed and named in the test, not left for the log's reader. Sized before the first cubin, by this file's own rule: one more call site for `permute`, not one more inlined copy, so the entry is of order 150-250 PTX lines against a file of 6,241. The shuffles take `unsigned long long` rather than `uint64_t` because the intrinsics have no overload for `unsigned long`. --- crypto/math-cuda/kernels/rpx.cu | 126 +++++++++++ crypto/math-cuda/src/device.rs | 6 + crypto/math-cuda/src/grinding.rs | 100 +++++++++ prover/tests/rpx_grind_counted.rs | 353 ++++++++++++++++++++++++++++++ 4 files changed, 585 insertions(+) create mode 100644 prover/tests/rpx_grind_counted.rs 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/device.rs b/crypto/math-cuda/src/device.rs index b71695a2e..529ae0618 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 @@ -810,6 +815,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 diff --git a/crypto/math-cuda/src/grinding.rs b/crypto/math-cuda/src/grinding.rs index 57272623b..31efb03f3 100644 --- a/crypto/math-cuda/src/grinding.rs +++ b/crypto/math-cuda/src/grinding.rs @@ -420,6 +420,106 @@ fn search(arm: Arm, inner: &[u64; 4], grinding_factor: u8, knobs: Knobs) -> Opti } } +/// ⛔ 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 diff --git a/prover/tests/rpx_grind_counted.rs b/prover/tests/rpx_grind_counted.rs new file mode 100644 index 000000000..e180460bf --- /dev/null +++ b/prover/tests/rpx_grind_counted.rs @@ -0,0 +1,353 @@ +//! ★ 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. Misses + // before the hitting block cost their whole block. + let ideal = (counts.launches - 1) * block + counts.nonce + stride; + 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." + ); +} From 2e6b887bc93849d50c5497daab25fa3bd7b32bd5 Mon Sep 17 00:00:00 2001 From: MauroFab Date: Sat, 19 Sep 2026 20:37:07 -0300 Subject: [PATCH 08/17] perf(math-cuda): keep a commitment's leaf layer, not its tree H4 kept the whole Merkle node array so an opening would not re-hash the leaves, measured it on the card, and lost about 15 s: the retention is one object per commitment IN THE GROUP, because every chain's commitment is built before any query index is drawn, and ten of those put the device at 96% -- after which allocations fail, commits fall back to the host, and the host grows about 1.5 GiB per fallen-back chain. That finding stands and is not edited away. What changes is which object. A tree is 2*num_leaves - 1 nodes; its leaf layer is num_leaves of them, half the bytes -- C * 2^(5-k) at fold width k, a quarter of a base codeword at the production k = 4 where H4 held half of one. And the leaf pass is the expensive part: a leaf absorbs a whole 2^k coset, two permutations on a base codeword and 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. Half the memory for most of the saving is a different trade from the one H4 measured. It can also decline, which H4 could not. The capture asks DeviceReservation::grow -- which already existed, unused, with a doc comment describing exactly this case -- and a refusal costs one leaf pass and nothing else, the behaviour of this file before the change. Allocate, then promise, then give the promise back if the allocation failed, so neither direction leaks. The commit cannot fail because of the cache, so commit_stacked's device attempt cannot start returning None, so the fallback cliff is unreachable rather than unmeasured. 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 the hash that built it; paths() takes the width as a parameter and the cache test opens one codeword at two widths on purpose. Served across widths this would hand back authentication paths that are internally consistent and wrong. The predicate is a free function so it takes unit cases on a machine with no device -- it is the one part of this whose failure is not slowness. Counters diverge where they used to agree: tree_builds counts trees assembled, leaf_hash_calls counts passes actually paid, and the two together assert the retention in both directions. The group-scale memory test is now two-sided against the form -- the layers must be held, and a whole tree must still never be -- where the old one-sided bound sat on its own edge. The base split prints admitted against refused retentions on every run, including the runs that retain nothing, because a refusal path that is silent is indistinguishable from a lever that never fired. No proof byte moves: the same leaves give the same tree, the same root and the same paths. --- crypto/math-cuda/src/whir.rs | 360 +++++++++++++++++-- crypto/math-cuda/tests/whir_tree_cache.rs | 90 ++++- crypto/multilinear/src/stacked_eval.rs | 16 + prover/src/lfm/per_table_aggregator_tests.rs | 30 ++ 4 files changed, 463 insertions(+), 33 deletions(-) diff --git a/crypto/math-cuda/src/whir.rs b/crypto/math-cuda/src/whir.rs index dfb8b8993..e00b1db26 100644 --- a/crypto/math-cuda/src/whir.rs +++ b/crypto/math-cuda/src/whir.rs @@ -5,7 +5,7 @@ //! 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}; use cudarc::driver::{CudaSlice, CudaStream, LaunchConfig, PushKernelArg}; @@ -41,6 +41,122 @@ 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); + +/// 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), + ) +} + +/// 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, +} + +/// 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,6 +173,14 @@ 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. /// @@ -104,40 +228,173 @@ 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, + ) { + 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); + *held = Some(RetainedLeaves { + nodes: copy, + log_folding, + hash, + num_leaves, + bytes, + }); + } + + /// 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 +426,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 +476,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,6 +777,8 @@ fn commit_from( elements: n, base: true, builds: BuildCount::default(), + leaf_passes: BuildCount::default(), + leaves: Arc::new(Mutex::new(None)), room: Arc::new(room), }; let root = codeword.commit(log_folding, hash)?; @@ -722,6 +1026,8 @@ 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(), diff --git a/crypto/math-cuda/tests/whir_tree_cache.rs b/crypto/math-cuda/tests/whir_tree_cache.rs index e3e001c27..1b9fd8e89 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: //! @@ -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" + ); } } @@ -365,20 +403,42 @@ fn a_group_holds_only_its_codewords_before_any_open() { 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); 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.", mib(taken), - mib(4 * codeword_bytes), + mib(expect), mib(bound), mib(tree_bytes), mib(4 * (codeword_bytes + tree_bytes)), ); + 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.", + mib(taken), + mib(floor), + mib(4 * leaf_bytes), + ); // The commitments are alive up to here, which is the whole point: a `drop` // any earlier and the assertion would be about a group that had already @@ -406,11 +466,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, 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/prover/src/lfm/per_table_aggregator_tests.rs b/prover/src/lfm/per_table_aggregator_tests.rs index 54591fb13..c1039015b 100644 --- a/prover/src/lfm/per_table_aggregator_tests.rs +++ b/prover/src/lfm/per_table_aggregator_tests.rs @@ -7456,6 +7456,36 @@ coset_gather {:.2}s ({:.0}%) · open_assemble {:.2}s ({:.0}%) · rebuild_calls { 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 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 {}{}", + mib(asked), + mib(got), + math_cuda::whir::leaf_hash_calls(), + math_cuda::whir::tree_builds(), + 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) From b713aed514136054c3a2fc5337f552e97ad80740 Mon Sep 17 00:00:00 2001 From: MauroFab Date: Sat, 19 Sep 2026 20:42:09 -0300 Subject: [PATCH 09/17] fix(bench): the ideal-work model double-counted the blocks a search missed `executed - ideal` came back NEGATIVE on the two arms where searches miss -- scan 1 read -4.17 strides -- and an overrun is not a quantity that can be negative. The cause is the model, not the counters. `ideal` added `(launches - 1) * block` on the reasoning that a missed block costs its whole `count`. It does, but `nonce` is ABSOLUTE: those nonces are already inside it, so the term counted them twice. The model is `nonce + stride`, full stop -- every nonce below the hit, plus one stride round for the threads that were mid-permutation when it landed. Scan 8 and scan 64 are untouched, because at those block sizes every seed in this bench hits on its first launch and the extra term was zero. Those are the two arms the verdict is written over, so the stage-B reading stands: the overrun is the same at both, and `ran_to_end` is 0 at scan 64. --- prover/tests/rpx_grind_counted.rs | 16 +++++++++++++--- 1 file changed, 13 insertions(+), 3 deletions(-) diff --git a/prover/tests/rpx_grind_counted.rs b/prover/tests/rpx_grind_counted.rs index e180460bf..2948dc3a9 100644 --- a/prover/tests/rpx_grind_counted.rs +++ b/prover/tests/rpx_grind_counted.rs @@ -220,9 +220,19 @@ fn what_the_slow_grind_launches_actually_execute() { ); // The model: every nonce below the hit, plus one stride round for - // the threads that were mid-permutation when it landed. Misses - // before the hitting block cost their whole block. - let ideal = (counts.launches - 1) * block + counts.nonce + stride; + // 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); From ccf0551ba7eb58e23039888be09b93ea016dfa33 Mon Sep 17 00:00:00 2001 From: MauroFab Date: Sat, 19 Sep 2026 20:45:10 -0300 Subject: [PATCH 10/17] test(math-cuda): the process-wide counters become an accounting identity `the_process_wide_counter_tracks_the_same_passes` asserted that a tree and a leaf pass are the same event -- `tree_builds() == leaf_hash_calls()`, and that an opening moves the global counter by one. With the leaf layer retained they are no longer the same event by design, so this test would have reddened on the box for the one reason the gate was not looking for: an invariant that expired when the code under it changed, in a file whose other tests were rewritten and this one was not. It now asserts the relationship that replaced it, in deltas because the three counters are process-wide and diverge on purpose: a commit assembles one tree and pays one pass; an opening assembles a tree and pays NOTHING, recording one saving instead; and over any window, trees == passes + savings. That identity 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 never assembled -- where the old equality could only fail in one. --- crypto/math-cuda/tests/whir_tree_cache.rs | 43 ++++++++++++++++------- 1 file changed, 31 insertions(+), 12 deletions(-) diff --git a/crypto/math-cuda/tests/whir_tree_cache.rs b/crypto/math-cuda/tests/whir_tree_cache.rs index 1b9fd8e89..dd8de10b1 100644 --- a/crypto/math-cuda/tests/whir_tree_cache.rs +++ b/crypto/math-cuda/tests/whir_tree_cache.rs @@ -70,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}; @@ -522,20 +522,39 @@ 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" - ); - assert_eq!( - codeword.tree_builds(), - leaf_hash_calls(), - "with one codeword in flight the two counters must agree" + 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 ); } From 90a38eccdca0c23ca195189b4dc55db9b10d39c8 Mon Sep 17 00:00:00 2001 From: MauroFab Date: Sun, 20 Sep 2026 13:54:04 -0300 Subject: [PATCH 11/17] fix(test): the host KAT punned its result type, and LP64 made it matter MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit `run_grind` declared its result `uint64_t` and handed the address to the kernel as `volatile unsigned long long *`. Those are the same type on Darwin/arm64 and different types of the same width on LP64 glibc, so on Linux the cast type-punned; 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, so layer 8's two checks per vector that expect the SENTINEL passed VACUOUSLY while the two that expect a found nonce failed. Six rows, at every sha back to the gated base c00342c1f, on a target that passes on a clang/arm64 laptop where the two types coincide. Measured on the box at this sha: -O2 6 FAILURE(S) -O2 -fno-strict-aliasing ALL HOST KAT CHECKS PASS -O0 ALL HOST KAT CHECKS PASS None of it was ever a statement about the device grind. This file is a HOST replay of the kernel source through `cuda_host_shim.h` — no nvcc, no cubin, no device — so the defect was in the harness holding the result, not in the kernel it tests. The production path re-validates every device nonce with the host predicate, and the block pins read `host fallbacks 0` throughout. `crypto/math-cuda/tests/host_kat/` holds exactly one instance of the pattern and this is it: the shim's `atomicMin` takes `unsigned long long *` as a parameter, and the kernels' casts there only drop `volatile` from an already-matching type. Test-only; no production code and no proof bytes move. It also unblocks this lineage's CI `host-kat` job, which has been failing since the device grind landed. --- .../math-cuda/tests/host_kat/rpx_host_kat.cpp | 26 ++++++++++++++++++- 1 file changed, 25 insertions(+), 1 deletion(-) 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; From febb5333d796b3d16142b998364115d9daa02f6a Mon Sep 17 00:00:00 2001 From: MauroFab Date: Sun, 20 Sep 2026 13:54:44 -0300 Subject: [PATCH 12/17] test(math-cuda): sample the group guard at one pool state, not two MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit `a_group_holds_only_its_codewords_before_any_open` took `free_before` AFTER a `drain_and_trim()` and `free_after` without one. `free_vram_bytes()` is the driver's count and the device pool is configured to retain all freed blocks, so that difference measured the code PLUS every transient the four commits made. A delta between two samples is about the code only if both are taken at the same pool state, and the test's own comment says exactly that about the first drain. It read 301,989,888 B — nine codewords on a bound of nine, where the model says eight are held. The old one-sided bound was `8 x 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. The slack was never sized against the transients either: `build_tree` allocates `(2L-1)*32` = 64 MiB, twice the bound's 32. The 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, and 5*(cw+tree) = 480 fits that run exactly but dies on the honest one, where 5*(cw+leaf) = 320 != 288. Every peak-demand model misses on BOTH sides (320 and 448 predicted). That is driver suballocation, not an accounting this tree keeps. So the driver's count is sampled at the same pool state on both sides, and the assertion messages now carry the CODE's own number beside it: the delta of `Backend::reserved_bytes()` across the window, the pool's share as the difference, and each held codeword's `reserved_bytes()`/`retained_leaf_bytes()`. A future failure says which of the two it is instead of posing the question. Pre-registered for the next card run: the honest case falls 288 -> ~256 and the whole-tree mutation 480 -> ~384. If the honest case still reads 288, the ninth block is the CODE and this reasoning is wrong. The assertion is not weakened: drained on both sides, four kept layers read ~256 below the 288 bound and four kept TREES still read ~384 above it, so the 128 MiB of separation the bound was designed around is restored rather than spent. Test-only. --- crypto/math-cuda/tests/whir_tree_cache.rs | 67 ++++++++++++++++++++++- 1 file changed, 65 insertions(+), 2 deletions(-) diff --git a/crypto/math-cuda/tests/whir_tree_cache.rs b/crypto/math-cuda/tests/whir_tree_cache.rs index dd8de10b1..9e01527f1 100644 --- a/crypto/math-cuda/tests/whir_tree_cache.rs +++ b/crypto/math-cuda/tests/whir_tree_cache.rs @@ -393,13 +393,47 @@ 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; @@ -417,27 +451,56 @@ fn a_group_holds_only_its_codewords_before_any_open() { 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(", ") + ) + }; assert!( taken < bound, "four unopened commitments took {} MiB from the device. Four codewords \ 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.", + leaf layer is held per commitment.\n {}", mib(taken), 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.", + 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` From 354bc6919b0a90c2e84aa45a8d4f03be7676e702 Mon Sep 17 00:00:00 2001 From: MauroFab Date: Sun, 20 Sep 2026 14:09:15 -0300 Subject: [PATCH 13/17] test(math-cuda): print the group guard's ledger on the honest path too MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The ledger the previous commit added is built inside the two assertion messages, so it appears only when the test FAILS. On the passing path its numbers were inferred rather than read, and what a pass alone establishes is `driver < bound` — pool share under one codeword — and nothing narrower. That is not enough for this guard. The mutation run that keeps a whole node array reads `driver 448 MiB · promised 383 MiB · pool share 64 MiB` even after the symmetric drain: a fragmentation floor of one largest transient, because a best-effort pool trim cannot release a chunk still backing a live allocation. If the honest path sits anywhere near that floor, this test has a margin of a few MiB and will flake, and a green run would never say so. 0 MiB and 31 MiB are the same observation today. So the line is printed unconditionally, under --nocapture, which is how the box gate runs this suite. Once its honest value is known a ceiling can be asserted against it, which is a gate change rather than a test change and is not made here. `promised` is worth having in the log for its own sake: at the mutation it read 383 MiB against a model of 4 x 96.00, exact to the MiB, which is what settled that the excess is the pool rather than a ninth object the code holds. Test-only; no assertion changes, no production code. --- crypto/math-cuda/tests/whir_tree_cache.rs | 16 ++++++++++++++++ 1 file changed, 16 insertions(+) diff --git a/crypto/math-cuda/tests/whir_tree_cache.rs b/crypto/math-cuda/tests/whir_tree_cache.rs index 9e01527f1..29b773a26 100644 --- a/crypto/math-cuda/tests/whir_tree_cache.rs +++ b/crypto/math-cuda/tests/whir_tree_cache.rs @@ -478,6 +478,22 @@ fn a_group_holds_only_its_codewords_before_any_open() { 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 \ From bc6ce9ea281c2e45b37cdfbd07f6b2c378b0d41e Mon Sep 17 00:00:00 2001 From: MauroFab Date: Sun, 20 Sep 2026 15:03:04 -0300 Subject: [PATCH 14/17] feat(math-cuda): count argue-surface device fallbacks (sumcheck/gkr/columns) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The only fallback number the WHIR campaign read was `multilinear::gpu::host_fallbacks()`, which has exactly one caller — the COMMIT path in `whir_chain.rs` — so every `host fallbacks 0` certified that no commitment fell back and said nothing about the per-table argument. wt16 read as a slot-level win for exactly that blind spot: the leaf-layer retention grew `be.reserved`, argue's `reserve` then refused, and its work moved to the host uncounted while `tree_rebuild` alone showed the saving. Add a process-wide counter in `device.rs` beside `reserve` (`device_fallbacks` / `reset_device_fallbacks` / `note_device_fallback`), always compiled and reading zero on a non-cuda build, bumped at the five argue-side `reserve`->`None` sites in math-cuda: `sumcheck.rs` (x3), `gkr.rs` (x1), `columns.rs` (x1). The doc states this scope precisely and records the multilinear residual (`gpu.rs:177`/`:1572`) as a follow-up that needs each caller traced, and why `gpu.rs:1551` — a speculative reserve whose `None` selects a lazy on-device path — must never be counted. Tests: a card test drives the cheapest site (`DeviceColumns::upload`) with the budget fully reserved (an atomic bump, no device memory) and asserts the counter reads one; a card-free unit test covers the counter API; a card-free source-count test asserts the call appears at exactly the five sites, so the four undriven sites fail without four card fixtures. --- crypto/math-cuda/src/columns.rs | 5 +- crypto/math-cuda/src/device.rs | 99 +++++++++++++++++++++++ crypto/math-cuda/src/gkr.rs | 1 + crypto/math-cuda/src/sumcheck.rs | 3 + crypto/math-cuda/tests/whir_tree_cache.rs | 62 ++++++++++++++ 5 files changed, 169 insertions(+), 1 deletion(-) 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 529ae0618..4838d3528 100644 --- a/crypto/math-cuda/src/device.rs +++ b/crypto/math-cuda/src/device.rs @@ -556,6 +556,53 @@ 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); +} + /// Allocates on `stream`, and if the device says no, gives the pool's retained /// blocks back and asks once more. /// @@ -1390,3 +1437,55 @@ impl Backend { Ok(PooledEvent { event: Some(ev) }) } } + +#[cfg(test)] +mod device_fallback_counter_tests { + use super::{device_fallbacks, note_device_fallback, reset_device_fallbacks}; + 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})" + ); + } +} 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/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/tests/whir_tree_cache.rs b/crypto/math-cuda/tests/whir_tree_cache.rs index 29b773a26..3a79a7540 100644 --- a/crypto/math-cuda/tests/whir_tree_cache.rs +++ b/crypto/math-cuda/tests/whir_tree_cache.rs @@ -637,3 +637,65 @@ fn the_process_wide_counter_tracks_the_same_passes() { 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!( + 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. +} From d508faf65c7ffeff547075fa84194a0bd3778977 Mon Sep 17 00:00:00 2001 From: MauroFab Date: Sun, 20 Sep 2026 15:19:07 -0300 Subject: [PATCH 15/17] feat(prover): print commit + device fallbacks at the WHIR tree's whole-run block MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The WHIR tree harness printed no fallback line at all, so the launcher had nothing to gate on and (v9) voided every run by reading a `host fallbacks` line that only the lb-class grind harness prints. Print BOTH fallback surfaces, whole-run scope, at the WHOLE-RUN block of the production-tree test: `commit fallbacks` (`multilinear::gpu::host_fallbacks`, the commit path) and `device fallbacks` (`math_cuda::device::device_fallbacks`, the argue surface added in the previous commit). The launcher (A-tree-whir.v10) refuses a block number unless both read zero and refuses loudly if either line is absent — closing the blind spot wt16 read as a slot-level win, where argue's per-table work fell to the host uncounted. --- prover/src/lfm/per_table_aggregator_tests.rs | 17 +++++++++++++++++ 1 file changed, 17 insertions(+) diff --git a/prover/src/lfm/per_table_aggregator_tests.rs b/prover/src/lfm/per_table_aggregator_tests.rs index c1039015b..a6048bc27 100644 --- a/prover/src/lfm/per_table_aggregator_tests.rs +++ b/prover/src/lfm/per_table_aggregator_tests.rs @@ -8148,6 +8148,23 @@ 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 WHIR tree at FIXTURE scale — the same driver, card-free, on a guest small From 9cd62d716c6e244f55e605d59530d78229481647 Mon Sep 17 00:00:00 2001 From: MauroFab Date: Sun, 20 Sep 2026 15:38:08 -0300 Subject: [PATCH 16/17] feat(math-cuda): live-retained-bytes footprint + be.reserved high-water instruments MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The evictable retention (round 2's fix) and any honest reading of the leaf-layer lever need two numbers the code lacked. RETAIN_BYTES_LIVE: the SIMULTANEOUS retained device footprint — fetch_add on admit, fetch_sub in a new Drop on RetainedLeaves — plus its per-run peak. Unlike the cumulative RETAIN_BYTES_ADMITTED (39,057 MiB in wt16), this rises and falls with the live layers, so it is the ~2 GiB that actually contends with argue for the budget. The peak is printed on the retention line; the instantaneous count is ~0 by the time the base readback prints (the codewords have dropped). reserved_high_water: the peak be.reserved, updated by note_reserved on every rise in Backend::reserve and DeviceReservation::grow. This is the reservation quantity argue's reserve is checked against, which the raw device trace cannot report — never-purge inflates the raw peak above the budget. Printed at the whole-run block. Tests: a card test asserts the live footprint rises with a held layer, the peak and high-water bound it, and — the balance the eviction relies on — the footprint returns to baseline when the codeword drops. Card-free unit tests cover the high-water's monotone max. --- crypto/math-cuda/src/device.rs | 59 +++++++++++++++++++- crypto/math-cuda/src/whir.rs | 46 +++++++++++++++ crypto/math-cuda/tests/whir_tree_cache.rs | 45 +++++++++++++++ prover/src/lfm/per_table_aggregator_tests.rs | 12 +++- 4 files changed, 160 insertions(+), 2 deletions(-) diff --git a/crypto/math-cuda/src/device.rs b/crypto/math-cuda/src/device.rs index 4838d3528..81f9e1efb 100644 --- a/crypto/math-cuda/src/device.rs +++ b/crypto/math-cuda/src/device.rs @@ -478,6 +478,7 @@ impl DeviceReservation { ) { Ok(_) => { self.bytes.fetch_add(extra, Ordering::Relaxed); + note_reserved(held + extra); return true; } Err(seen) => held = seen, @@ -603,6 +604,34 @@ 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); +} + /// Allocates on `stream`, and if the device says no, gives the pool's retained /// blocks back and asks once more. /// @@ -1019,6 +1048,7 @@ impl Backend { Ordering::Relaxed, ) { Ok(_) => { + note_reserved(held + bytes); return Some(DeviceReservation { bytes: AtomicU64::new(bytes), }); @@ -1440,7 +1470,10 @@ impl Backend { #[cfg(test)] mod device_fallback_counter_tests { - use super::{device_fallbacks, note_device_fallback, reset_device_fallbacks}; + 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 @@ -1488,4 +1521,28 @@ mod device_fallback_counter_tests { 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/whir.rs b/crypto/math-cuda/src/whir.rs index e00b1db26..2c6501337 100644 --- a/crypto/math-cuda/src/whir.rs +++ b/crypto/math-cuda/src/whir.rs @@ -70,6 +70,19 @@ 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. @@ -84,6 +97,24 @@ pub fn retention_report() -> (u64, u64, u64, u64, u64, u64) { ) } +/// ★ 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 @@ -108,6 +139,16 @@ struct RetainedLeaves { 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); + } +} + /// 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, @@ -364,6 +405,11 @@ impl DeviceCodeword { } 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, diff --git a/crypto/math-cuda/tests/whir_tree_cache.rs b/crypto/math-cuda/tests/whir_tree_cache.rs index 3a79a7540..58c1ad5df 100644 --- a/crypto/math-cuda/tests/whir_tree_cache.rs +++ b/crypto/math-cuda/tests/whir_tree_cache.rs @@ -699,3 +699,48 @@ fn an_argue_reservation_refusal_bumps_the_device_fallback_counter() { // `_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" + ); +} diff --git a/prover/src/lfm/per_table_aggregator_tests.rs b/prover/src/lfm/per_table_aggregator_tests.rs index a6048bc27..b0a9753b5 100644 --- a/prover/src/lfm/per_table_aggregator_tests.rs +++ b/prover/src/lfm/per_table_aggregator_tests.rs @@ -7470,11 +7470,12 @@ coset_gather {:.2}s ({:.0}%) · open_assemble {:.2}s ({:.0}%) · rebuild_calls { let (admitted, refused, asked, got, headroom, saved) = math_cuda::whir::retention_report(); 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 {}{}", + " 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{}", mib(asked), mib(got), math_cuda::whir::leaf_hash_calls(), math_cuda::whir::tree_builds(), + mib(math_cuda::whir::retained_bytes_peak()), if refused > 0 { format!( " · FIRST REFUSAL at {:.0} MiB of budget headroom", @@ -8165,6 +8166,15 @@ fn the_whir_production_tree_composes_to_a_root() { " 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 From faf8123e372c34357ad6d7fce33c1a6cdd4a4b60 Mon Sep 17 00:00:00 2001 From: MauroFab Date: Sun, 20 Sep 2026 15:54:42 -0300 Subject: [PATCH 17/17] =?UTF-8?q?feat(math-cuda):=20evictable=20leaf-layer?= =?UTF-8?q?=20retention=20=E2=80=94=20argue=20reclaims=20spare=20bytes=20i?= =?UTF-8?q?nstead=20of=20falling=20to=20host?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit wt16 read the leaf-layer retention net-negative at the block: it grew be.reserved during commit and the openings, argue's per-table reserve then refused and fell to the host uncounted (+11.3 s argue vs the -9.45 s tree_rebuild saved). The cause was a single byte budget raced between the retention and argue. Make the retention SUBORDINATE to argue by eviction, so it holds only genuinely spare bytes and gives them back the instant a real caller needs them: - Backend::reserve, on a budget miss, consults an evictor ONCE before returning None and retries after it frees. grow (the retention's own capture) never evicts, so the retention can only ever yield to argue, never displace it. - whir installs the evictor (a fn pointer, via OnceLock) and keeps a registry of Weak handles to each codeword's (leaves mutex, reservation), registered once at first capture and pruned when the codeword drops. Eviction takes a layer out under its own mutex (a concurrent serve sees None and rebuilds -> no UAF; the evicted nodes' free is stream-ordered on the codeword's own stream, cited to cudarc core.rs:776-795 in the code), drops it, and shrinks the reservation -- freeing the BUDGET while never-purge keeps the raw bytes pooled. The evictor never calls reserve/grow (no reentrancy); lock order is registry -> one layer at a time, be.reserved lock-free -- acyclic. LFM_WHIR_RETENTION=0 disables capture entirely, so the ABBA control arm runs off the same binary as the retention arm. RETAIN_EVICTED / RETAIN_BYTES_EVICTED and the live-footprint peak are on the retention print. Tests: a card test fills the budget below a held layer's size and asserts the next reserve SUCCEEDS by evicting it (the layer reads None after) -- the arm that fails when the evictor is disabled. --- crypto/math-cuda/src/device.rs | 46 ++++++ crypto/math-cuda/src/whir.rs | 154 ++++++++++++++++++- crypto/math-cuda/tests/whir_tree_cache.rs | 64 ++++++++ prover/src/lfm/per_table_aggregator_tests.rs | 4 +- 4 files changed, 264 insertions(+), 4 deletions(-) diff --git a/crypto/math-cuda/src/device.rs b/crypto/math-cuda/src/device.rs index 81f9e1efb..9d010151a 100644 --- a/crypto/math-cuda/src/device.rs +++ b/crypto/math-cuda/src/device.rs @@ -632,6 +632,36 @@ 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. /// @@ -1036,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( diff --git a/crypto/math-cuda/src/whir.rs b/crypto/math-cuda/src/whir.rs index 2c6501337..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, Mutex}; +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. /// @@ -149,6 +149,128 @@ impl Drop for RetainedLeaves { } } +/// 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, @@ -228,6 +350,11 @@ pub struct DeviceCodeword { /// 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 { @@ -379,6 +506,15 @@ impl DeviceCodeword { 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; }; @@ -417,6 +553,14 @@ impl DeviceCodeword { 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 @@ -826,6 +970,7 @@ fn commit_from( 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)) @@ -1077,6 +1222,9 @@ pub fn fold_resident( // 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/whir_tree_cache.rs b/crypto/math-cuda/tests/whir_tree_cache.rs index 58c1ad5df..a6494ae49 100644 --- a/crypto/math-cuda/tests/whir_tree_cache.rs +++ b/crypto/math-cuda/tests/whir_tree_cache.rs @@ -744,3 +744,67 @@ fn the_live_footprint_and_reserved_high_water_track_the_retention() { 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/prover/src/lfm/per_table_aggregator_tests.rs b/prover/src/lfm/per_table_aggregator_tests.rs index b0a9753b5..8b523b951 100644 --- a/prover/src/lfm/per_table_aggregator_tests.rs +++ b/prover/src/lfm/per_table_aggregator_tests.rs @@ -7468,14 +7468,16 @@ coset_gather {:.2}s ({:.0}%) · open_assemble {:.2}s ({:.0}%) · rebuild_calls { // 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{}", + " 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",