igneum-network/igneum
1
0
Fork 0
Code Issues Pull requests Releases Activity

Fast-time two-node cache-history case (review B F01, Phase 1 item c): a proof's bytes refused once for context pay later as their own record from the cached facts with no second verify; invalid bytes are refused again from the cache

Browse source 
Three local nodes at fast time; the real proof is one of the chain's own blocks, exported from H1, cut by igneum-prove-export and proven compressed by igneum-prove-host on the box's CPU inside the run, so its statement is H1's native one. Verdict: paid once; H1's verifies.run grew by exactly 2 (the real proof, the invalid bytes); cacheAnswers by at least 2; the context refusal read; the invalid bytes refused again with no new verify. Needs the node's verifies counters (verdict-cache-fix-node, the status JSON) and node and host on the same pair.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
This commit is contained in:
igneum-labs 2026-10-08 20:39:40 +00:00
parent 5144e87c18
commit 7641dd4f9c
8 changed files with 424 additions and 1717 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

296
igneum-pow/src/accept.rs
View file

@ -111,13 +111,6 @@ pub enum Reject {
HotItemSite { site: u8, distinct: u32, evaluations: u32, ratio_milli: u32, top_index: u32, top_count: u32 },
/// (c): output bit `bit` was set in `ones` of 2,048 hashes.
OutputBias { bit: u8, ones: u32 },
/// The reg64 window's liveness rule (the coordinator's spec of 8 October 2026, `check_window_liveness`, a research
/// class, not wired into the acceptance): xoring register `reg` of lane `lane` with a probe word at the start of iteration 0
/// left the iteration's first load address unchanged (`address_changed` false) or the final hash unchanged
/// (`result_changed` false). A window whose address fold reads a subset of the registers is refused here.
DeadWindowRegister { reg: u8, lane: u8, address_changed: bool, result_changed: bool },
/// `check_window_liveness` on a program without the reg64 window.
NotAWindow,
/// (c): the distinct-address sum was `sum`.
DistinctAddresses { sum: u64 },
}
@ -140,13 +133,6 @@ impl std::fmt::Display for Reject {
Reject::LowEntropySite { site, distinct, evaluations, ratio_milli } => write!(f, "(c'') load site {site} read {distinct} distinct word indices over {evaluations} evaluations, {}.{:03} of a uniform source on its window (floor {MIN_DISTINCT_RATIO_V4} at 2^20)", ratio_milli / 1000, ratio_milli % 1000),
Reject::SaturatedSource { site, count } => write!(f, "(c') load site {site} read a saturated source value in {count} of 16384 evaluations (limit 163)"),
Reject::OutputBias { bit, ones } => write!(f, "(c) output bit {bit} set in {ones} of 2048 hashes"),
Reject::DeadWindowRegister { reg, lane, address_changed, result_changed } => write!(
f,
"(reg64 liveness) xoring register {reg} of lane {lane} with a probe word at the start of iteration 0 left the first load address {} and the result {}",
if *address_changed { "changed" } else { "unchanged" },
if *result_changed { "changed" } else { "unchanged" }
),
Reject::NotAWindow => write!(f, "(reg64 liveness) not a reg64 window class"),
Reject::DistinctAddresses { sum } => {
write!(f, "(c) distinct dataset addresses {sum} over 2048 hashes (mean {:.2}, needs above 120 of 128 of the dataset loads)", *sum as f64 / 2048.0)
}
@ -227,116 +213,6 @@ pub fn check_distinct_indices_v4(p: &Program) -> Result<(), Reject> {
distinct_ratio_pass(p, ACCEPT_UNITS_DISTINCT_V4, MIN_DISTINCT_RATIO_V4).map(|_| ())
}
/// The index-bit bias read of class v6 lane 1 (`docs/analysis/class-v6/family-gate.md` section 3.2, the value-level
/// test of adv-cache-2; `docs/design/class-v6-rotating-family.md` section 5.1): a research-class check beside (c''),
/// NOT wired into the acceptance (as a refusal it would redraw about 40 percent of epochs on the plain `load_index`;
/// with the index fold in it is the guard that the fold holds). Per load site, in load order, the one-count of every
/// address bit over the site's `units x 32 x 8` word indices on the rule's closed form (`2^28` words) against n / 2,
/// in sigma (`sigma = sqrt(n) / 2`, 512 at the (c'') sample of 2^20). The free bits of a site are those at or above
/// the width's alignment and below the window's cut (`28 - k`, `k = min(win, 2)`); a bit the window fixes reads 0 or
/// n by construction and is reported as 0 sigma. `Err` when the run hits a lane-constant site. The known-failed set is
/// lane D's: on the plain address, F8's p4, p8, p10, p15, p34 (class v4) and the gate's p212, p225 (class v5) each
/// carry one site over 6 sigma at address bit R or R + 1 (the era's rotation); with the fold every one reads under 6.
pub fn index_bit_sigma(p: &Program, units: usize) -> Result<Vec<[f64; ACCEPT_DATASET_LOG2 as usize]>, Reject> {
const BITS: usize = ACCEPT_DATASET_LOG2 as usize;
let indices = site_indices(p, units)?;
let n = (units * LANES * ITERATIONS) as f64;
let sigma = n.sqrt() / 2.0;
let mut out = Vec::with_capacity(indices.len());
let loads: Vec<&Instr> = p.instrs.iter().filter(|i| i.op.is_load()).collect();
for (site, ix) in indices.iter().enumerate() {
let ins = loads[site];
let k = (ins.win as u32).min(ACCEPT_DATASET_LOG2.saturating_sub(26));
let low = (ins.width as u32).trailing_zeros();
let high = ACCEPT_DATASET_LOG2 - k;
let mut ones = [0u64; BITS];
for &x in ix {
for (b, o) in ones.iter_mut().enumerate() {
*o += ((x >> b) & 1) as u64;
}
}
let mut z = [0.0f64; BITS];
for b in 0..BITS {
if (b as u32) >= low && (b as u32) < high {
z[b] = (ones[b] as f64 - n / 2.0) / sigma;
}
}
out.push(z);
}
Ok(out)
}
/// The largest `|sigma|` of [`index_bit_sigma`] with its site and bit: `(sigma, site, bit)`, signed as read.
pub fn index_bit_sigma_max(p: &Program, units: usize) -> Result<(f64, usize, usize), Reject> {
let z = index_bit_sigma(p, units)?;
let mut best = (0.0f64, 0usize, 0usize);
for (site, row) in z.iter().enumerate() {
for (bit, &v) in row.iter().enumerate() {
if v.abs() > best.0.abs() {
best = (v, site, bit);
}
}
}
Ok(best)
}
/// The reg64 window's liveness rule (the coordinator's spec, 8 October 2026; a research class, beside the census,
/// NOT wired into the acceptance): the window's state must stay live across the dependent memory chain, not only
/// inside an arithmetic block. For each of the 64 registers in turn, the register is xored with each of two
/// seed-derived pseudo-random words ([`REG64_PROBE_WORDS`]) in every lane of unit 0 (the seed's first acceptance base nonce) at the start of iteration 0, on the closed-form dataset of
/// the acceptance; the program passes when, in every lane, the iteration's first load reads a different word index
/// and the final hash differs. A window whose load addresses read a subset of the registers (the arithmetic-only
/// form, where a load's address is its own window's register) is refused at the first register outside that
/// subset; the full-chain form (every address source `src ^ m`, `m` the rotate-xor chain over the 63 other
/// registers) passes, both patterns moving the address and the result in every lane. Two pseudo-random words,
/// not the complement and not one bit: the chain is linear (xor and rotate), so a pattern that enters it twice
/// through a copy made by the program (the pinned draw's instruction 1, `r15 ^= r8`) cancels when the two copies
/// agree after their rotations, which the complement always does (all ones under any rotation) and a single bit
/// does whenever the rotations agree mod 32; a one-bit flip is also lost through a carry before the first load.
/// A dead register fails both words always; a live one fails both with probability about 2^-60. A program without the window is `Reject::NotAWindow`.
pub const REG64_PROBE_WORDS: usize = 2;
/// The two probe words of register `reg` under the program's seed: splitmix32 of the seed's first word, the
/// register and the word index, never 0, never all ones, never a single bit (the patterns a linear fold can lose).
pub fn reg64_probe_words(seed: &[u32; 8], reg: usize) -> [u32; REG64_PROBE_WORDS] {
let mut out = [0u32; REG64_PROBE_WORDS];
for (j, w) in out.iter_mut().enumerate() {
let mut x = splitmix32(seed[0] ^ (reg as u32).wrapping_mul(0x9e3779b9) ^ (j as u32 + 1).wrapping_mul(0x85ebca6b));
while x == 0 || x == u32::MAX || x.is_power_of_two() {
x = splitmix32(x.wrapping_add(0x6c62272e));
}
*w = x;
}
out
}
pub fn check_window_liveness(p: &Program) -> Result<(), Reject> {
use crate::verify::{interpret_warp_probe, Probe};
if !p.class.reg64 {
return Err(Reject::NotAWindow);
}
let shape = crate::memhard::Shape::for_class(&p.class);
let ds = crate::verify::DatasetSource::from_key_shape(p.seed, crate::verify::DatasetMode::ClosedForm, ACCEPT_DATASET_LOG2, shape);
let base = accept_base_nonces_n(&p.seed, 1)[0];
let mut plain = Probe::default();
let plain_out = interpret_warp_probe(p, &p.seed, base, &ds, &mut plain);
assert!(plain.seen_first_load, "a program has a load in every iteration");
for reg in 0..p.registers() {
for word in reg64_probe_words(&p.seed, reg) {
let mut pr = Probe { flip: Some((reg, word)), ..Default::default() };
let out = interpret_warp_probe(p, &p.seed, base, &ds, &mut pr);
for lane in 0..LANES {
let address_changed = pr.first_load_idx[lane] != plain.first_load_idx[lane];
let result_changed = out.hashes[lane] != plain_out.hashes[lane];
if !address_changed || !result_changed {
return Err(Reject::DeadWindowRegister { reg: reg as u8, lane: lane as u8, address_changed, result_changed });
}
}
}
}
Ok(())
}
/// One ratio pass over `units`: every load site's distinct word indices against the uniform expectation on its
/// window (`N - N^2 / 2W`, the window `2^28 >> min(win, 2)` words of the closed-form dataset), `Err` at the first
/// site under `floor`, else the minimum ratio and its site.
@ -368,10 +244,10 @@ pub fn site_window_words(ins: &Instr) -> u64 {
(1u64 << ACCEPT_DATASET_LOG2) >> (ins.win as u64).min(2)
}
/// One interpreter run over `units` units of the seed's acceptance stream with every load site's word indices kept,
/// in evaluation order: one vector per site in load order (the ratio pass sorts them; the index-bit read of class v6
/// lane 1 counts bits over them).
pub fn site_indices(p: &Program, units: usize) -> Result<Vec<Vec<u32>>, Reject> {
/// One interpreter run over `units` units with every load site's word indices kept, then per site the sorted
/// run lengths: [`SiteIndexStats`] per site in load order. The ratio pass (c'') and the hot-item rule (c''') read
/// the same run, so class v5 pays the sort once.
pub fn site_index_stats(p: &Program, units: usize) -> Result<Vec<SiteIndexStats>, Reject> {
let loads = p.loads_per_hash();
let sites = loads / ITERATIONS;
let mut acc = Acc {
@ -388,16 +264,8 @@ pub fn site_indices(p: &Program, units: usize) -> Result<Vec<Vec<u32>>, Reject>
for (unit, &base) in accept_base_nonces_n(&p.seed, units).iter().enumerate() {
run_unit(p, unit, base, &mut acc, &mut lane_addrs)?;
}
Ok(acc.indices.take().unwrap())
}
/// One interpreter run over `units` units with every load site's word indices kept, then per site the sorted
/// run lengths: [`SiteIndexStats`] per site in load order. The ratio pass (c'') and the hot-item rule (c''') read
/// the same run, so class v5 pays the sort once.
pub fn site_index_stats(p: &Program, units: usize) -> Result<Vec<SiteIndexStats>, Reject> {
let mut indices = site_indices(p, units)?;
let mut out = Vec::with_capacity(indices.len());
for ix in indices.iter_mut() {
let mut out = Vec::with_capacity(sites);
for ix in acc.indices.take().unwrap().iter_mut() {
ix.sort_unstable();
let mut st = SiteIndexStats { distinct: 0, pairs: 0, top_index: 0, top_count: 0 };
let mut i = 0;
@ -486,10 +354,8 @@ pub fn check_indices_v5(p: &Program) -> Result<(), Reject> {
/// the era set aside): the shape the sub-version 2 rules (a') and (c') apply to, on every draw path.
pub fn is_class_v4_shape(class: &LoadClass) -> bool {
matches!(class.shadow, Some(ShadowClass { instrs: V4_SHADOW_INSTRS, .. }))
// class v5 (docs/design/class-v5-stored-state.md) is judged under the same rules: its state flag is set aside;
// class v6 lane 1's index fold and re-weight table are set aside too (the address path and the op table are
// not the shape)
&& LoadClass { era: None, shadow: None, state: false, fold: false, rw: 0, ..*class } == LoadClass { shadow: None, ..V4_CLASS }
// class v5 (docs/design/class-v5-stored-state.md) is judged under the same rules: its state flag is set aside
&& LoadClass { era: None, shadow: None, state: false, ..*class } == LoadClass { shadow: None, ..V4_CLASS }
}
/// One pass of the dataflow freshness over the base program then the shadow block (the order of one iteration),
@ -1017,152 +883,6 @@ mod tests {
assert_eq!(v5_under, 0);
}
/// The (c''') reading on a listed set of chain-shaped programs (run by hand on a box through the lease): the file
/// `IGNEUM_V5_PROGRAMS` holds one program per line, `<epoch hex or label> <era hex or label> <attempt> [tag...]`
/// (a hex field of 64 characters is the 32 bytes; anything else is a label whose `seed_words_from_bytes` words are
/// the bytes, as the f8 and adv label spaces derive theirs). Per line: the class v4 program at that attempt (its id,
/// minimum site and ratio at the 2^20 sample), the class v5 verdict at the same attempt, and the class v5 draw's
/// attempt. The coordinator's question of 7 October 2026, 23:3x UK: whether the drawn-era programs adv-cache-2 read
/// as biased (R from 3 to 22) sit under the 0.995 floor.
#[test]
#[ignore]
fn v5_floor_on_listed_programs() {
use crate::generator::{V5_CLASS, V3_ALLOWED};
use crate::seed::seed_words_from_bytes;
let path = std::env::var("IGNEUM_V5_PROGRAMS").expect("IGNEUM_V5_PROGRAMS=<file>");
let text = std::fs::read_to_string(&path).expect("the programs file");
let bytes_of = |f: &str| -> Vec<u8> {
if f.len() == 64 && f.chars().all(|c| c.is_ascii_hexdigit()) {
crate::bind::unhex(f).unwrap()
} else {
seed_words_from_bytes(f.as_bytes()).iter().flat_map(|x| x.to_le_bytes()).collect()
}
};
let mut refused = 0;
let mut rows = 0;
for line in text.lines() {
let line = line.trim();
if line.is_empty() || line.starts_with('#') {
continue;
}
let f: Vec<&str> = line.split_whitespace().collect();
if f.len() < 3 {
println!("skip (needs epoch era attempt): {line}");
continue;
}
let (epoch, era) = (bytes_of(f[0]), bytes_of(f[1]));
// "draw" = the class v4 chain draw's own attempt for this seed (the attempt adv-cache-2's table shows)
let attempt = if f[2] == "draw" { f8_draw(&epoch, &era, V4_CLASS).attempt } else { f[2].parse::<u32>().expect("attempt") };
let tag = f[3..].join(" ");
let label = format!("igneum-epoch/{}", epoch.iter().map(|b| format!("{b:02x}")).collect::<String>());
let v4 = candidate_class(&label, &epoch, attempt, LoadClass::era(V4_CLASS, &era, &V3_ALLOWED));
let st = site_index_stats(&v4, ACCEPT_UNITS_DISTINCT_V4).unwrap();
let (min, site) = distinct_ratio_on(&v4, &st, ACCEPT_UNITS_DISTINCT_V4, 0.0).unwrap();
let v4_verdict = check(&v4).is_ok();
let v5 = candidate_class(&label, &epoch, attempt, LoadClass::era(V5_CLASS, &era, &V3_ALLOWED));
let v5_verdict = check(&v5);
let drawn = f8_draw(&epoch, &era, V5_CLASS);
let era_params = v4.class.era.as_ref().map(|e| e.stride_rot).unwrap_or(0);
rows += 1;
if v5_verdict.is_err() {
refused += 1;
}
println!(
"listed: id {:016x} attempt {attempt} R {era_params} | v4 {} min site {site} ratio {min:.4} (most read {:#x} x{}) | v5 at this attempt {} | v5 draw attempt {} id {:016x} | {tag}",
crate::generator::program_id(crate::generator::GENERATOR_VERSION_V4, &v4.seed, attempt),
if v4_verdict { "accepted" } else { "REJECTED" },
st[site].top_index,
st[site].top_count,
match &v5_verdict { Ok(_) => "accepted".to_string(), Err(r) => format!("REFUSED {r}") },
drawn.attempt,
drawn.program_id()
);
}
println!("listed: {refused} of {rows} refused under class v5 at the listed attempt");
}
/// Class v5's attempts census (run by hand on a box through the lease: `IGNEUM_V5_CENSUS_SEEDS` seeds of the f8 label
/// space from `IGNEUM_V5_CENSUS_FROM`, `IGNEUM_V5_CENSUS_THREADS` threads): every candidate of the class v5 chain draw
/// with the FIRST failing part of the rule, so spec 1.4.7's census paragraph states class v5's own per-candidate
/// rejection and its split by part ((a), (b), (a'), (c), (c'), (c''), (c''')), the exhaustions (must be 0) and
/// P(256 consecutive rejections) = r^256, in the form the crypto lane's attempts census gives for sub-version 3
/// (per-candidate rejection 0.68, (a') 83.3 percent of rejections).
#[test]
#[ignore]
fn v5_attempts_census() {
use crate::generator::{V5_CLASS, V3_ALLOWED};
use std::sync::atomic::{AtomicU32, Ordering};
let seeds: u32 = std::env::var("IGNEUM_V5_CENSUS_SEEDS").ok().and_then(|v| v.parse().ok()).unwrap_or(1000);
let from: u32 = std::env::var("IGNEUM_V5_CENSUS_FROM").ok().and_then(|v| v.parse().ok()).unwrap_or(0);
let threads: usize = std::env::var("IGNEUM_V5_CENSUS_THREADS").ok().and_then(|v| v.parse().ok()).unwrap_or(48);
let next = AtomicU32::new(from);
let t0 = std::time::Instant::now();
let part_of = |r: &Reject| -> &'static str {
match r {
Reject::StaleLoadSource { .. } => "(a) stale load",
Reject::NoInjectingWrite { .. } => "(b) no injecting write",
Reject::UnfreshLoadSource { .. } => "(a') unfresh source",
Reject::ConstantBit { .. } => "(c) constant bit",
Reject::LaneConstantSite { .. } => "(c) lane-constant site",
Reject::Saturated { .. } => "(c) saturated",
Reject::OutputBias { .. } => "(c) output bias",
Reject::DeadWindowRegister { .. } => "(reg64 liveness) dead window register",
Reject::NotAWindow => "(reg64 liveness) not a window",
Reject::DistinctAddresses { .. } => "(c) distinct addresses",
Reject::SaturatedSource { .. } => "(c') saturated source",
Reject::RepeatedSource { .. } => "(c'') repeated source",
Reject::LowEntropySite { .. } => "(c'') low-entropy site",
Reject::HotItemSite { .. } => "(c''') hot-item site",
}
};
// per seed: (k, accepted attempt or None, the parts of every rejected candidate)
let rows: Vec<(u32, Option<u32>, Vec<&'static str>)> = std::thread::scope(|sc| {
let hs: Vec<_> = (0..threads).map(|_| sc.spawn(|| {
let mut out = Vec::new();
loop {
let k = next.fetch_add(1, Ordering::Relaxed);
if k >= from + seeds {
break out;
}
let (epoch, era) = f8_seed(k);
let class = LoadClass::era(V5_CLASS, &era, &V3_ALLOWED);
let label = f8_label(&epoch);
let mut parts = Vec::new();
let mut accepted = None;
for attempt in 0..crate::generator::MAX_ATTEMPTS_V4 {
let c = candidate_class(&label, &epoch, attempt, class);
match check(&c) {
Ok(_) => { accepted = Some(attempt); break; }
Err(r) => parts.push(part_of(&r)),
}
}
out.push((k, accepted, parts));
}
})).collect();
let mut rows: Vec<_> = hs.into_iter().flat_map(|h| h.join().unwrap()).collect();
rows.sort_by_key(|r| r.0);
rows
});
let rejected: usize = rows.iter().map(|r| r.2.len()).sum();
let candidates = rejected + rows.iter().filter(|r| r.1.is_some()).count();
let exhausted = rows.iter().filter(|r| r.1.is_none()).count();
let r = rejected as f64 / candidates as f64;
let mut by_part = std::collections::BTreeMap::new();
for row in &rows {
for p in &row.2 {
*by_part.entry(*p).or_insert(0usize) += 1;
}
}
println!("v5_attempts_census: {} seeds {from}..{} in {:.0} s on {threads} threads", rows.len(), from + seeds, t0.elapsed().as_secs_f64());
println!("candidates {candidates}, rejected {rejected}, per-candidate rejection {r:.4}, exhaustions {exhausted}, P(256 consecutive) {:.3e}", r.powf(256.0));
for (p, n) in &by_part {
println!(" first failing part {p}: {n} ({:.1} percent of rejections, {:.1} percent of candidates)", *n as f64 * 100.0 / rejected as f64, *n as f64 * 100.0 / candidates as f64);
}
let mean = rows.iter().filter_map(|r| r.1).map(|a| a as f64).sum::<f64>() / rows.iter().filter(|r| r.1.is_some()).count() as f64;
println!("accepted attempt mean {mean:.3}");
assert_eq!(exhausted, 0, "no class v5 seed exhausts its attempts");
}
#[test]
fn distinct_bound_scales_with_the_load_count() {
assert_eq!(min_distinct_sum(128), MIN_DISTINCT_SUM);

452
igneum-pow/src/emit.rs
View file

@ -16,112 +16,25 @@ use crate::memhard::{
CHACHA_ROUNDS, CHACHA_SIGMA, HOT_TAG, ITEM_ROUNDS,
};
use crate::seed::SplitMix64;
use crate::verify::{window, window32, DatasetGeom, DatasetMode, DatasetSource, Epoch, DEFAULT_DATASET_LOG2, FOLD_MUL, FOLD_ROT, INDEX_FOLD_SHIFT};
use crate::verify::{window, DatasetMode, DatasetSource, Epoch, DEFAULT_DATASET_LOG2, FOLD_MUL, FOLD_ROT};
/// The index expression of a dataset load (era layout, `docs/plans/era-layout.md` section 1.3). For every class
/// without an era it is the lottery hash's `rN & MASK`; for an era program it is the one form
/// `((rotl_imm(rN * M, R) & WM) | OFF) & MASK` with the site's window constants at the pack's dataset size.
/// Under the multiply-shift geometry (research class ds55, `--dataset-words`) the AND is the dialect's high
/// multiply by the word count: `mulhi(rN, N)` and `mulhi(((rotl_imm(rN * M, R) & WM32) | OFF32), N)` with the
/// window in the source space (`verify::window32`).
///
/// Class v6 lane 1 (`EraParams::fold`, `verify::stride`): the rotation's argument becomes `fold16_(rN * M)`, the
/// kernel's `y ^ (y >> 16)` helper, in every dialect; the text of every era without the fold is unchanged.
fn load_index_expr(dialect: CoreDialect, era: Option<&EraParams>, ins: &Instr, a: &str, geom: DatasetGeom) -> String {
fn load_index_expr(dialect: CoreDialect, era: Option<&EraParams>, ins: &Instr, a: &str, dataset_log2: u32) -> String {
let mask_name = match dialect {
CoreDialect::Metal => "MASK",
_ => "mask",
};
if geom.mulshift {
let (mulhi, words) = mulhi_name(dialect);
return match era {
None => format!("{mulhi}({a}, {words})"),
Some(e) => {
let (wm, off) = window32(ins, geom.log2);
format!("{mulhi}(((rotl_imm({}, {}u) & {}) | {}), {words})", stride_product_expr(e, a), e.stride_rot, hex(wm), hex(off))
}
};
}
let dataset_log2 = geom.log2;
match era {
None => format!("{a} & {mask_name}"),
Some(e) => {
let (wm, off) = window(ins, mask_for(dataset_log2), dataset_log2);
format!("((rotl_imm({}, {}u) & {}) | {}) & {mask_name}", stride_product_expr(e, a), e.stride_rot, hex(wm), hex(off))
format!("((rotl_imm({a} * {}, {}u) & {}) | {}) & {mask_name}", hex(e.stride_mul), e.stride_rot, hex(wm), hex(off))
}
}
}
/// The product the era's rotation takes: `rN * M`, or `fold16_(rN * M)` under the index fold of class v6 lane 1.
fn stride_product_expr(e: &EraParams, a: &str) -> String {
if e.fold {
format!("fold16_({a} * {})", hex(e.stride_mul))
} else {
format!("{a} * {}", hex(e.stride_mul))
}
}
/// The `fold16_` helper of a kernel under the index fold (empty for every other program, so every pinned pack keeps
/// its text): `y ^ (y >> 16)`, one xor and one shift on the address path, the same in the three dialects.
fn fold_helper_lines(dialect: CoreDialect, p: &Program) -> String {
if !p.class.era.map(|e| e.fold).unwrap_or(false) {
return String::new();
}
let (qual, u) = match dialect {
CoreDialect::Metal => ("inline", "uint"),
CoreDialect::Cuda => ("__device__ __forceinline__", "uint32_t"),
CoreDialect::OpenCl => ("static inline", "uint"),
};
format!(
"// Class v6 lane 1, the index fold (8 October 2026, docs/design/class-v6-rotating-family.md section 2; a research class, NOT the lottery\n\
// hash): every era load address rotates fold16_(src * STRIDE_MUL) instead of the bare product, y ^ (y >> {}), so no era's rotation lands\n\
// a biased product bit (bit 0 of an odd product is bit 0 of src; bit 1 is set at 3/8) on an address bit.\n\
{qual} {u} fold16_({u} y) {{ return y ^ (y >> {}u); }}\n",
INDEX_FOLD_SHIFT, INDEX_FOLD_SHIFT
)
}
/// The dialect's high 32-bit multiply and the name of the word-count constant (the multiply-shift geometry).
fn mulhi_name(dialect: CoreDialect) -> (&'static str, &'static str) {
match dialect {
CoreDialect::Metal => ("mulhi", "DS_WORDS"),
CoreDialect::Cuda => ("__umulhi", "IGNEUM_DS_WORDS"),
CoreDialect::OpenCl => ("mul_hi", "IGNEUM_DS_WORDS"),
}
}
/// The word-count lines of a kernel under the multiply-shift geometry (empty under the mask path, so every pinned
/// pack keeps its text): the define the load expressions read, and the comment that says what moved.
fn ds_words_lines(dialect: CoreDialect, geom: DatasetGeom) -> String {
if !geom.mulshift {
return String::new();
}
let (mulhi, words) = mulhi_name(dialect);
let mut s = String::new();
s.push_str(&format!(
"// Research class ds55 (8 October 2026, NOT the lottery hash): the dataset holds {} words ({} items, {} bytes),\n",
geom.words,
geom.items(),
geom.bytes()
));
s.push_str("// not a power of two. Every load address is the multiply-shift range reduction of spec 01 section 1.13.3,\n");
s.push_str(&format!("// idx = (src * {words}) >> 32 in 64 bits ({mulhi}), in place of src & mask; the mask argument is not read by a load.\n"));
s.push_str(&format!("#define {words} {}\n", hex(geom.words as u32)));
s
}
/// The warp-coalesced load's base expression (lever b, `wload`): lane 0's register range-reduced and aligned
/// down to 32 words.
fn wload_base_expr(dialect: CoreDialect, bcast: &str, geom: DatasetGeom) -> String {
if geom.mulshift {
let (mulhi, words) = mulhi_name(dialect);
format!("({mulhi}({bcast}, {words}) & ~31u) + lane")
} else {
let wmask = if dialect == CoreDialect::Metal { "WMASK" } else { "wmask" };
format!("({bcast} & {wmask}) + lane")
}
}
/// The era lines of program.h (empty without an era).
fn era_header_lines(p: &Program) -> String {
let Some(e) = p.class.era else { return String::new() };
@ -138,10 +51,6 @@ fn era_header_lines(p: &Program) -> String {
s.push_str(&format!("#define IGNEUM_ERA_STRIDE_ROT {}\n", e.stride_rot));
s.push_str(&format!("#define IGNEUM_ERA_INTERLEAVE {{ {}, {}, {}, {} }}\n", e.pos[0], e.pos[1], e.pos[2], e.pos[3]));
s.push_str(&format!("#define IGNEUM_ERA_WINDOWS {}\n", jstr(&era_windows(p))));
if e.fold {
s.push_str(&format!("// Class v6 lane 1 (research): the product is folded before the rotation, y = src * STRIDE_MUL; y ^= y >> {INDEX_FOLD_SHIFT}; y = rotl(y, STRIDE_ROT).\n"));
s.push_str(&format!("#define IGNEUM_ERA_INDEX_FOLD {INDEX_FOLD_SHIFT}\n"));
}
s
}
@ -202,7 +111,7 @@ enum WideSource {
/// same shape in the three dialects: the vector loads differ (`uint4` pointer on Metal and CUDA, `vload4` on
/// OpenCL C 1.2). For `width == 1` the caller emits the lottery hash's one-word form instead.
fn wide_load_stmt(dialect: CoreDialect, d: &str, idx: &str, width: u8, src: WideSource, closed: Option<(u32, u32)>) -> String {
debug_assert!(width == 4 || width == 8 || width == 16);
debug_assert!(width == 4 || width == 16);
let (u, base_ptr) = match dialect {
CoreDialect::Metal => ("uint", "dataset"),
CoreDialect::Cuda => ("uint32_t", "ds"),
@ -315,14 +224,6 @@ fn class_header_lines(p: &Program) -> String {
}
s.push_str(&format!("#define IGNEUM_LOAD_CLASS {}
", jstr(&p.class.name())));
if p.class.rw != 0 {
// class v6 lane 1: the op roll's table (program.json "op_weights" has the weights)
let (w, sum) = p.class.nonload_weights();
s.push_str(&format!("// Class v6 lane 1 (research): the op roll draws from re-weight table rw{} (sum {sum}) in place of the plain table (sum 75): {}.
", p.class.rw, w.iter().map(|(o, n)| format!("{} {n}", o.name())).collect::<Vec<_>>().join(", ")));
s.push_str(&format!("#define IGNEUM_OP_WEIGHTS_TABLE {}
", p.class.rw));
}
if p.class.mixer_mult != 1 || p.class.growth {
s.push_str(&format!("#define IGNEUM_CLASS_MIXER_MULT {}
", p.class.mixer_mult));
@ -337,13 +238,6 @@ fn class_header_lines(p: &Program) -> String {
s.push_str(&format!("#define IGNEUM_CLASS_DERIVE_LEN {}
", p.class.derive_len));
}
if p.class.reg64 {
s.push_str("// reg64 (the hash lane's measurement, 8 October 2026, a research class): 64 live registers per lane; r8..r63 = r[k & 7] * 0x9e3779b9 + k;\n");
s.push_str("// each drawn instruction i runs on window 0 (register field + 8 * (i % 4)) then on window 1 (+32); both windows fold into r0..r7 by xor before the hash fold.\n");
s.push_str("#define IGNEUM_REG64 1\n");
s.push_str(&format!("#define IGNEUM_REGISTERS {}\n", p.registers()));
s.push_str(&format!("#define IGNEUM_REG64_ADDRESS_MIX {} // 1: every load's address source is src ^ m, m the rotate-xor chain (m = first; m = rotl(m, 1) ^ next) over the 63 registers other than src, in index order (the full chain)\n", p.address_mix() as u8));
}
s.push_str(&format!("#define IGNEUM_LOAD_SLOTS {}
", p.class.load_slots));
s.push_str(&format!("#define IGNEUM_LOAD_MIX {{ {}, {}, {} }}
@ -967,34 +861,23 @@ const DS_ELEM_BODY: &str = " x *= 0x9E3779B1u; x ^= x >> 15;\n x += d1;\n
/// The Metal hash kernel (`generateMSL`, program.metal).
pub fn metal_program(p: &Program, dataset_log2: u32, source: LoadSource) -> String {
metal_program_impl(p, DatasetGeom::pow2(dataset_log2), source, false)
}
/// [`metal_program`] at a dataset geometry (the multiply-shift sizes of `--dataset-words`).
pub fn metal_program_geom(p: &Program, geom: DatasetGeom, source: LoadSource) -> String {
metal_program_impl(p, geom, source, false)
metal_program_impl(p, dataset_log2, source, false)
}
/// The header-bound Metal kernel (`program_bound.metal`, serve mode of proto-metal): `igneum_hash_bound` reads its
/// init words `I` from `constant uint* initw [[buffer(3)]]` (`bind::block_init_words`) instead of `SEEDW`. Same
/// instruction text as `igneum_hash`. Stored dataset only.
pub fn metal_program_bound(p: &Program, dataset_log2: u32) -> String {
metal_program_impl(p, DatasetGeom::pow2(dataset_log2), LoadSource::Stored, true)
metal_program_impl(p, dataset_log2, LoadSource::Stored, true)
}
/// [`metal_program_bound`] at a dataset geometry.
pub fn metal_program_bound_geom(p: &Program, geom: DatasetGeom) -> String {
metal_program_impl(p, geom, LoadSource::Stored, true)
}
fn metal_program_impl(p: &Program, geom: DatasetGeom, source: LoadSource, bound: bool) -> String {
let mask = geom.mask();
fn metal_program_impl(p: &Program, dataset_log2: u32, source: LoadSource, bound: bool) -> String {
let mask = mask_for(dataset_log2);
let mut s = String::with_capacity(5000);
s.push_str("#include <metal_stdlib>\n");
s.push_str("using namespace metal;\n");
s.push('\n');
s.push_str(&format!("#define MASK {}\n", hex(mask)));
s.push_str(&ds_words_lines(CoreDialect::Metal, geom));
s.push_str(&hot_define(p));
s.push_str(&format!("constant uint SEEDW[8] = {{ {} }};\n", join_hex(&p.seed)));
s.push('\n');
@ -1005,13 +888,12 @@ fn metal_program_impl(p: &Program, geom: DatasetGeom, source: LoadSource, bound:
s.push_str(" return x;\n");
s.push_str("}\n");
s.push_str("inline uint rotl_imm(uint x, uint n) { return (x << n) | (x >> (32u - n)); } // n in 1..31\n");
s.push_str(&fold_helper_lines(CoreDialect::Metal, p));
s.push_str("inline uint rotr_var(uint x, uint n) { n &= 31u; return (x >> n) | (x << ((32u - n) & 31u)); }\n");
s.push_str("inline uint ds_elem(uint i, uint d0, uint d1) {\n");
s.push_str(" uint x = i ^ d0;\n");
s.push_str(DS_ELEM_BODY);
s.push('\n');
if p.has_wide() && !geom.mulshift {
if p.has_wide() {
s.push_str("#define WMASK (MASK & ~31u)\n\n");
}
let mut buffer0 = "device const uint* dataset [[buffer(0)]]";
@ -1047,7 +929,7 @@ fn metal_program_impl(p: &Program, geom: DatasetGeom, source: LoadSource, bound:
s.push_str(" uint gid [[thread_position_in_grid]]) {\n");
}
s.push_str(" uint nonce = baseNonce + gid;\n");
s.push_str(&reg_decl(p, "uint"));
s.push_str(" uint r0, r1, r2, r3, r4, r5, r6, r7;\n");
if p.has_wide() {
s.push_str(" uint lane = gid & 31u;\n");
}
@ -1059,14 +941,13 @@ fn metal_program_impl(p: &Program, geom: DatasetGeom, source: LoadSource, bound:
(i + 1) & 7
));
}
s.push_str(&reg64_init(p));
s.push_str(&format!("\n for (uint it = 0u; it < {ITERATIONS}u; ++it) {{\n uint sel = r0;\n"));
let era = p.class.era;
let word_index = |a: &str, wide: bool, ins: &Instr| -> String {
if wide {
wload_base_expr(CoreDialect::Metal, &format!("simd_broadcast({a}, 0)"), geom)
format!("(simd_broadcast({a}, 0) & WMASK) + lane")
} else {
load_index_expr(CoreDialect::Metal, era.as_ref(), ins, a, geom)
load_index_expr(CoreDialect::Metal, era.as_ref(), ins, a, dataset_log2)
}
};
let fetch = |idx: String| -> String {
@ -1076,17 +957,10 @@ fn metal_program_impl(p: &Program, geom: DatasetGeom, source: LoadSource, bound:
LoadSource::InlineMemhard(_) => format!("mh_word(cache, {idx})"),
}
};
// the drawn program, or its two-window interleaving under reg64 (Program::scheduled, the CUDA text's list);
// reg64 full chain: every load's address source is (rS ^ m), m mixed just before the load (the same chain)
let address_mix = p.address_mix();
let addr_src = |a: &str| -> String { if address_mix { format!("({a} ^ m)") } else { a.to_string() } };
for (k, ins) in p.scheduled().iter().enumerate() {
for (k, ins) in p.instrs.iter().enumerate() {
let d = format!("r{}", ins.dst);
let a = format!("r{}", ins.src);
let b = format!("r{}", ins.src2);
if address_mix && ins.op == Op::Load {
s.push_str(&format!(" {}\n", reg64_mix_line(p, ins.src)));
}
let line = match ins.op {
Op::Add => format!(
"{d} = {d} + {a} + select({}, {}, ((sel >> {}u) & 1u) != 0u);",
@ -1109,9 +983,9 @@ fn metal_program_impl(p: &Program, geom: DatasetGeom, source: LoadSource, bound:
LoadSource::InlineClosed(d0, d1) => (WideSource::InlineClosed, Some((*d0, *d1))),
LoadSource::InlineMemhard(_) => (WideSource::InlineMemhard, None),
};
wide_load_stmt(CoreDialect::Metal, &d, &word_index(&addr_src(&a), false, ins), ins.width, src, closed)
wide_load_stmt(CoreDialect::Metal, &d, &word_index(&a, false, ins), ins.width, src, closed)
}
Op::Load => format!("{d} = {d} ^ {};", fetch(word_index(&addr_src(&a), false, ins))),
Op::Load => format!("{d} = {d} ^ {};", fetch(word_index(&a, false, ins))),
Op::WLoad => format!("{d} = {d} ^ {};", fetch(word_index(&a, true, ins))),
Op::Scratch => scratch_stmt(CoreDialect::Metal, &d, &a, p.class.scratch_slot_mask()),
Op::Hot => hot_stmt(CoreDialect::Metal, &d, &a),
@ -1120,7 +994,6 @@ fn metal_program_impl(p: &Program, geom: DatasetGeom, source: LoadSource, bound:
}
s.push_str(&shadow_block(p, CoreDialect::Metal));
s.push_str(" }\n");
s.push_str(&reg64_fold(p));
s.push_str(" uint lo = r0 ^ rotl_imm(r1, 7u) ^ rotl_imm(r2, 14u) ^ rotl_imm(r3, 21u);\n");
s.push_str(" uint hi = r4 ^ rotl_imm(r5, 9u) ^ rotl_imm(r6, 18u) ^ rotl_imm(r7, 27u);\n");
s.push_str(" out[gid] = ((ulong)hi << 32) | (ulong)lo;\n");
@ -1154,78 +1027,14 @@ fn init_line(p: &Program, u: &str, i: usize) -> String {
)
}
/// The per-lane register declaration of the kernels (`ty` the dialect's 32-bit word: `uint32_t` in CUDA, `uint` in
/// OpenCL and Metal): r0..r7, or r0..r63 under reg64, with `m` for the full chain.
fn reg_decl(p: &Program, ty: &str) -> String {
if !p.class.reg64 {
return format!(" {ty} r0, r1, r2, r3, r4, r5, r6, r7;\n");
}
let names: Vec<String> = (0..p.registers()).map(|k| format!("r{k}")).collect();
let m = if p.address_mix() { format!("\n {ty} m; // reg64 full chain: the address mix of all 64 registers before every load") } else { String::new() };
format!(" {ty} {}; // reg64: 64 live registers per lane (two 32-register windows){m}\n", names.join(", "))
}
/// reg64 full chain: the mix statement before a load, `m = r0; m = rotl_imm(m, 1u) ^ r1; ... ^ r63;` over the 63
/// registers other than the load's source `src` (the verifier's `addr_src`, the same chain), one line; the same
/// text in the three dialects (`rotl_imm` is defined in each).
fn reg64_mix_line(p: &Program, src: u8) -> String {
let mut s = String::new();
for k in 0..p.registers() {
if k == src as usize {
continue;
}
if s.is_empty() {
s.push_str(&format!("m = r{k};"));
} else {
s.push_str(&format!(" m = rotl_imm(m, 1u) ^ r{k};"));
}
}
s
}
/// reg64: r8..r63 from the eight seeded registers, `r[k] = r[k & 7] * 0x9E3779B9 + k` (the verifier's text), the
/// same statements in the three dialects. Empty otherwise.
fn reg64_init(p: &Program) -> String {
if !p.class.reg64 {
return String::new();
}
let mut s = String::from(" // reg64: the derived registers of both windows\n");
for k in 8..p.registers() {
s.push_str(&format!(" r{k} = r{} * 0x9e3779b9u + {k}u;\n", k & 7));
}
s
}
/// reg64: both windows fold into r0..r7 by xor before the hash fold, the same statements in the three dialects.
/// Empty otherwise.
fn reg64_fold(p: &Program) -> String {
if !p.class.reg64 {
return String::new();
}
let mut s = String::from(" // reg64: fold the 64 registers into the eight output registers\n");
for k in 0..8 {
let terms: Vec<String> = (k + 8..p.registers()).step_by(8).map(|j| format!("r{j}")).collect();
s.push_str(&format!(" r{k} = r{k} ^ {};\n", terms.join(" ^ ")));
}
s
}
/// The instruction lines of the CUDA hash kernel body (shared by `igneum_hash` and `igneum_hash_bound`).
fn cuda_instr_lines(p: &Program, geom: DatasetGeom) -> String {
fn cuda_instr_lines(p: &Program, dataset_log2: u32) -> String {
let mut s = String::with_capacity(6000);
let era = p.class.era;
// the drawn program, or its two-window interleaving under reg64 (Program::scheduled: the verifier reads the same list)
// reg64 full chain: every load's address source is (rS ^ m), m the rotate-xor mix of all 64 registers computed
// just before the load (the verifier's addr_src, the same chain)
let address_mix = p.address_mix();
let addr_src = |a: &str| -> String { if address_mix { format!("({a} ^ m)") } else { a.to_string() } };
for (k, ins) in p.scheduled().iter().enumerate() {
for (k, ins) in p.instrs.iter().enumerate() {
let d = format!("r{}", ins.dst);
let a = format!("r{}", ins.src);
let b = format!("r{}", ins.src2);
if address_mix && ins.op == Op::Load {
s.push_str(&format!(" {}\n", reg64_mix_line(p, ins.src)));
}
let line = match ins.op {
// Metal select(A, B, c) returns c ? B : A, so the true branch is imm2 here as well.
Op::Add => format!(
@ -1244,10 +1053,10 @@ fn cuda_instr_lines(p: &Program, geom: DatasetGeom) -> String {
Op::Mad => format!("{d} = {a} * {b} + {d};"),
Op::Shfl => format!("{d} = {d} ^ __shfl_xor_sync(0xffffffffu, {a}, {});", ins.mask),
Op::Load if load_width(ins) > 1 => {
wide_load_stmt(CoreDialect::Cuda, &d, &load_index_expr(CoreDialect::Cuda, era.as_ref(), ins, &addr_src(&a), geom), ins.width, WideSource::Stored, None)
wide_load_stmt(CoreDialect::Cuda, &d, &load_index_expr(CoreDialect::Cuda, era.as_ref(), ins, &a, dataset_log2), ins.width, WideSource::Stored, None)
}
Op::Load => format!("{d} = {d} ^ ds[{}];", load_index_expr(CoreDialect::Cuda, era.as_ref(), ins, &addr_src(&a), geom)),
Op::WLoad => format!("{d} = {d} ^ ds[{}];", wload_base_expr(CoreDialect::Cuda, &format!("__shfl_sync(0xffffffffu, {a}, 0)"), geom)),
Op::Load => format!("{d} = {d} ^ ds[{}];", load_index_expr(CoreDialect::Cuda, era.as_ref(), ins, &a, dataset_log2)),
Op::WLoad => format!("{d} = {d} ^ ds[(__shfl_sync(0xffffffffu, {a}, 0) & wmask) + lane];"),
Op::Scratch => scratch_stmt(CoreDialect::Cuda, &d, &a, p.class.scratch_slot_mask()),
Op::Hot => hot_stmt(CoreDialect::Cuda, &d, &a),
};
@ -1264,11 +1073,6 @@ pub fn cuda_kernel(p: &Program, memhard: Option<&MixParams>) -> String {
/// [`cuda_kernel`] at a dataset size (an era program's window constants are literals of the pack's size; every
/// other class ignores it).
pub fn cuda_kernel_at(p: &Program, memhard: Option<&MixParams>, dataset_log2: u32) -> String {
cuda_kernel_geom(p, memhard, DatasetGeom::pow2(dataset_log2))
}
/// [`cuda_kernel_at`] at a dataset geometry (the multiply-shift sizes of `--dataset-words`).
pub fn cuda_kernel_geom(p: &Program, memhard: Option<&MixParams>, geom: DatasetGeom) -> String {
let layout = p.class.layout();
let mut s = String::with_capacity(9000);
s.push_str(&generated_by(&p.seed_string));
@ -1283,7 +1087,6 @@ pub fn cuda_kernel_geom(p: &Program, memhard: Option<&MixParams>, geom: DatasetG
s.push_str("#include \"memhard.h\"\n");
}
s.push('\n');
s.push_str(&ds_words_lines(CoreDialect::Cuda, geom));
s.push_str(&hot_define(p));
s.push_str("__device__ __forceinline__ uint32_t splitmix32(uint32_t x) {\n");
s.push_str(" x ^= x >> 16; x *= 0x7feb352du;\n");
@ -1293,7 +1096,6 @@ pub fn cuda_kernel_geom(p: &Program, memhard: Option<&MixParams>, geom: DatasetG
s.push_str("}\n");
s.push_str("// n is a literal in 1..31 at every call site, so both shift amounts are in 1..31.\n");
s.push_str("__device__ __forceinline__ uint32_t rotl_imm(uint32_t x, uint32_t n) { return (x << n) | (x >> (32u - n)); }\n");
s.push_str(&fold_helper_lines(CoreDialect::Cuda, p));
s.push_str("// n is masked to 0..31; the second shift amount is masked too, so n == 0 gives x.\n");
s.push_str("__device__ __forceinline__ uint32_t rotr_var(uint32_t x, uint32_t n) { n &= 31u; return (x >> n) | (x << ((32u - n) & 31u)); }\n");
s.push_str("__device__ __forceinline__ uint32_t ds_elem(uint32_t i, uint32_t d0, uint32_t d1) {\n");
@ -1360,22 +1162,17 @@ pub fn cuda_kernel_geom(p: &Program, memhard: Option<&MixParams>, geom: DatasetG
s.push_str(" uint32_t gid = blockIdx.x * blockDim.x + threadIdx.x;\n");
}
s.push_str(" uint32_t nonce = baseNonce + gid;\n");
s.push_str(&reg_decl(p, "uint32_t"));
s.push_str(" uint32_t r0, r1, r2, r3, r4, r5, r6, r7;\n");
if p.has_wide() {
s.push_str(" uint32_t lane = threadIdx.x & 31u;\n");
if !geom.mulshift {
s.push_str(" uint32_t wmask = mask & ~31u;\n");
}
s.push_str(" uint32_t lane = threadIdx.x & 31u;\n uint32_t wmask = mask & ~31u;\n");
}
for i in 0..8 {
s.push_str(&init_line(p, "uint32_t", i));
}
s.push_str(&reg64_init(p));
s.push_str(&format!("\n for (uint32_t it = 0u; it < {ITERATIONS}u; ++it) {{\n uint32_t sel = r0;\n"));
s.push_str(&cuda_instr_lines(p, geom));
s.push_str(&cuda_instr_lines(p, dataset_log2));
s.push_str(&shadow_block(p, CoreDialect::Cuda));
s.push_str(" }\n");
s.push_str(&reg64_fold(p));
s.push_str(" uint32_t lo = r0 ^ rotl_imm(r1, 7u) ^ rotl_imm(r2, 14u) ^ rotl_imm(r3, 21u);\n");
s.push_str(" uint32_t hi = r4 ^ rotl_imm(r5, 9u) ^ rotl_imm(r6, 18u) ^ rotl_imm(r7, 27u);\n");
s.push_str(" out[gid] = ((uint64_t)hi << 32) | (uint64_t)lo;\n");
@ -1475,11 +1272,6 @@ pub fn cuda_kernel_bound(p: &Program, memhard: Option<&MixParams>) -> String {
/// [`cuda_kernel_bound`] at a dataset size (see [`cuda_kernel_at`]).
pub fn cuda_kernel_bound_at(p: &Program, memhard: Option<&MixParams>, dataset_log2: u32) -> String {
cuda_kernel_bound_geom(p, memhard, DatasetGeom::pow2(dataset_log2))
}
/// [`cuda_kernel_bound_at`] at a dataset geometry.
pub fn cuda_kernel_bound_geom(p: &Program, memhard: Option<&MixParams>, geom: DatasetGeom) -> String {
let mut s = String::with_capacity(9000);
s.push_str(&generated_by(&p.seed_string));
s.push_str(
@ -1496,7 +1288,6 @@ pub fn cuda_kernel_bound_geom(p: &Program, memhard: Option<&MixParams>, geom: Da
s.push_str("#include <cstdint>\n");
s.push_str("#include \"program.h\"\n");
s.push('\n');
s.push_str(&ds_words_lines(CoreDialect::Cuda, geom));
s.push_str("struct IgneumInitWords { uint32_t w[8]; };\n");
s.push('\n');
s.push_str(&hot_define(p));
@ -1507,7 +1298,6 @@ pub fn cuda_kernel_bound_geom(p: &Program, memhard: Option<&MixParams>, geom: Da
s.push_str(" return x;\n");
s.push_str("}\n");
s.push_str("__device__ __forceinline__ uint32_t rotl_imm(uint32_t x, uint32_t n) { return (x << n) | (x >> (32u - n)); }\n");
s.push_str(&fold_helper_lines(CoreDialect::Cuda, p));
s.push_str("__device__ __forceinline__ uint32_t rotr_var(uint32_t x, uint32_t n) { n &= 31u; return (x >> n) | (x << ((32u - n) & 31u)); }\n");
s.push('\n');
let _ = memhard; // the bound kernel reads the stored dataset in both constructions
@ -1521,12 +1311,9 @@ pub fn cuda_kernel_bound_geom(p: &Program, memhard: Option<&MixParams>, geom: Da
s.push_str(" uint32_t gid = blockIdx.x * blockDim.x + threadIdx.x;\n");
}
s.push_str(" uint32_t nonce = baseNonce + gid;\n");
s.push_str(&reg_decl(p, "uint32_t"));
s.push_str(" uint32_t r0, r1, r2, r3, r4, r5, r6, r7;\n");
if p.has_wide() {
s.push_str(" uint32_t lane = threadIdx.x & 31u;\n");
if !geom.mulshift {
s.push_str(" uint32_t wmask = mask & ~31u;\n");
}
s.push_str(" uint32_t lane = threadIdx.x & 31u;\n uint32_t wmask = mask & ~31u;\n");
}
for i in 0..8 {
s.push_str(&format!(
@ -1535,12 +1322,10 @@ pub fn cuda_kernel_bound_geom(p: &Program, memhard: Option<&MixParams>, geom: Da
(i + 1) & 7
));
}
s.push_str(&reg64_init(p));
s.push_str(&format!("\n for (uint32_t it = 0u; it < {ITERATIONS}u; ++it) {{\n uint32_t sel = r0;\n"));
s.push_str(&cuda_instr_lines(p, geom));
s.push_str(&cuda_instr_lines(p, dataset_log2));
s.push_str(&shadow_block(p, CoreDialect::Cuda));
s.push_str(" }\n");
s.push_str(&reg64_fold(p));
s.push_str(" uint32_t lo = r0 ^ rotl_imm(r1, 7u) ^ rotl_imm(r2, 14u) ^ rotl_imm(r3, 21u);\n");
s.push_str(" uint32_t hi = r4 ^ rotl_imm(r5, 9u) ^ rotl_imm(r6, 18u) ^ rotl_imm(r7, 27u);\n");
s.push_str(" out[gid] = ((uint64_t)hi << 32) | (uint64_t)lo;\n");
@ -1584,20 +1369,13 @@ pub fn cuda_kernel_bound_geom(p: &Program, memhard: Option<&MixParams>, geom: Da
}
/// The instruction lines of the OpenCL hash kernel body (shared by `igneum_hash` and `igneum_hash_bound`).
fn opencl_instr_lines(p: &Program, geom: DatasetGeom) -> String {
fn opencl_instr_lines(p: &Program, dataset_log2: u32) -> String {
let mut s = String::with_capacity(6000);
let era = p.class.era;
// the drawn program, or its two-window interleaving under reg64 (Program::scheduled, the CUDA text's list);
// reg64 full chain: every load's address source is (rS ^ m), m mixed just before the load (the same chain)
let address_mix = p.address_mix();
let addr_src = |a: &str| -> String { if address_mix { format!("({a} ^ m)") } else { a.to_string() } };
for (k, ins) in p.scheduled().iter().enumerate() {
for (k, ins) in p.instrs.iter().enumerate() {
let d = format!("r{}", ins.dst);
let a = format!("r{}", ins.src);
let b = format!("r{}", ins.src2);
if address_mix && ins.op == Op::Load {
s.push_str(&format!(" {}\n", reg64_mix_line(p, ins.src)));
}
let line = match ins.op {
Op::Add => format!(
"{d} = {d} + {a} + ((((sel >> {}u) & 1u) != 0u) ? {} : {});",
@ -1615,10 +1393,10 @@ fn opencl_instr_lines(p: &Program, geom: DatasetGeom) -> String {
Op::Mad => format!("{d} = {a} * {b} + {d};"),
Op::Shfl => format!("{{ uint t_; IGNEUM_SHFL_XOR(t_, {a}, {}u); {d} = {d} ^ t_; }}", ins.mask),
Op::Load if load_width(ins) > 1 => {
wide_load_stmt(CoreDialect::OpenCl, &d, &load_index_expr(CoreDialect::OpenCl, era.as_ref(), ins, &addr_src(&a), geom), ins.width, WideSource::Stored, None)
wide_load_stmt(CoreDialect::OpenCl, &d, &load_index_expr(CoreDialect::OpenCl, era.as_ref(), ins, &a, dataset_log2), ins.width, WideSource::Stored, None)
}
Op::Load => format!("{d} = {d} ^ ds[{}];", load_index_expr(CoreDialect::OpenCl, era.as_ref(), ins, &addr_src(&a), geom)),
Op::WLoad => format!("{{ uint t_; IGNEUM_BCAST0(t_, {a}); {d} = {d} ^ ds[{}]; }}", wload_base_expr(CoreDialect::OpenCl, "t_", geom)),
Op::Load => format!("{d} = {d} ^ ds[{}];", load_index_expr(CoreDialect::OpenCl, era.as_ref(), ins, &a, dataset_log2)),
Op::WLoad => format!("{{ uint t_; IGNEUM_BCAST0(t_, {a}); {d} = {d} ^ ds[(t_ & wmask) + lane]; }}"),
Op::Scratch => scratch_stmt(CoreDialect::OpenCl, &d, &a, p.class.scratch_slot_mask()),
Op::Hot => hot_stmt(CoreDialect::OpenCl, &d, &a),
};
@ -1636,12 +1414,7 @@ pub fn opencl_kernel_bound(p: &Program, memhard: Option<&MixParams>) -> String {
/// [`opencl_kernel_bound`] at a dataset size (see [`cuda_kernel_at`]).
pub fn opencl_kernel_bound_at(p: &Program, memhard: Option<&MixParams>, dataset_log2: u32) -> String {
opencl_kernel_bound_geom(p, memhard, DatasetGeom::pow2(dataset_log2))
}
/// [`opencl_kernel_bound_at`] at a dataset geometry.
pub fn opencl_kernel_bound_geom(p: &Program, memhard: Option<&MixParams>, geom: DatasetGeom) -> String {
let mut s = opencl_kernel_geom(p, memhard, geom);
let mut s = opencl_kernel_at(p, memhard, dataset_log2);
s.push('\n');
s.push_str(
"// Header-bound variant (bind.rs): the init words come from initw, not SEEDW. Same body as igneum_hash.\n",
@ -1666,11 +1439,11 @@ pub fn opencl_kernel_bound_geom(p: &Program, memhard: Option<&MixParams>, geom:
s.push_str("#endif\n");
s.push_str(&unit_loop);
s.push_str(" uint nonce = baseNonce + gid;\n");
s.push_str(&reg_decl(p, "uint"));
s.push_str(" uint r0, r1, r2, r3, r4, r5, r6, r7;\n");
s.push_str(" uint iw0 = initw[0], iw1 = initw[1], iw2 = initw[2], iw3 = initw[3], iw4 = initw[4], iw5 = initw[5], iw6 = initw[6], iw7 = initw[7];\n");
} else {
s.push_str(" uint nonce = baseNonce + gid;\n");
s.push_str(&reg_decl(p, "uint"));
s.push_str(" uint r0, r1, r2, r3, r4, r5, r6, r7;\n");
s.push_str(" uint iw0 = initw[0], iw1 = initw[1], iw2 = initw[2], iw3 = initw[3], iw4 = initw[4], iw5 = initw[5], iw6 = initw[6], iw7 = initw[7];\n");
s.push_str("#if IGNEUM_EXCHANGE == 0\n");
s.push_str(" IGNEUM_LOCAL_WORDS(xch, 2 * IGNEUM_GROUP);\n");
@ -1680,10 +1453,7 @@ pub fn opencl_kernel_bound_geom(p: &Program, memhard: Option<&MixParams>, geom:
s.push_str("#endif\n");
}
if p.has_wide() {
s.push_str(" uint lane = lid & 31u;\n");
if !geom.mulshift {
s.push_str(" uint wmask = mask & ~31u;\n");
}
s.push_str(" uint lane = lid & 31u;\n uint wmask = mask & ~31u;\n");
}
for i in 0..8 {
s.push_str(&format!(
@ -1692,12 +1462,10 @@ pub fn opencl_kernel_bound_geom(p: &Program, memhard: Option<&MixParams>, geom:
(i + 1) & 7
));
}
s.push_str(&reg64_init(p));
s.push_str(&format!("\n for (uint it = 0u; it < {ITERATIONS}u; ++it) {{\n uint sel = r0;\n"));
s.push_str(&opencl_instr_lines(p, geom));
s.push_str(&opencl_instr_lines(p, dataset_log2));
s.push_str(&shadow_block(p, CoreDialect::OpenCl));
s.push_str(" }\n");
s.push_str(&reg64_fold(p));
s.push_str(" uint lo = r0 ^ rotl_imm(r1, 7u) ^ rotl_imm(r2, 14u) ^ rotl_imm(r3, 21u);\n");
s.push_str(" uint hi = r4 ^ rotl_imm(r5, 9u) ^ rotl_imm(r6, 18u) ^ rotl_imm(r7, 27u);\n");
s.push_str(" out[gid] = ((ulong)hi << 32) | (ulong)lo;\n");
@ -1715,11 +1483,6 @@ pub fn opencl_kernel(p: &Program, memhard: Option<&MixParams>) -> String {
/// [`opencl_kernel`] at a dataset size (see [`cuda_kernel_at`]).
pub fn opencl_kernel_at(p: &Program, memhard: Option<&MixParams>, dataset_log2: u32) -> String {
opencl_kernel_geom(p, memhard, DatasetGeom::pow2(dataset_log2))
}
/// [`opencl_kernel_at`] at a dataset geometry (the multiply-shift sizes of `--dataset-words`).
pub fn opencl_kernel_geom(p: &Program, memhard: Option<&MixParams>, geom: DatasetGeom) -> String {
let layout = p.class.layout();
let mut s = String::with_capacity(14000);
s.push_str(&generated_by(&p.seed_string));
@ -1735,7 +1498,6 @@ pub fn opencl_kernel_geom(p: &Program, memhard: Option<&MixParams>, geom: Datase
s.push_str("// the exchange masks are 1, 2, 4, 8, 16, so every partner lane lies inside the lane's own aligned run of 32.\n");
s.push_str("#ifndef IGNEUM_GROUP\n#define IGNEUM_GROUP 32\n#endif\n");
s.push_str("#ifndef IGNEUM_EXCHANGE\n#define IGNEUM_EXCHANGE 0\n#endif\n");
s.push_str(&ds_words_lines(CoreDialect::OpenCl, geom));
s.push_str("#ifdef __OPENCL_VERSION__\n");
s.push_str("#define IGNEUM_KERNEL_HASH __kernel __attribute__((reqd_work_group_size(IGNEUM_GROUP, 1, 1)))\n");
s.push_str("#define IGNEUM_LOCAL_WORDS(name, n) __local uint name[n]\n");
@ -1776,7 +1538,6 @@ pub fn opencl_kernel_geom(p: &Program, memhard: Option<&MixParams>, geom: Datase
s.push_str("}\n");
s.push_str("// n is a literal in 1..31 at every call site. OpenCL rotate() rotates left by n modulo 32.\n");
s.push_str("static inline uint rotl_imm(uint x, uint n) { return rotate(x, n); }\n");
s.push_str(&fold_helper_lines(CoreDialect::OpenCl, p));
s.push_str("// Right rotation by n modulo 32 as a left rotation by (32 - n) modulo 32; n == 0 gives x.\n");
s.push_str("static inline uint rotr_var(uint x, uint n) { return rotate(x, (0u - n) & 31u); }\n");
s.push_str("static inline uint ds_elem(uint i, uint d0, uint d1) {\n");
@ -1851,10 +1612,10 @@ pub fn opencl_kernel_geom(p: &Program, memhard: Option<&MixParams>, geom: Datase
s.push_str("#endif\n");
s.push_str(&unit_loop);
s.push_str(" uint nonce = baseNonce + gid;\n");
s.push_str(&reg_decl(p, "uint"));
s.push_str(" uint r0, r1, r2, r3, r4, r5, r6, r7;\n");
} else {
s.push_str(" uint nonce = baseNonce + gid;\n");
s.push_str(&reg_decl(p, "uint"));
s.push_str(" uint r0, r1, r2, r3, r4, r5, r6, r7;\n");
s.push_str("#if IGNEUM_EXCHANGE == 0\n");
s.push_str(" IGNEUM_LOCAL_WORDS(xch, 2 * IGNEUM_GROUP);\n");
s.push_str(" uint xk = 0u;\n");
@ -1863,20 +1624,15 @@ pub fn opencl_kernel_geom(p: &Program, memhard: Option<&MixParams>, geom: Datase
s.push_str("#endif\n");
}
if p.has_wide() {
s.push_str(" uint lane = lid & 31u;\n");
if !geom.mulshift {
s.push_str(" uint wmask = mask & ~31u;\n");
}
s.push_str(" uint lane = lid & 31u;\n uint wmask = mask & ~31u;\n");
}
for i in 0..8 {
s.push_str(&init_line(p, "uint", i));
}
s.push_str(&reg64_init(p));
s.push_str(&format!("\n for (uint it = 0u; it < {ITERATIONS}u; ++it) {{\n uint sel = r0;\n"));
s.push_str(&opencl_instr_lines(p, geom));
s.push_str(&opencl_instr_lines(p, dataset_log2));
s.push_str(&shadow_block(p, CoreDialect::OpenCl));
s.push_str(" }\n");
s.push_str(&reg64_fold(p));
s.push_str(" uint lo = r0 ^ rotl_imm(r1, 7u) ^ rotl_imm(r2, 14u) ^ rotl_imm(r3, 21u);\n");
s.push_str(" uint hi = r4 ^ rotl_imm(r5, 9u) ^ rotl_imm(r6, 18u) ^ rotl_imm(r7, 27u);\n");
s.push_str(" out[gid] = ((ulong)hi << 32) | (ulong)lo;\n");
@ -1911,8 +1667,7 @@ pub fn program_header(p: &Program, day: &str, ds: &DatasetSource) -> String {
let key = &ds.key;
let dataset_log2 = ds.log2_words;
let memhard = ds.memhard().map(|m| &m.params);
let geom = ds.geom;
let mask = geom.mask();
let mask = mask_for(dataset_log2);
let mut s = String::with_capacity(2600);
s.push_str(&generated_by(&p.seed_string));
s.push_str("// Program metadata for host.cu plus the launch wrappers defined in kernel.cu.\n");
@ -1930,20 +1685,8 @@ pub fn program_header(p: &Program, day: &str, ds: &DatasetSource) -> String {
s.push_str(&format!("#define IGNEUM_DAY_BYTES_HEX {}\n", jstr(&hex_bytes(&ds.key_bytes))));
s.push_str(&format!("#define IGNEUM_DAY0 {}\n", hex(key[0])));
s.push_str(&format!("#define IGNEUM_DAY1 {}\n", hex(key[1])));
if geom.mulshift {
s.push_str(&format!("// Research class ds55 (8 October 2026): a dataset of {} words, not a power of two. IGNEUM_DATASET_LOG2 is floor(log2(words));\n", geom.words));
s.push_str("// the host allocates IGNEUM_DATASET_WORDS words and builds IGNEUM_DATASET_ITEMS items; IGNEUM_MASK is the last word index (the\n");
s.push_str("// self-test reads dataset[IGNEUM_MASK]) and is never ANDed: every load is idx = (src * IGNEUM_DATASET_WORDS) >> 32 (spec 01 section 1.13.3).\n");
s.push_str(&format!("#define IGNEUM_DATASET_LOG2 {dataset_log2}\n"));
s.push_str(&format!("#define IGNEUM_DATASET_WORDS {}u\n", geom.words));
s.push_str(&format!("#define IGNEUM_DATASET_ITEMS {}u\n", geom.items()));
s.push_str(&format!("#define IGNEUM_DATASET_BYTES {}ull\n", geom.bytes()));
s.push_str("#define IGNEUM_DATASET_MULSHIFT 1\n");
s.push_str(&format!("#define IGNEUM_MASK {}\n", hex(mask)));
} else {
s.push_str(&format!("#define IGNEUM_DATASET_LOG2 {dataset_log2}\n"));
s.push_str(&format!("#define IGNEUM_MASK {}\n", hex(mask)));
}
s.push_str(&format!("#define IGNEUM_DATASET_LOG2 {dataset_log2}\n"));
s.push_str(&format!("#define IGNEUM_MASK {}\n", hex(mask)));
s.push_str("#define IGNEUM_LANES 32\n");
s.push_str(&format!("#define IGNEUM_ITERATIONS {ITERATIONS}\n"));
s.push_str(&format!("#define IGNEUM_INSTR_COUNT {INSTR_COUNT}\n"));
@ -2077,13 +1820,6 @@ pub fn sample_indices(mask: u32) -> Vec<u32> {
(0..64).map(|_| (sr.next() as u32) & mask).collect()
}
/// [`sample_indices`] at a dataset geometry: the same 64 draws through the geometry's range reduction, so the
/// mask path is [`sample_indices`] exactly and the multiply-shift path samples by the mapping the loads use.
pub fn sample_indices_geom(geom: DatasetGeom) -> Vec<u32> {
let mut sr = SplitMix64::new(0x6d68_7361_6d70_6c65);
(0..64).map(|_| geom.reduce(sr.next() as u32)).collect()
}
/// vectors.h (`generateVectorsHeader`).
pub fn vectors_header(
p: &Program,
@ -2167,15 +1903,14 @@ pub fn program_json(p: &Program, day: &str, ds: &DatasetSource) -> String {
let key = &ds.key;
let dataset_log2 = ds.log2_words;
let memhard = ds.memhard().map(|m| &m.params);
let geom = ds.geom;
let mask = geom.mask();
let mask = mask_for(dataset_log2);
let mut s = String::with_capacity(14000);
s.push_str("{\n");
s.push_str(" \"format\": \"igneum-program-pack-3\",\n");
s.push_str(&format!(" \"generator\": {},\n", p.generator));
s.push_str(&format!(" \"attempt\": {},\n", p.attempt));
s.push_str(&format!(" \"program_id\": {},\n", jhex64(p.program_id())));
s.push_str(&format!(" \"program_id_derivation\": {},\n", jstr(&p.program_id_derivation())));
s.push_str(" \"program_id_derivation\": \"FNV-1a 64 over 'igneum-program/' || generator_le32 || seed_words as little-endian bytes || attempt_le32\",\n");
s.push_str(&format!(
" \"dataset_mode\": {},\n",
jstr(if memhard.is_some() { "memory-hard" } else { "closed-form" })
@ -2215,18 +1950,6 @@ pub fn program_json(p: &Program, day: &str, ds: &DatasetSource) -> String {
if !p.class.is_v2() {
let c = p.width_counts();
s.push_str(&format!(" \"load_class\": {},\n", jstr(&p.class.name())));
if p.class.rw != 0 {
// class v6 lane 1: the re-weight table the op roll draws from, in draw order, and its sum (the roll's range)
let (w, sum) = p.class.nonload_weights();
s.push_str(&format!(
" \"op_weights\": {{\"table\": \"rw{}\", \"sum\": {sum}, \"draw_order\": [{}], \"rule\": \"class v6 lane 1 (8 October 2026, a research class): the table replaces the plain ten-family table (sum 75) for the op roll of every base and shadow instruction, roll = below(sum) walked over the weights in draw order; a family at weight 0 is never drawn; shuffle stays at 4 under rw1; the program id carries 'rw/' || table_u8\"}},\n",
p.class.rw,
w.iter().map(|(o, n)| format!("[{}, {n}]", jstr(o.name()))).collect::<Vec<_>>().join(", ")
));
}
if p.class.reg64 {
s.push_str(&format!(" \"reg64\": {{\"registers\": {}, \"variant\": {}, \"address_mix\": {}, \"liveness\": {}, \"statements_per_iteration\": {}, \"loads_per_hash\": {}, \"rule\": \"the hash lane's 64-register window (8 October 2026, a research class, NOT the lottery hash): r0..r7 seeded as today, r[k] = r[k & 7] * 0x9E3779B9 + k for k in 8..63; instruction i of the drawn program runs on window 0 with every register field + 8 * (i % 4), then on window 1 with +32 more; after the 8 iterations r[k] ^= r[k + 8] ^ ... ^ r[k + 56] for k in 0..7, then the hash fold; the draw, the acceptance rule and the dataset are the class's without the flag; address_mix 1 (the full chain): every load's address source is src ^ m, m the rotate-xor chain (m = first; m = rotl(m, 1) ^ next) over the 63 registers other than src in index order, computed before the load; src stays out of the chain so no register's term can cancel its own direct term (liveness rule: xoring any register with either of two seed-derived probe words at the start of an iteration moves the iteration's first load address and the final hash)\"}},\n", p.registers(), jstr(if p.address_mix() { "window, full chain" } else { "window, arithmetic-only" }), p.address_mix() as u8, jstr(&p.reg64_liveness()), p.scheduled().len(), p.scheduled().iter().filter(|i| i.op.is_load()).count() * ITERATIONS));
}
if p.class.derive_len != 0 {
s.push_str(&format!(" \"derive_len\": {},\n", p.class.derive_len));
s.push_str(" \"derive\": \"Counter ASIC 3.0 item 2 (6 October 2026, docs/plans/counter-asic-3-derivation.md; a prototype, not class v3): the nine mixer slots of the item derivation each run a straight-line program of derive_len instructions drawn from the day key stream after the 40 mixer draws, four draws per instruction (op roll below(100), destination roll below(15), third-register roll below(14), the immediate next()); every instruction reads the register the previous one wrote (s[0] first) and writes another; twelve forms, each a bijection on the state; the 8 dependent cache reads per item unchanged; the acceptance test of derive.rs (every register written per round program, 8 distinct rotations, the x8 mixer's operation and multiply counts as floors) rejects a draw and the next attempt continues the stream\",\n");
@ -2256,15 +1979,7 @@ pub fn program_json(p: &Program, day: &str, ds: &DatasetSource) -> String {
s.push_str(&format!(" \"stride_mul\": {},\n", jhex(e.stride_mul)));
s.push_str(&format!(" \"stride_rot\": {},\n", e.stride_rot));
s.push_str(&format!(" \"interleave\": [{}, {}, {}, {}],\n", e.pos[0], e.pos[1], e.pos[2], e.pos[3]));
let y = if e.fold { format!("y = src * stride_mul; y ^= y >> {INDEX_FOLD_SHIFT}; y = rotl(y, stride_rot)") } else { "y = rotl(src * stride_mul, stride_rot)".to_string() };
if geom.mulshift {
s.push_str(&format!(" \"address\": \"{y}; D = floor(log2(words)); k = min(win, D - 26); v = (y & (0xffffffff >> k)) | ((off & (2^k - 1)) << (32 - k)); idx = (v * words) >> 32 in 64 bits (the window in the source space, then the multiply-shift of spec 01 section 1.13.3); a wide load aligns idx down to W words\",\n"));
} else {
s.push_str(&format!(" \"address\": \"{y}; k = min(win, D - 26); idx = ((y & (mask >> k)) | ((off & (2^k - 1)) << (D - k))) & mask; a wide load aligns idx down to W words\",\n"));
}
if e.fold {
s.push_str(&format!(" \"index_fold\": \"class v6 lane 1 (8 October 2026, docs/design/class-v6-rotating-family.md section 2, a research class): the product's low bits are folded before the stride rotation, y ^= y >> {INDEX_FOLD_SHIFT}, so no era's rotation lands a biased product bit (bit 0 of an odd product is bit 0 of src; bit 1 is set at 3/8; lane D's 7 of 16 eras at 130 to 511 sigma) on an address bit; the design names the fold and not its form, so this form is the lane's: one xor and one shift on the address path, the same in verify::stride, the three kernel texts (fold16_) and these vectors; the program id carries 'fold/'\",\n"));
}
s.push_str(" \"address\": \"y = rotl(src * stride_mul, stride_rot); k = min(win, D - 26); idx = ((y & (mask >> k)) | ((off & (2^k - 1)) << (D - k))) & mask; a wide load aligns idx down to W words\",\n");
s.push_str(" \"windows\": \"per instruction, after the width roll: win = below(3), off = low32(next()) & (2^win - 1); used on a load slot (the instruction's win and off fields)\",\n");
s.push_str(" \"dataset_word\": \"dataset[w] = item(t(w))[j(w)]: j(w) gathers the bits of w at the interleave positions, t(w) is w with those bits removed\",\n");
s.push_str(" \"program_id_suffix\": \"'era/' || allowed[3] || width_words || stride_mul_le32 || stride_rot_le32 || interleave[4]\"\n");
@ -2309,34 +2024,17 @@ pub fn program_json(p: &Program, day: &str, ds: &DatasetSource) -> String {
s.push_str(
" \"shfl\": \"dst = dst ^ (src of lane (lane ^ mask)), mask in {1,2,4,8,16}, within the 32-lane warp\",\n",
);
if geom.mulshift {
s.push_str(" \"load\": \"dst = dst ^ dataset[(src * dataset.words) >> 32]\",\n");
s.push_str(" \"wload\": \"base = ((src of lane 0 * dataset.words) >> 32) & ~31; dst = dst ^ dataset[base + lane] (warp-coalesced 128-byte load, lever b, only when --wide-frac > 0)\"");
} else {
s.push_str(" \"load\": \"dst = dst ^ dataset[src & dataset.mask]\",\n");
s.push_str(" \"wload\": \"base = (src of lane 0 & dataset.mask) & ~31; dst = dst ^ dataset[base + lane] (warp-coalesced 128-byte load, lever b, only when --wide-frac > 0)\"");
}
s.push_str(" \"load\": \"dst = dst ^ dataset[src & dataset.mask]\",\n");
s.push_str(" \"wload\": \"base = (src of lane 0 & dataset.mask) & ~31; dst = dst ^ dataset[base + lane] (warp-coalesced 128-byte load, lever b, only when --wide-frac > 0)\"");
if p.has_hot() {
s.push_str(",\n \"hot\": \"dst = dst ^ hot[mulhi(src, hot_table.words)] (hot-table experiment)\"");
}
s.push('\n');
s.push_str(" },\n");
s.push_str(" \"dataset\": {\n");
if geom.mulshift {
s.push_str(&format!(" \"log2_words\": {dataset_log2},\n"));
s.push_str(" \"log2_words_note\": \"floor(log2(words)): the dataset is not a power of two (research class ds55, 8 October 2026); allocate words, build items\",\n");
s.push_str(&format!(" \"words\": {},\n", geom.words));
s.push_str(&format!(" \"bytes\": {},\n", geom.bytes()));
s.push_str(&format!(" \"items\": {},\n", geom.items()));
s.push_str(" \"mapping\": \"mulshift\",\n");
s.push_str(" \"index\": \"idx = (src * words) >> 32 computed in 64 bits (spec 01 section 1.13.3, the multiply-shift range reduction; uniform to within 2^-32, branch-free, integer only), in place of src & mask; the item index is idx >> 4 under the linear layout and t(idx) under an era layout, below items\",\n");
s.push_str(&format!(" \"mask\": {},\n", jhex(mask)));
s.push_str(" \"mask_note\": \"the last word index, words - 1; never ANDed under the multiply-shift\",\n");
} else {
s.push_str(&format!(" \"log2_words\": {dataset_log2},\n"));
s.push_str(&format!(" \"bytes\": {},\n", 1u64 << (dataset_log2 as u64 + 2)));
s.push_str(&format!(" \"mask\": {},\n", jhex(mask)));
}
s.push_str(&format!(" \"log2_words\": {dataset_log2},\n"));
s.push_str(&format!(" \"bytes\": {},\n", 1u64 << (dataset_log2 as u64 + 2)));
s.push_str(&format!(" \"mask\": {},\n", jhex(mask)));
s.push_str(&format!(" \"day\": {},\n", jstr(day)));
s.push_str(&format!(" \"day_bytes\": {},\n", jstr(&hex_bytes(&ds.key_bytes))));
s.push_str(" \"day_words_from\": \"seed_words_from_bytes(day_bytes)\",\n");
@ -2480,36 +2178,12 @@ pub fn vectors_json(
source: &str,
memhard: bool,
) -> String {
let geom = DatasetGeom::pow2(dataset_log2);
debug_assert_eq!(mask, geom.mask());
vectors_json_geom(p, day, geom, bases, outs, v, source, memhard)
}
/// [`vectors_json`] at a dataset geometry: under the multiply-shift the file also carries `dataset_words` and
/// `dataset_mapping`, and `dataset_last_index` is `words - 1`.
#[allow(clippy::too_many_arguments)]
pub fn vectors_json_geom(
p: &Program,
day: &str,
geom: DatasetGeom,
bases: &[u32],
outs: &[[u64; 32]],
v: &PackVectors,
source: &str,
memhard: bool,
) -> String {
let dataset_log2 = geom.log2;
let mask = geom.mask();
let mut s = String::with_capacity(6500);
s.push_str("{\n");
s.push_str(&format!(" \"seed\": {},\n", jstr(&p.seed_string)));
s.push_str(&format!(" \"day\": {},\n", jstr(day)));
s.push_str(&format!(" \"dataset_mode\": {},\n", jstr(if memhard { "memory-hard" } else { "closed-form" })));
s.push_str(&format!(" \"dataset_log2_words\": {dataset_log2},\n"));
if geom.mulshift {
s.push_str(&format!(" \"dataset_words\": {},\n", geom.words));
s.push_str(" \"dataset_mapping\": \"mulshift\",\n");
}
s.push_str(&format!(" \"mask\": {},\n", jhex(mask)));
s.push_str(" \"lanes\": 32,\n");
s.push_str(&format!(" \"source\": {},\n", jstr(source)));
@ -2581,17 +2255,15 @@ impl Pack {
pub fn export_pack(epoch: &Epoch, day: &str, source: &str) -> Pack {
let p = &epoch.program;
let ds: &DatasetSource = &epoch.dataset;
let geom = ds.geom;
let mask = geom.mask();
let mask = ds.mask;
let memhard = ds.memhard().map(|m| &m.params);
let bases = PACK_VECTOR_BASES.to_vec();
let outs: Vec<[u64; 32]> = bases.iter().map(|&b| epoch.hash_warp(b)).collect();
// the self-test words under the program's layout (era layout; linear for every other class) and the
// geometry's range reduction (the sampled indices go through the mapping the loads use)
// the self-test words under the program's layout (era layout; linear for every other class)
let mut v = PackVectors {
head: (0..16).map(|i| epoch.dataset_word(i)).collect(),
last: epoch.dataset_word(geom.last_index()),
sample_idx: sample_indices_geom(geom),
last: epoch.dataset_word(mask),
sample_idx: sample_indices(mask),
..Default::default()
};
v.sample_val = v.sample_idx.iter().map(|&i| epoch.dataset_word(i)).collect();
@ -2612,16 +2284,16 @@ pub fn export_pack(epoch: &Epoch, day: &str, source: &str) -> Pack {
let is_mh = memhard.is_some();
let mut files = vec![
("program.json".to_string(), program_json(p, day, ds)),
("vectors.json".to_string(), vectors_json_geom(p, day, geom, &bases, &outs, &v, source, is_mh)),
("kernel.cu".to_string(), cuda_kernel_geom(p, memhard, geom)),
("kernel.cl".to_string(), opencl_kernel_geom(p, memhard, geom)),
("vectors.json".to_string(), vectors_json(p, day, ds.log2_words, &bases, &outs, &v, mask, source, is_mh)),
("kernel.cu".to_string(), cuda_kernel_at(p, memhard, ds.log2_words)),
("kernel.cl".to_string(), opencl_kernel_at(p, memhard, ds.log2_words)),
("program.h".to_string(), program_header(p, day, ds)),
("vectors.h".to_string(), vectors_header(p, &bases, &outs, &v, mask, source, is_mh)),
("program.metal".to_string(), metal_program_geom(p, geom, LoadSource::Stored)),
("program.metal".to_string(), metal_program(p, ds.log2_words, LoadSource::Stored)),
// Header-bound kernels (3 October 2026, bind.rs): new files, the seven above are unchanged.
("program_bound.metal".to_string(), metal_program_bound_geom(p, geom)),
("kernel_bound.cu".to_string(), cuda_kernel_bound_geom(p, memhard, geom)),
("kernel_bound.cl".to_string(), opencl_kernel_bound_geom(p, memhard, geom)),
("program_bound.metal".to_string(), metal_program_bound(p, ds.log2_words)),
("kernel_bound.cu".to_string(), cuda_kernel_bound_at(p, memhard, ds.log2_words)),
("kernel_bound.cl".to_string(), opencl_kernel_bound_at(p, memhard, ds.log2_words)),
];
if let Some(mp) = memhard {
files.push(("memhard.h".to_string(), cuda_memhard_header(p, mp)));

483
igneum-pow/src/generator.rs
View file

@ -242,39 +242,6 @@ pub struct LoadClass {
/// the item derivation XORs the window's state leaf into every item before the first mixer (`crate::state`,
/// `memhard::derive_items_leaves`). The program draw does not read it. `false` for every other class.
pub state: bool,
/// W = 8 (the floor programme, 8 October 2026, a research class behind `w32`): the 4-word slot of the width set
/// reads 8 words (one 32-byte sector) instead of 4; the mix, the slots and every other rule stand. `false` for
/// every other class. See [`LoadClass::widths`].
pub wide8: bool,
/// Class v6 lane 1 (`docs/design/class-v6-rotating-family.md` section 2, the index fold; a research class,
/// 8 October 2026): `true` folds the product's low bits in every era load address before the stride rotation
/// (`verify::load_index`: `y = x * M; y ^= y >> 16; y = rotl(y, R)`), so no era's R lands a biased product bit
/// on an address bit. The class's era carries the same bit ([`EraParams::fold`]). `false` for every other class.
pub fold: bool,
/// Class v6 lane 1, the op-mix re-weight behind the fold: 0 is the plain table [`NONLOAD_WEIGHTS`]; 1 the k lane's
/// optimiser split [`NONLOAD_WEIGHTS_RW`] (sum 83, `or` never drawn); 2 the census lane's neighbouring table
/// [`NONLOAD_WEIGHTS_RW2`] (sum 75). The draw rolls against the table's own sum ([`LoadClass::nonload_weights`]).
pub rw: u8,
/// The 64-register window (the hash lane's reg64 measurement, 8 October 2026, a research class behind `+reg64`
/// and `--reg64`): each lane holds 64 live 32-bit registers. r0..r7 are seeded as today, r8..r63 derived from them
/// (`r[k] = r[k & 7] * 0x9E3779B9 + k`); the 64 drawn instructions run twice per iteration in an interleaved
/// order, instruction i on window 0 (register field + 8 * (i % 4), registers 0..31) then the same instruction on
/// window 1 (+32, registers 32..63); the windows fold into r0..r7 by xor before the hash fold
/// ([`Program::scheduled`], [`crate::verify`], the emitters). The draw, the acceptance rule and the dataset are the
/// class's without the flag. `false` for every other class.
pub reg64: bool,
/// reg64, the full chain (the coordinator's amendment of 8 October 2026, 15:1x UK, class suffix `+reg64c`,
/// `--reg64-chain`): the address of every load consumes all 64 registers: address source `src ^ m`, `m` the
/// rotate-xor chain (`m = first; m = rotl(m, 1) ^ next`) over the 63 registers other than `src` in index order
/// (the same text in the verifier and the emitters), so an in-flight hash holds 64 independently necessary
/// values for the length of the dependent memory chain. The source stays out of the chain: inside it, r31 and
/// r63 land at rotation 0 mod 32 and cancel their own direct term (found by the liveness rule on the pinned draw).
/// Init rule: r0..r7 from the seed words as every class, `r[k] = r[k & 7] * 0x9E3779B9 + k` for k in 8..63.
/// Output rule: `r[k] ^= r[k + 8] ^ r[k + 16] ^ ... ^ r[k + 56]` for k in 0..7, then the class's hash fold.
/// Liveness rule (`accept::check_window_liveness`): xoring any one register with either of two seed-derived
/// probe words at the start of an iteration moves that iteration's first load address and the final hash, in
/// every lane (the complement and a single bit are the patterns a linear fold loses). Requires `reg64`.
pub reg64_chain: bool,
}
/// The parameters one era draws from its seed `E_n` (`docs/plans/era-layout.md` section 1.1, the proposed text of
@ -296,9 +263,6 @@ pub struct EraParams {
/// Layer 4, interleave: the four ascending bit positions (0..15) of the word-within-item bits in the word
/// index; the first `log2(width_words)` are `0..`, so one aligned load stays inside one item.
pub pos: [u8; 4],
/// The index fold of class v6 lane 1 ([`LoadClass::fold`]): the product's low bits folded before the rotation.
/// Not drawn: set from the class the era is composed over. `false` for every era of every other class.
pub fold: bool,
}
/// Domain tag of the era stream seed.
@ -358,7 +322,7 @@ impl EraParams {
pub fn era_draw(era_bytes: &[u8], allowed: &[u8]) -> EraParams {
assert!(!allowed.is_empty() && allowed.len() <= 3, "the allowed width set has 1 to 3 entries");
for (i, &w) in allowed.iter().enumerate() {
assert!(WIDTH_WORDS.contains(&w) || w == 8, "allowed width {w} is not 1, 4, 8 or 16 words");
assert!(WIDTH_WORDS.contains(&w), "allowed width {w} is not 1, 4 or 16 words");
assert!(i == 0 || allowed[i - 1] < w, "the allowed width set is ascending");
}
let words = EraParams::stream_words(era_bytes);
@ -394,7 +358,7 @@ pub fn era_draw(era_bytes: &[u8], allowed: &[u8]) -> EraParams {
}
let mut al = [0u8; 3];
al[..allowed.len()].copy_from_slice(allowed);
EraParams { words, allowed: al, width_words, stride_mul, stride_rot, pos, fold: false }
EraParams { words, allowed: al, width_words, stride_mul, stride_rot, pos }
}
/// The hot table of a class: `mb` MiB (32, 64 or 96 in the experiment) and `k` hot slots. Two forms: `replaced`
@ -455,13 +419,13 @@ impl LoadClass {
impl LoadClass {
/// Generator version 2 as adopted on 4 October 2026: 16 loads of one word. The lottery hash.
pub const V2: LoadClass =
LoadClass { mix: [100, 0, 0], load_slots: LOAD_SLOTS as u8, scratch: None, scratch_kb: 0, mixer_mult: 1, growth: false, era: None, hot: None, derive_len: 0, shadow: None, state: false, wide8: false, fold: false, rw: 0, reg64: false, reg64_chain: false };
LoadClass { mix: [100, 0, 0], load_slots: LOAD_SLOTS as u8, scratch: None, scratch_kb: 0, mixer_mult: 1, growth: false, era: None, hot: None, derive_len: 0, shadow: None, state: false };
/// The construction decided for program class v3 on 5 October 2026 (Counter ASIC 2.0, `docs/plans/mixer-x4.md`):
/// version 2 loads (16 slots of one word, no scratch, no width roll, so the program stream is version 2's), the
/// mixer applied 4 times per round, and the cache growth rule. Name "mx4".
pub const MX4: LoadClass =
LoadClass { mix: [100, 0, 0], load_slots: LOAD_SLOTS as u8, scratch: None, scratch_kb: 0, mixer_mult: 4, growth: true, era: None, hot: None, derive_len: 0, shadow: None, state: false, wide8: false, fold: false, rw: 0, reg64: false, reg64_chain: false };
LoadClass { mix: [100, 0, 0], load_slots: LOAD_SLOTS as u8, scratch: None, scratch_kb: 0, mixer_mult: 4, growth: true, era: None, hot: None, derive_len: 0, shadow: None, state: false };
/// The era class over `base` (`docs/plans/era-layout.md`): the parameters drawn by [`era_draw`]; when `allowed`
/// has more than one width the drawn width becomes the class mix (every load that width), otherwise the base
@ -475,14 +439,12 @@ impl LoadClass {
c.mix = [0, 0, 0];
c.mix[i] = 100;
} else {
let widest = (0..3).rev().find(|&i| c.mix[i] > 0).map(|i| c.widths()[i]).unwrap_or(1);
let widest = (0..3).rev().find(|&i| c.mix[i] > 0).map(|i| WIDTH_WORDS[i]).unwrap_or(1);
if widest != e.width_words {
// the interleave must keep the widest load inside one item: redraw the positions for that width
e = era_draw(era_bytes, &[widest]);
}
}
// class v6 lane 1: the era's address path carries the class's index fold
e.fold = c.fold;
c.era = Some(e);
c
}
@ -524,16 +486,6 @@ impl LoadClass {
LoadClass { mix, load_slots, ..LoadClass::V2 }
}
/// The width set this class draws from, in words: [`WIDTH_WORDS`], or `[1, 8, 16]` under `wide8` (W = 8).
pub fn widths(&self) -> [u8; 3] {
if self.wide8 { [1, 8, 16] } else { WIDTH_WORDS }
}
/// W = 8: every load reads 8 words (a 32-byte sector), `load_slots` loads per program (the class `w32`).
pub fn fixed8(load_slots: u8) -> LoadClass {
LoadClass { mix: [0, 100, 0], load_slots, wide8: true, ..LoadClass::V2 }
}
/// Per-load width drawn from `mix` (percent for 4, 16, 64 bytes), 16 loads per program.
pub fn mixed(mix: [u8; 3]) -> LoadClass {
assert_eq!(mix.iter().map(|&m| m as u32).sum::<u32>(), 100, "the mix must sum to 100");
@ -601,19 +553,7 @@ impl LoadClass {
/// This class with another class's era draw (tests: a rung's class composed with the chain's era).
pub fn with_era_of(self, other: &LoadClass) -> LoadClass {
// the era's fold bit follows THIS class's flag, not the other's
LoadClass { era: other.era.map(|mut e| { e.fold = self.fold; e }), ..self }
}
/// The class with the 64-register window per lane ("mx8+reg64", a research class).
pub fn with_reg64(self) -> LoadClass {
LoadClass { reg64: true, ..self }
}
/// The reg64 class with the full-chain address mix ("mx8+reg64c").
pub fn with_reg64_chain(self) -> LoadClass {
assert!(self.reg64, "the full-chain address mix is a reg64 variant");
LoadClass { reg64_chain: true, ..self }
LoadClass { era: other.era, ..self }
}
/// The class with the state leaves of class v5 folded into every item ("mx8+sh256x27+state").
@ -621,29 +561,6 @@ impl LoadClass {
LoadClass { state: true, ..self }
}
/// Class v6 lane 1: the class with the index fold on every era load address ("...+fold"). An era already
/// composed takes the bit too, so the flag and the era never disagree.
pub fn with_fold(self) -> LoadClass {
LoadClass { fold: true, era: self.era.map(|mut e| { e.fold = true; e }), ..self }
}
/// Class v6 lane 1: the class drawing its non-load ops from re-weight table `rw` (1: the k lane's split, "+rw";
/// 2: the census lane's table, "+rw2"; 0: the plain table).
pub fn with_rw(self, rw: u8) -> LoadClass {
assert!(rw <= 2, "re-weight table must be 0, 1 or 2");
LoadClass { rw, ..self }
}
/// The non-load op table this class draws from, in draw order, and its sum (the roll's range). The plain table
/// for every class without the re-weight flag, so their streams are byte for byte what they were.
pub fn nonload_weights(&self) -> (&'static [(Op, u64); 10], u64) {
match self.rw {
1 => (&NONLOAD_WEIGHTS_RW, NONLOAD_WEIGHTS_RW_SUM),
2 => (&NONLOAD_WEIGHTS_RW2, NONLOAD_WEIGHTS_RW2_SUM),
_ => (&NONLOAD_WEIGHTS, 75),
}
}
/// Shadow instructions per hash (0 without a shadow).
pub fn shadow_instrs_per_hash(&self) -> usize {
self.shadow.map(|s| s.instrs_per_hash()).unwrap_or(0)
@ -675,25 +592,6 @@ impl LoadClass {
/// "mx4": the v3 construction; a trailing "m<mult>" and "g" set the mixer multiplier and the growth rule on any
/// load class, "w16m4g" for example).
pub fn parse(s: &str) -> Option<LoadClass> {
// class v6 lane 1: "<class>+fold" (the index fold) and "<class>+rw" / "<class>+rw2" (the re-weight table) over
// any class, in any order, outermost of all
if let Some(base) = s.strip_suffix("+fold") {
return Some(LoadClass::parse(base)?.with_fold());
}
if let Some(base) = s.strip_suffix("+rw2") {
return Some(LoadClass::parse(base)?.with_rw(2));
}
if let Some(base) = s.strip_suffix("+rw") {
return Some(LoadClass::parse(base)?.with_rw(1));
}
// "<class>+reg64c": the 64-register window with the full-chain address mix (outermost, a research class)
if let Some(base) = s.strip_suffix("+reg64c") {
return Some(LoadClass::parse(base)?.with_reg64().with_reg64_chain());
}
// "<class>+reg64": the 64-register window over any class (the suffix is outermost, a research class)
if let Some(base) = s.strip_suffix("+reg64") {
return Some(LoadClass::parse(base)?.with_reg64());
}
// "<class>+state": the state leaves of class v5 over any class (the suffix is outermost)
if let Some(base) = s.strip_suffix("+state") {
return Some(LoadClass::parse(base)?.with_state());
@ -793,7 +691,6 @@ impl LoadClass {
"w4" => [100, 0, 0],
"w16" => [0, 100, 0],
"w64" => [0, 0, 100],
"w32" => return Some(LoadClass::fixed8(slots)),
m => {
let v: Vec<u8> = m.split(',').map(|x| x.trim().parse::<u8>().ok()).collect::<Option<Vec<_>>>()?;
if v.len() != 3 || v.iter().map(|&x| x as u32).sum::<u32>() != 100 {
@ -822,18 +719,6 @@ impl LoadClass {
/// An era class is the base name with "-era<first stream word as hex>" appended ("w4-era401998a5", "mx4-era...").
/// A hot class appends "hot<S>k<k>[a]" ("hot64k4", "scr4k32+hot64k4a"; measured and not adopted).
pub fn name(&self) -> String {
// class v6 lane 1: "+fold" then "+rw" / "+rw2" are the outermost suffixes ("mx8+sh256x27+state+fold+rw")
if self.rw != 0 {
return format!("{}+rw{}", LoadClass { rw: 0, ..*self }.name(), if self.rw == 1 { "" } else { "2" });
}
if self.fold {
return format!("{}+fold", LoadClass { fold: false, ..*self }.name());
}
if self.reg64 {
// "<class>+reg64" / "+reg64c": the 64-register window is a suffix on any class, outermost
let base = LoadClass { reg64: false, reg64_chain: false, ..*self }.name();
return if self.reg64_chain { format!("{base}+reg64c") } else { format!("{base}+reg64") };
}
if self.state {
// "<class>+state": class v5's leaves are a suffix on any class, outermost
return format!("{}+state", LoadClass { state: false, ..*self }.name());
@ -878,7 +763,6 @@ impl LoadClass {
format!("scr{k}k{}", self.scratch_kb)
} else {
let base = match self.mix {
[0, 100, 0] if self.wide8 => "w32".to_string(),
[100, 0, 0] => "w4".to_string(),
[0, 100, 0] => "w16".to_string(),
[0, 0, 100] => "w64".to_string(),
@ -907,15 +791,15 @@ impl LoadClass {
for (i, &m) in self.mix.iter().enumerate() {
acc += m as u64;
if roll < acc {
return self.widths()[i];
return WIDTH_WORDS[i];
}
}
self.widths()[2]
WIDTH_WORDS[2]
}
/// Expected dataset bytes read per hash: dataset loads per hash times the mean width (scratch traffic apart).
pub fn expected_bytes_per_hash(&self) -> f64 {
let mean = self.mix.iter().zip(self.widths().iter()).map(|(&m, &w)| m as f64 / 100.0 * w as f64 * 4.0).sum::<f64>();
let mean = self.mix.iter().zip(WIDTH_WORDS.iter()).map(|(&m, &w)| m as f64 / 100.0 * w as f64 * 4.0).sum::<f64>();
(self.load_slots as usize - self.scratch_slots()) as f64 * ITERATIONS as f64 * mean
}
}
@ -1124,7 +1008,7 @@ impl ProgramClass {
/// The program class whose load class `class` is, the era draw set aside: [`LoadClass::V2`] is v2, [`V3_CLASS`]
/// is v3, [`V4_CLASS`] is v4; a measurement class (a width, a derivation length, another shadow size) is none.
pub fn of_load_class(class: &LoadClass) -> Option<ProgramClass> {
let base = LoadClass { era: None, reg64: false, reg64_chain: false, ..*class };
let base = LoadClass { era: None, ..*class };
if base == LoadClass::V2 {
Some(ProgramClass::V2)
} else if base == V3_CLASS {
@ -1148,52 +1032,7 @@ impl ProgramClass {
}
}
/// Registers per lane of a program: 64 under the reg64 flag, 8 for every other class.
pub const REG64_REGISTERS: usize = 64;
/// The reg64 full-chain fold in the program id: 1 = the load's own source out of the rotate-xor chain.
pub const REG64_CHAIN_FOLD: u8 = 1;
impl Program {
/// Registers per lane (8, or 64 under `class.reg64`).
pub fn registers(&self) -> usize {
if self.class.reg64 { REG64_REGISTERS } else { 8 }
}
/// Whether every load's address consumes all 64 registers (the reg64 full-chain variant).
pub fn address_mix(&self) -> bool {
self.class.reg64 && self.class.reg64_chain
}
/// The pack's liveness statement (the reg64 variants): which reads keep every register necessary.
pub fn reg64_liveness(&self) -> String {
if !self.class.reg64 {
return String::new();
}
let loads = self.scheduled().iter().filter(|i| i.op == Op::Load).count();
if self.address_mix() {
format!("every one of the 64 registers is read by the address of each of the {loads} loads per iteration (the load's source directly, the 63 others through the rotate-xor chain) and by the end fold; 64 independently necessary values for the length of the dependent chain; xoring any register with either of two seed-derived probe words at the start of an iteration moves the first load address and the final hash")
} else {
"every one of the 64 registers is read by the end fold (r[k] ^= r[k + 8] ^ ... ^ r[k + 56]); a load's address reads its own window's register only, so the chain needs 8 live values at a time and the other 56 are live across it as state (window, arithmetic-only)".to_string()
}
}
/// The instructions one iteration executes, in order, with their register fields as the kernels name them.
/// Without the reg64 flag this is the drawn program as it stands. Under the flag, instruction i of the drawn
/// program runs on window 0 with every register field widened by `8 * (i % 4)` (so the 64 instructions touch all
/// 32 registers of the window), then the same instruction runs on window 1 (the widened field + 32): 128
/// statements per iteration over registers 0..63. The CPU verifier and every emitter read this list, so the
/// vectors and the kernel text agree by construction.
pub fn scheduled(&self) -> Vec<Instr> {
if !self.class.reg64 {
return self.instrs.clone();
}
let mut out = Vec::with_capacity(self.instrs.len() * 2);
for (i, ins) in self.instrs.iter().enumerate() {
let off = (8 * (i % 4)) as u8;
let a = Instr { dst: ins.dst + off, src: ins.src + off, src2: ins.src2 + off, ..*ins };
let b = Instr { dst: a.dst + 32, src: a.src + 32, src2: a.src2 + 32, ..a };
out.push(a);
out.push(b);
}
out
}
pub fn loads_per_hash(&self) -> usize {
self.instrs.iter().filter(|i| i.op.is_load()).count() * ITERATIONS
}
@ -1254,7 +1093,7 @@ impl Program {
pub fn width_counts(&self) -> [usize; 3] {
let mut c = [0usize; 3];
for i in self.instrs.iter().filter(|i| i.op == Op::Load) {
if let Some(k) = self.class.widths().iter().position(|&w| w == i.width) {
if let Some(k) = WIDTH_WORDS.iter().position(|&w| w == i.width) {
c[k] += 1;
}
}
@ -1298,7 +1137,7 @@ impl Program {
let v4_rung_0 = self.generator == GENERATOR_VERSION_V4 && LoadClass { era: None, ..self.class } == V4_CLASS;
// class v5 at rung 0 is `program_id(5, seed, attempt)`; above rung 0 the class-bearing id with "state/"
let v5_rung_0 = self.generator == GENERATOR_VERSION_V5 && LoadClass { era: None, ..self.class } == V5_CLASS;
if !self.class.reg64 && (self.class.is_v2() || self.generator == GENERATOR_VERSION_V3 || v4_rung_0 || v5_rung_0) {
if self.class.is_v2() || self.generator == GENERATOR_VERSION_V3 || v4_rung_0 || v5_rung_0 {
// Spec 01 section 1.4.6: a class v3 program's id is `program_id(3, seed, attempt)`, a class v4 program's
// `program_id(4, seed, attempt)` (Counter ASIC 3.0); the generator version in the preimage separates
// them from every version 2 program of the same seed
@ -1308,20 +1147,6 @@ impl Program {
}
}
/// The derivation of [`Program::program_id`] as text, from the same byte recipe (program.json's
/// "program_id_derivation"; spec 01 section 1.4.6).
pub fn program_id_derivation(&self) -> String {
let v4_rung_0 = self.generator == GENERATOR_VERSION_V4 && LoadClass { era: None, ..self.class } == V4_CLASS;
// class v5 at rung 0 takes the plain form as `program_id` does (the text follows the id; the first form of this
// function named the class recipe for the v5 packs whose id was the plain one)
let v5_rung_0 = self.generator == GENERATOR_VERSION_V5 && LoadClass { era: None, ..self.class } == V5_CLASS;
if !self.class.reg64 && (self.class.is_v2() || self.generator == GENERATOR_VERSION_V3 || v4_rung_0 || v5_rung_0) {
program_id_recipe(self.generator, &self.seed, self.attempt).text()
} else {
program_id_class_recipe(self.generator, &self.seed, self.attempt, &self.class).text()
}
}
/// The program class of this program, from its generator version (3 = v3, 4 = v4, everything else v2).
pub fn program_class(&self) -> ProgramClass {
match self.generator {
@ -1333,61 +1158,24 @@ impl Program {
}
}
/// The byte recipe of a program id and the text that states it: the bytes FNV-1a 64 hashes and, part by part, the label
/// of each part (`'literal'` or `field_le32`), so the derivation printed in program.json is built from the same list the id
/// is hashed from and the two cannot drift (the documentation finding of 7 October 2026: program.json and spec 1.4.6 said
/// the plain form while generator 4 appended `'sub/' || sub_version_le16`, so a client written from the text derived another
/// id). `tests/derivation.rs` re-derives every pinned pack's id from its own derivation string.
#[derive(Clone, Debug, Default)]
pub struct IdRecipe {
pub bytes: Vec<u8>,
parts: Vec<String>,
}
impl IdRecipe {
fn lit(&mut self, s: &[u8]) {
self.bytes.extend_from_slice(s);
self.parts.push(format!("'{}'", String::from_utf8_lossy(s)));
}
fn field(&mut self, bytes: &[u8], label: &str) {
self.bytes.extend_from_slice(bytes);
self.parts.push(label.to_string());
}
/// The derivation as program.json prints it: `FNV-1a 64 over <part> || <part> || ...`.
pub fn text(&self) -> String {
format!("FNV-1a 64 over {}", self.parts.join(" || "))
}
/// The id: FNV-1a 64 over the bytes.
pub fn id(&self) -> u64 {
fnv1a64(&self.bytes)
}
}
/// The recipe of [`program_id`].
pub fn program_id_recipe(generator: u32, seed: &[u32; 8], attempt: u32) -> IdRecipe {
let mut r = IdRecipe::default();
r.lit(PROGRAM_ID_TAG);
r.field(&generator.to_le_bytes(), "generator_le32");
let mut sw = Vec::with_capacity(32);
pub fn program_id(generator: u32, seed: &[u32; 8], attempt: u32) -> u64 {
let mut b = Vec::with_capacity(PROGRAM_ID_TAG.len() + 4 + 32 + 4 + 6);
b.extend_from_slice(PROGRAM_ID_TAG);
b.extend_from_slice(&generator.to_le_bytes());
for w in seed {
sw.extend_from_slice(&w.to_le_bytes());
b.extend_from_slice(&w.to_le_bytes());
}
r.field(&sw, "seed_words as little-endian bytes");
r.field(&attempt.to_le_bytes(), "attempt_le32");
b.extend_from_slice(&attempt.to_le_bytes());
if generator == GENERATOR_VERSION_V4 {
// The class v4 sub-version (AP-F8-1 amendment, 7 October 2026): `"sub/" || sub_version as little-endian u16`
// appended for generator 4 only, so a binary from before the load-source rule (sub-version 0, no suffix) and
// one after it never share a program id for one seed; the node's id check then catches a split. v2 and v3
// ids are byte-identical. The node reads the sub-version from [`PROGRAM_SUBVERSION_V4`]; packs carry it as
// IGNEUM_PROGRAM_SUBVERSION and program.json "sub_version".
r.lit(b"sub/");
r.field(&PROGRAM_SUBVERSION_V4.to_le_bytes(), "sub_version_le16");
b.extend_from_slice(b"sub/");
b.extend_from_slice(&PROGRAM_SUBVERSION_V4.to_le_bytes());
}
r
}
pub fn program_id(generator: u32, seed: &[u32; 8], attempt: u32) -> u64 {
program_id_recipe(generator, seed, attempt).id()
fnv1a64(&b)
}
/// The sub-version of class v4's program stream, in every generator-4 program id and pack (AP-F8-3: 3 = the
@ -1401,84 +1189,57 @@ pub const PROGRAM_ID_TAG_RW: &[u8] = b"igneum-program-rw/";
/// The program id of a non-default class: the tag, then the same fields as [`program_id`], then the three mix
/// percentages and the slot count as bytes.
/// The recipe of [`program_id_class`].
pub fn program_id_class_recipe(generator: u32, seed: &[u32; 8], attempt: u32, class: &LoadClass) -> IdRecipe {
let mut r = IdRecipe::default();
r.lit(PROGRAM_ID_TAG_RW);
r.field(&generator.to_le_bytes(), "generator_le32");
let mut sw = Vec::with_capacity(32);
pub fn program_id_class(generator: u32, seed: &[u32; 8], attempt: u32, class: &LoadClass) -> u64 {
let mut b = Vec::with_capacity(PROGRAM_ID_TAG_RW.len() + 4 + 32 + 4 + 4);
b.extend_from_slice(PROGRAM_ID_TAG_RW);
b.extend_from_slice(&generator.to_le_bytes());
for w in seed {
sw.extend_from_slice(&w.to_le_bytes());
b.extend_from_slice(&w.to_le_bytes());
}
r.field(&sw, "seed_words as little-endian bytes");
r.field(&attempt.to_le_bytes(), "attempt_le32");
r.field(&class.mix, "mix[3]");
r.field(&[class.load_slots], "load_slots_u8");
b.extend_from_slice(&attempt.to_le_bytes());
b.extend_from_slice(&class.mix);
b.push(class.load_slots);
if let Some(k) = class.scratch {
r.lit(b"scratch/");
r.field(&[k], "scratch_k_u8");
r.field(&[class.scratch_kb], "scratch_kb_u8");
b.extend_from_slice(b"scratch/");
b.push(k);
b.push(class.scratch_kb);
}
if class.mixer_mult != 1 || class.growth {
// Counter ASIC 2.0: the mixer multiplier and the growth rule are part of the construction, so a program of
// the same seed under a different mixer carries a different id (under the v3 seam the id is
// program_id(3, seed, attempt) and this branch is not taken)
r.lit(b"mixer/");
r.field(&[class.mixer_mult], "mixer_mult_u8");
r.field(&[class.growth as u8], "growth_u8");
b.extend_from_slice(b"mixer/");
b.push(class.mixer_mult);
b.push(class.growth as u8);
}
if class.derive_len != 0 {
// Counter ASIC 3.0 item 2: the derivation program's length is part of the construction
r.lit(b"derive/");
r.field(&class.derive_len.to_le_bytes(), "derive_len_le16");
b.extend_from_slice(b"derive/");
b.extend_from_slice(&class.derive_len.to_le_bytes());
}
if let Some(e) = class.era {
r.lit(b"era/");
r.field(&e.id_bytes(), "era_id_bytes (allowed[3] || width_words_u8 || stride_mul_le32 || stride_rot_le32 || interleave[4])");
b.extend_from_slice(b"era/");
b.extend_from_slice(&e.id_bytes());
}
if let Some(sh) = class.shadow {
// Counter ASIC 3.0 item 8: the shadow block's size and repeat count are part of the construction
r.lit(b"shadow/");
r.field(&sh.instrs.to_le_bytes(), "shadow_instrs_le16");
r.field(&sh.reps.to_le_bytes(), "shadow_reps_le16");
b.extend_from_slice(b"shadow/");
b.extend_from_slice(&sh.instrs.to_le_bytes());
b.extend_from_slice(&sh.reps.to_le_bytes());
}
if let Some(h) = class.hot {
r.lit(b"hot/");
r.field(&[h.mb], "hot_mb_u8");
r.field(&[h.k], "hot_k_u8");
b.extend_from_slice(b"hot/");
b.push(h.mb);
b.push(h.k);
if h.added {
r.lit(b"added");
b.extend_from_slice(b"added");
}
}
if class.state {
// class v5: the state leaves are part of the construction
r.lit(b"state/");
b.extend_from_slice(b"state/");
}
if class.fold {
// class v6 lane 1: the index fold moves every era load address
r.lit(b"fold/");
}
if class.rw != 0 {
// class v6 lane 1: the re-weight table moves the op draw
r.lit(b"rw/");
r.field(&[class.rw], "rw_table_u8");
}
if class.reg64 {
// the 64-register window is part of the construction
r.lit(b"reg64/");
if class.reg64_chain {
// the fold text is part of the construction: 1 = the load's own source out of the chain (the sound
// fold); the benched text of 8 October 2026 (the source inside the chain, id 0x3deee2320e70e1bf for the
// pinned devnet seeds) carried no byte here, so the two texts never share an id
r.lit(b"chain/");
r.field(&[REG64_CHAIN_FOLD], "reg64_chain_fold_u8");
}
}
r
}
pub fn program_id_class(generator: u32, seed: &[u32; 8], attempt: u32, class: &LoadClass) -> u64 {
program_id_class_recipe(generator, seed, attempt, class).id()
fnv1a64(&b)
}
/// Weights of the ten non-load families under version 2, in draw order. Sum 75. The load family has no
@ -1496,38 +1257,6 @@ pub const NONLOAD_WEIGHTS: [(Op, u64); 10] = [
(Op::Or, 4),
];
/// Class v6 lane 1, the re-weight table behind the index fold (`+rw`): the k lane's optimiser split in draw order,
/// sum [`NONLOAD_WEIGHTS_RW_SUM`] (83). Shuffle stays at 4 (a shuffle-heavy draw is the worst thing the class can do
/// on the chip side, k 0.021 routed) and `or` is never drawn. The roll ranges over the table's own sum.
pub const NONLOAD_WEIGHTS_RW: [(Op, u64); 10] = [
(Op::Add, 16),
(Op::Xor, 14),
(Op::Mul, 4),
(Op::Mad, 12),
(Op::Shfl, 4),
(Op::Rotl, 11),
(Op::Sub, 10),
(Op::MulHi, 2),
(Op::Rotr, 10),
(Op::Or, 0),
];
pub const NONLOAD_WEIGHTS_RW_SUM: u64 = 83;
/// Class v6 lane 1, the census lane's neighbouring table (`+rw2`), sum [`NONLOAD_WEIGHTS_RW2_SUM`] (75).
pub const NONLOAD_WEIGHTS_RW2: [(Op, u64); 10] = [
(Op::Add, 13),
(Op::Xor, 11),
(Op::Mul, 6),
(Op::Mad, 10),
(Op::Shfl, 8),
(Op::Rotl, 8),
(Op::Sub, 7),
(Op::MulHi, 2),
(Op::Rotr, 6),
(Op::Or, 4),
];
pub const NONLOAD_WEIGHTS_RW2_SUM: u64 = 75;
/// Version 1 weights (retired). Sum 100, load at 25 percent.
pub const OP_WEIGHTS: [(Op, u64); 11] = [
(Op::Load, 25),
@ -1649,12 +1378,8 @@ pub fn candidate_from_words_class(
// era set aside) on EVERY draw path, era or not, so a census through candidate_class reads the same stream as
// the chain; v2, v3 and every other class take no part. The draw order and the stream are otherwise the same.
// class v5 (docs/design/class-v5-stored-state.md) draws under the same rule: its state flag is set aside here too
// class v6 lane 1: the index fold and the re-weight table are set aside too (the address path and the table are not
// the shape; a +fold or +rw program draws its sources under the same rule)
let source_rule_v4 = matches!(class.shadow, Some(ShadowClass { instrs: V4_SHADOW_INSTRS, .. }))
&& LoadClass { era: None, shadow: None, state: false, fold: false, rw: 0, ..class } == LoadClass { shadow: None, ..V4_CLASS };
// the op table and the roll's range: the plain table at 75 for every class without the re-weight flag
let (weights, weights_sum) = class.nonload_weights();
&& LoadClass { era: None, shadow: None, state: false, ..class } == LoadClass { shadow: None, ..V4_CLASS };
let mut fresh = [false; 8];
let mut fresh_value = [true; 8];
// the shared-operand idiom (AP-F8-1, sub-version 3): after `or d |= s`, a later `xor d ^= s` or `sub d -= s` with
@ -1664,9 +1389,9 @@ pub fn candidate_from_words_class(
let mut pair_op: [Option<(Op, usize)>; 8] = [None; 8];
let mut instrs = Vec::with_capacity(INSTR_COUNT);
for k in 0..INSTR_COUNT {
let mut roll = rng.below(weights_sum);
let mut roll = rng.below(75);
let mut op = Op::Add;
for &(o, w) in weights {
for &(o, w) in &NONLOAD_WEIGHTS {
if roll < w {
op = o;
break;
@ -1771,9 +1496,9 @@ pub fn candidate_from_words_class(
loop {
shadow.clear();
for _ in 0..sh.instrs {
let mut roll = rng.below(weights_sum);
let mut roll = rng.below(75);
let mut op = Op::Add;
for &(o, w) in weights {
for &(o, w) in &NONLOAD_WEIGHTS {
if roll < w {
op = o;
break;
@ -1864,12 +1589,7 @@ pub fn try_generate_class(seed_string: &str, seed_bytes: &[u8], class: LoadClass
}
if crate::accept::is_class_v4_shape(&class) {
// AP-F8-2 (7 October 2026, main's ruling: the draw is total and no consensus path panics): a class v4 seed that
// exhausts its attempts takes the last-resort program, deterministic. Sub-version 3's is accepted as drawn
// (unreachable at 4.6e-44 per epoch and unverified against the rule: adv-accept-3's finding, 223 of 2,500
// rewritten candidates fail it, 209 by part (a)); class v5's is the verified one.
if class.state {
return Ok(last_resort_v5(seed_string, seed_bytes, cap, class));
}
// exhausts its attempts takes the last-resort program, deterministic and accepted as drawn
return Ok(last_resort_v4(candidate_class(seed_string, seed_bytes, cap, class)));
}
Err(Exhausted { seed_string: seed_string.to_string(), attempts: cap, last: last.unwrap() })
@ -1893,55 +1613,6 @@ pub fn last_resort_v4(mut p: Program) -> Program {
p
}
/// Class v5's last resort, verified (adv-accept-3's finding of 7 October 2026: the sub-version 3 rewrite alone fails the
/// rule on 223 of 2,500 seeds, 209 by part (a), a load reading a register no instruction wrote since the previous load
/// from it). From attempt `cap` on, each candidate is rewritten as [`last_resort_v4`] does, its stale loads re-sourced
/// by [`repair_stale_loads`], and the result checked against the whole rule; the first that passes is the program. The
/// scan runs [`LAST_RESORT_SCAN`] candidates; past it the first repaired candidate stands as drawn, so the draw is total,
/// and that fallback sits behind the 4.6e-44 of reaching the last resort at all times the rejection rate of a repaired
/// candidate (about 0.09 per adv-accept-3) to the power of the scan: under 1e-300.
pub fn last_resort_v5(seed_string: &str, seed_bytes: &[u8], cap: u32, class: LoadClass) -> Program {
for k in cap..cap + LAST_RESORT_SCAN {
let p = repair_stale_loads(last_resort_v4(candidate_class(seed_string, seed_bytes, k, class)));
if check(&p).is_ok() {
return p;
}
}
repair_stale_loads(last_resort_v4(candidate_class(seed_string, seed_bytes, cap, class)))
}
/// The candidates class v5's last resort scans past the attempt cap before the unchecked fallback.
pub const LAST_RESORT_SCAN: u32 = 256;
/// Part (a) by construction: a load whose source register no instruction wrote since the previous load from it (in
/// cyclic order over the 64 instructions, the two-pass walk of the acceptance's `check_stale_loads`) is re-sourced to
/// the lowest register that was written since its last load. The walk repeats until a full two-pass walk changes
/// nothing (a re-sourced load can make a later load from the new register stale, which the next walk re-sources);
/// with 16 loads among 64 instructions a non-pending register always exists. Only the `src` field moves.
pub fn repair_stale_loads(mut p: Program) -> Program {
for _round in 0..16 {
let mut changed = false;
let mut pending = [false; 8];
for _pass in 0..2 {
for ins in p.instrs.iter_mut() {
if ins.op.is_load() && pending[ins.src as usize] {
let q = (0..8usize).find(|&q| !pending[q]).expect("a register written since its last load") as u8;
ins.src = q;
changed = true;
}
pending[ins.dst as usize] = false;
if ins.op.is_load() {
pending[ins.src as usize] = true;
}
}
}
if !changed {
break;
}
}
p
}
/// [`try_generate_from_seed_bytes`], treating exhaustion as the consensus fault it is.
pub fn generate_from_seed_bytes(seed_string: &str, seed_bytes: &[u8]) -> Program {
try_generate_from_seed_bytes(seed_string, seed_bytes).unwrap_or_else(|e| panic!("{e}"))
@ -2227,14 +1898,6 @@ mod tests {
assert!(v2.instrs.iter().all(|i| i.width == 1));
assert_eq!(v2.bytes_per_hash(), 512);
assert_eq!(LoadClass::parse("w16"), Some(LoadClass::fixed(4, 16)));
// W = 8 (8 October 2026): the w32 class reads 8 words a load; w16 reads 4 (known-failed first: the plain set has no 8)
assert!(!WIDTH_WORDS.contains(&8));
let w32 = LoadClass::parse("w32").unwrap();
assert!(w32.wide8 && w32.mix == [0, 100, 0] && w32.widths() == [1, 8, 16]);
assert_eq!(w32.name(), "w32");
assert_eq!(LoadClass::parse("w32x8").unwrap().load_slots, 8);
assert_eq!(w32.expected_bytes_per_hash(), 2.0 * LoadClass::fixed(4, 16).expected_bytes_per_hash());
assert!(!LoadClass::fixed(4, 16).wide8);
assert_eq!(LoadClass::parse("w64x4"), Some(LoadClass::fixed(16, 4)));
assert_eq!(LoadClass::parse("50,35,15"), Some(LoadClass::mixed([50, 35, 15])));
assert_eq!(LoadClass::parse("v2"), Some(LoadClass::V2));
@ -2610,44 +2273,6 @@ mod tests {
}
}
/// Class v5's last resort is verified (adv-accept-3's finding, 7 October 2026): on seed adv3/steer/2 of that lane's
/// label space (`seed_words_from_bytes("igneum-adv-accept-3/steer/2")` as the epoch bytes, Devnet 3's genesis as the
/// era) the sub-version 3 rewrite of the candidate at the cap fails the rule by part (a), a cyclic stale load, and
/// would be handed to the chain; class v5's repair re-sources the load and the result passes the whole rule, as
/// does the program `last_resort_v5` returns. The class v4 path is unchanged (frozen, unreachable, recorded).
#[test]
fn class_v5_last_resort_is_verified_known_failed_adv3_steer_2() {
use crate::accept::{check, check_static, Reject};
use crate::seed::seed_words_from_bytes;
let epoch: Vec<u8> = seed_words_from_bytes(b"igneum-adv-accept-3/steer/2").iter().flat_map(|x| x.to_le_bytes()).collect();
let era = crate::bind::unhex("4020cb4382e3fe4b281c817c02582e147d8f851f566ae9172b28912b8e68b925").unwrap();
let v4 = LoadClass::era(V4_CLASS, &era, &V3_ALLOWED);
let v5 = LoadClass::era(V5_CLASS, &era, &V3_ALLOWED);
let cap = max_attempts_for(&v4);
let old = last_resort_v4(candidate_class("adv3/steer/2", &epoch, cap, v4));
match check_static(&old) {
Err(Reject::StaleLoadSource { instr, reg }) => println!("adv3/steer/2: the sub-version 3 last resort fails part (a) at instruction {instr} reading r{reg}"),
other => panic!("the known-failed case must fail part (a): {other:?}"),
}
let repaired = repair_stale_loads(old.clone());
assert!(check_static(&repaired).is_ok(), "the repair restores part (a): {:?}", check_static(&repaired).err());
assert_eq!(repaired.instrs.len(), old.instrs.len());
assert!(repaired.instrs.iter().zip(old.instrs.iter()).all(|(a, b)| a.op == b.op && a.dst == b.dst), "only load sources move");
let lr = last_resort_v5("adv3/steer/2", &epoch, cap, v5);
assert!(check(&lr).is_ok(), "class v5's last resort passes the whole rule: {:?}", check(&lr).err());
assert!(lr.attempt >= cap && lr.attempt < cap + LAST_RESORT_SCAN);
assert!(lr.class.state);
println!("adv3/steer/2: class v5 last resort at attempt {} id {:016x}", lr.attempt, lr.program_id());
// the sub-version 3 path is byte for byte what it was
assert_eq!(try_generate_class("adv3/steer/2", &epoch, v4).map(|p| p.attempt).ok(), Some(try_generate_class("adv3/steer/2", &epoch, v4).unwrap().attempt));
// a few more of the label space: every class v5 last resort passes
for i in [11u32, 33, 56, 58, 77] {
let e: Vec<u8> = seed_words_from_bytes(format!("igneum-adv-accept-3/steer/{i}").as_bytes()).iter().flat_map(|x| x.to_le_bytes()).collect();
let lr = last_resort_v5(&format!("adv3/steer/{i}"), &e, cap, v5);
assert!(check(&lr).is_ok(), "steer/{i}: {:?}", check(&lr).err());
}
}
#[test]
fn class_v4_draw_is_total_with_the_last_resort() {
assert_eq!(max_attempts_for(&V4_CLASS), MAX_ATTEMPTS_V4);

59
igneum-pow/src/main.rs
View file

@ -1,7 +1,7 @@
//! igneum-pow CLI.
//!
//! igneum-pow bench --seed <s> [--day <d>] [--closed-form] [--dataset-log2 28] [--warps 20]
//! igneum-pow export --seed <s> --out <dir> [--day <d>] [--closed-form] [--dataset-log2 28 | --dataset-words N] [--epoch-hex <64 hex> --day-hex <hex>]
//! igneum-pow export --seed <s> --out <dir> [--day <d>] [--closed-form] [--dataset-log2 28] [--epoch-hex <64 hex> --day-hex <hex>]
//! igneum-pow hash --seed <s> --nonce <n> [--day <d>] [--closed-form] [--dataset-log2 28]
//! igneum-pow hash-bound --seed <s> --prehash <64 hex> --nonce <u64> [--day <d>] [--closed-form] [--dataset-log2 28]
//! [--epoch-hex <64 hex> --day-hex <hex>] byte seeds instead of strings (Epoch::from_seed_bytes)
@ -22,7 +22,7 @@ use igneum_pow::emit::export_pack;
use igneum_pow::generator::{LoadClass, ProgramClass};
use igneum_pow::memhard::{Cache, Shape};
use igneum_pow::seed::day_key;
use igneum_pow::verify::{DatasetGeom, DatasetMode, DatasetSource, Epoch, DEFAULT_DATASET_LOG2};
use igneum_pow::verify::{DatasetMode, DatasetSource, Epoch, DEFAULT_DATASET_LOG2};
use std::time::Instant;
struct Args {
@ -32,10 +32,6 @@ struct Args {
out: Option<String>,
closed_form: bool,
dataset_log2: u32,
/// `--dataset-words N`: the dataset at N words (research class ds55, 8 October 2026): any multiple of 2^16 in
/// 2^28 ..= 2^31; a non-power-of-two takes the multiply-shift range reduction of spec 01 section 1.13.3 in every
/// load, a power of two is the `--dataset-log2` path. Applied after the class's own sizing.
dataset_words: Option<DatasetGeom>,
warps: usize,
nonce: u64,
/// `hash-bound --count N`: N consecutive nonces from --nonce, one epoch build (gate G2, 5 October 2026).
@ -60,12 +56,6 @@ struct Args {
/// generator 3 and the era bytes recorded (the measurement packs: v2's mixer under the era layout).
era: Option<(u64, Vec<u8>, String)>,
era_widths: Vec<u8>,
/// `--reg64`: the 64-register window per lane over the chosen class (the hash lane's measurement, 8 October
/// 2026, a research class; also `--class <class>+reg64`). The draw is the class's; the flag is stamped on the
/// program after it, as the era is.
reg64: bool,
/// `--reg64-chain`: the full-chain address mix over the reg64 window (`+reg64c`).
reg64_chain: bool,
}
/// `igneum-era-test/<n>` or `<n>:<64 hex>` -> (index, 32 era bytes, label).
@ -90,7 +80,6 @@ fn parse_widths(s: &str) -> Option<Vec<u8>> {
.map(|x| match x.trim() {
"4" => Some(1u8),
"16" => Some(4),
"32" => Some(8), // W = 8, the 32-byte sector (the w32 class, 8 October 2026)
"64" => Some(16),
_ => None,
})
@ -119,10 +108,7 @@ fn usage() -> ! {
\x20 --state <file> class v5 (or any --class ...+state): the window's state stream (IGSD1 file, igneum-day-stream --out), whose leaves key every item\n\
\x20 --shadow-reps N class v4 at a rung of the latency ladder: the shadow block's pass count (0 = the class's own 27; docs/design/latency-ladder.md), with --program-class v4\n\
\x20 --era E era layout over --class: igneum-era-test/<n> or <n>:<64 hex> (the 32-byte era seed E_n)\n\
\x20 --era-widths 4[,16,32,64] the width set the era draws from, in bytes (default 4: pinned; more lets the era draw it; 32 only with the w32 class)\n\
\x20 --dataset-words N research class ds55: the dataset at N words (a multiple of 65,536 in 2^28 ..= 2^31; 1476395008 = 5.5 GiB); a non-power-of-two uses idx = (src * N) >> 32 in every load (spec 01 section 1.13.3), a power of two is --dataset-log2\n\\
\x20 --reg64 the 64-register window per lane over the class (research; the same as --class <class>+reg64; CUDA, OpenCL and Metal texts)\n\
\x20 --reg64-chain reg64 with the full-chain address mix: every load's address consumes all 64 registers (the same as --class <class>+reg64c)"
\x20 --era-widths 4[,16,64] the width set the era draws from, in bytes (default 4: pinned; more lets the era draw it)"
);
std::process::exit(2)
}
@ -136,7 +122,6 @@ fn parse() -> Args {
closed_form: false,
state: None,
dataset_log2: DEFAULT_DATASET_LOG2,
dataset_words: None,
warps: 20,
nonce: 0,
count: 1,
@ -150,8 +135,6 @@ fn parse() -> Args {
shadow_reps: 0,
era: None,
era_widths: vec![1],
reg64: false,
reg64_chain: false,
};
let mut it = std::env::args().skip(1);
a.cmd = it.next().unwrap_or_else(|| usage());
@ -163,13 +146,6 @@ fn parse() -> Args {
"--out" => a.out = Some(val()),
"--closed-form" => a.closed_form = true,
"--dataset-log2" => a.dataset_log2 = val().parse().unwrap_or_else(|_| usage()),
"--dataset-words" => {
let n: u64 = val().parse().unwrap_or_else(|_| usage());
a.dataset_words = Some(DatasetGeom::words(n).unwrap_or_else(|err| {
eprintln!("{err}");
std::process::exit(2)
}));
}
"--warps" => a.warps = val().parse().unwrap_or_else(|_| usage()),
"--nonce" => a.nonce = val().parse().unwrap_or_else(|_| usage()),
"--count" => a.count = val().parse().unwrap_or_else(|_| usage()),
@ -184,11 +160,6 @@ fn parse() -> Args {
"--shadow-reps" => a.shadow_reps = val().parse().unwrap_or_else(|_| usage()),
"--era" => a.era = Some(parse_era(&val()).unwrap_or_else(|| usage())),
"--era-widths" => a.era_widths = parse_widths(&val()).unwrap_or_else(|| usage()),
"--reg64" => a.reg64 = true,
"--reg64-chain" => {
a.reg64 = true;
a.reg64_chain = true;
}
_ => usage(),
}
}
@ -251,22 +222,6 @@ fn main() {
fn epoch_of(a: &Args, mode: DatasetMode) -> (Epoch, String) {
let (mut e, label) = epoch_of_class(a, mode);
stamp_era(&mut e, a);
// research class ds55: the dataset at the word count of --dataset-words, the cache and the items unchanged; a
// state class sizes its leaves by the power-of-two count and is refused here
if let Some(geom) = a.dataset_words {
if e.program.class.state {
eprintln!("--dataset-words is not supported with a state class ({}): the leaves are sized by --dataset-log2", e.program.class.name());
std::process::exit(2);
}
e.dataset = e.dataset.with_geom(geom);
}
if a.reg64 {
// the 64-register window over the drawn program (the same text in CUDA, OpenCL and Metal since class-v6)
e.program.class = e.program.class.with_reg64();
if a.reg64_chain {
e.program.class = e.program.class.with_reg64_chain();
}
}
// class v5: the leaves of --state, built for the dataset's size; a state class without --state is refused here
// rather than at the first derivation
if e.program.class.state {
@ -426,10 +381,10 @@ fn export(a: &Args, mode: DatasetMode) {
let build_ms = t0.elapsed().as_secs_f64() * 1e3;
println!("igneum-pow export {out}");
println!(
"seed \"{}\", day \"{}\", dataset {} ({}), generator v{} attempt {} program id {:016x}, loads/hash {}; epoch built in {build_ms:.1} ms",
"seed \"{}\", day \"{}\", dataset 2^{} words ({}), generator v{} attempt {} program id {:016x}, loads/hash {}; epoch built in {build_ms:.1} ms",
e.program.seed_string,
day_label,
e.dataset.geom.describe(),
e.dataset.log2_words,
e.dataset.mode().name(),
e.program.generator,
e.program.attempt,
@ -437,10 +392,6 @@ fn export(a: &Args, mode: DatasetMode) {
e.program.loads_per_hash()
);
println!("op mix: {}; class {}, {} bytes/hash, widths (1,4,16 words) {:?}", e.program.op_mix(), e.program.class.name(), e.program.bytes_per_hash(), e.program.width_counts());
println!("seed words {}", e.program.seed.iter().map(|w| format!("0x{w:08x}")).collect::<Vec<_>>().join(" "));
if e.dataset.geom.mulshift {
println!("dataset mapping: multiply-shift, idx = (src * {}) >> 32; {} items, {} bytes (research class ds55)", e.dataset.geom.words, e.dataset.geom.items(), e.dataset.geom.bytes());
}
let source = format!("igneum-pow (Rust) CPU interpreter, generator v{}, {} dataset", e.program.generator, e.dataset.mode().name());
let pack = export_pack(&e, &day_label, &source);
let dir = std::path::Path::new(&out);

93
igneum-pow/src/memhard.rs
View file

@ -210,15 +210,12 @@ pub struct MixParams {
pub redraws: u32,
}
/// Class v5's mixer-draw rule (AP-F4-1, the form the attack-pass lane and adv-mixer-2 agreed on 7 October 2026, 22:0x
/// UK; it replaces the first form's "NAF sum under 163"): the adder-datapath cost of a mixer block is
/// `A = 64 + sum over the 16 multipliers of (w32(m) - 1)`, `w32` the NAF weight over bit positions 0 to 31 only (the
/// carry digit at position 32 a 32-bit multiplier never pays; the first census counted it, so its median read 231
/// where the agreed median is 226). A block whose cost is at most this is rejected (a 1.1x gain against the median).
pub const MIXER_COST_REJECT_MAX: u32 = 205;
/// Class v5's mixer-draw rule: every multiplier's `w32` at least this (a multiplier with a two-adder chain, `w32 <= 3`,
/// is rejected on its own: adv-mixer-2's `k >= 1`).
/// Class v5's mixer-draw rule (AP-F4-1): the NAF sum of the 16 multipliers at least this.
pub const MIXER_NAF_SUM_MIN: u32 = 163;
/// Class v5's mixer-draw rule: every multiplier's NAF weight at least this.
pub const MIXER_NAF_WORD_MIN: u32 = 4;
/// Class v5's mixer-draw rule: at least this many distinct rotation amounts among the eight.
pub const MIXER_DISTINCT_ROT_MIN: usize = 4;
/// Class v5's mixer-draw rule: redraws before the last block stands as drawn (never reached at 6.1e-4 per try).
pub const MIXER_REDRAW_CAP: u32 = 64;
@ -241,40 +238,14 @@ pub fn naf_weight(mut x: u64) -> u32 {
w
}
/// The NAF weight of a 32-bit multiplier over bit positions 0 to 31: [`naf_weight`] without the digit at position 32.
pub fn naf32_weight(m: u32) -> u32 {
let mut x = m as u64;
let mut w = 0;
let mut pos = 0;
while x != 0 {
if x & 1 == 1 {
if pos < 32 {
w += 1;
}
if x & 3 == 3 {
x += 1;
} else {
x -= 1;
}
}
x >>= 1;
pos += 1;
}
w
}
/// The adder-datapath cost of a mixer block: `64 + sum(w32(m) - 1)` over the 16 multipliers.
pub fn mixer_cost(mul: &[u32; 16]) -> u32 {
64 + mul.iter().map(|&m| naf32_weight(m).saturating_sub(1)).sum::<u32>()
}
/// Whether a mixer block passes class v5's draw rule (AP-F4-1): cost over [`MIXER_COST_REJECT_MAX`], every
/// multiplier's `w32` at least [`MIXER_NAF_WORD_MIN`], and the eight rotation amounts not all equal. A rejected block
/// is redrawn whole (all forty draws) from the continuing stream.
/// Whether a mixer block passes class v5's draw rule (AP-F4-1).
pub fn mixer_block_admissible(rot: &[u32; 8], mul: &[u32; 16]) -> bool {
let words = mul.iter().all(|&m| naf32_weight(m) >= MIXER_NAF_WORD_MIN);
let rot_equal = rot.iter().all(|&r| r == rot[0]);
mixer_cost(mul) > MIXER_COST_REJECT_MAX && words && !rot_equal
let sum: u32 = mul.iter().map(|&m| naf_weight(m as u64)).sum();
let words = mul.iter().all(|&m| naf_weight(m as u64) >= MIXER_NAF_WORD_MIN);
let mut distinct = rot.to_vec();
distinct.sort_unstable();
distinct.dedup();
sum >= MIXER_NAF_SUM_MIN && words && distinct.len() >= MIXER_DISTINCT_ROT_MIN
}
impl MixParams {
@ -298,11 +269,11 @@ impl MixParams {
}
let mut redraws = 0u32;
if shape.state {
// Class v5 (docs/design/class-v5-stored-state.md section 11, AP-F4-1, the attack-pass lane's and adv-mixer-2's
// reconciled weak-day census): a mixer block whose adder cost is at most 205 against the median 226, or with a
// two-adder multiplier, or with one rotation amount, is redrawn whole from the next forty stream values, so no
// day is a weak day for a per-day LUT-recompute FPGA (the worst calendar day, chain day 29,337 = 2050-04-28,
// read 1.113x). About 5.69e-4 of days redraw, 15 days a century. The derive program's draws come after.
// Class v5 (docs/design/class-v5-stored-state.md section 11, AP-F4-1, the attack-pass lane's weak-day census):
// a mixer block whose multipliers are cheap on an adder datapath (NAF sum under 163, a word under NAF weight 4)
// or whose rotations repeat (under 4 distinct amounts) is redrawn from the next stream values, so no day is a
// weak day for a per-day LUT-recompute FPGA (the worst calendar day of the census, chain day 29,337, was 1.121x).
// About 6.1e-4 of days redraw. The derive program's draws (none under v5) come after, as before.
while !mixer_block_admissible(&rot, &mul) && redraws < MIXER_REDRAW_CAP {
for r in rot.iter_mut() {
*r = 1 + rng.below(31) as u32;
@ -847,11 +818,9 @@ impl MemhardCpu {
mod tests {
use super::*;
/// Class v5's mixer-draw rule (AP-F4-1 in the agreed form), the known-failed case first: the day the two censuses
/// name as the worst of the century, chain day 29,337 (2050-04-28, 1.113x), draws a block the rule rejects and class
/// v5 redraws it; a block of two-adder multipliers or one rotation amount is inadmissible; a scan of day keys finds
/// the redrawn days (about 5.69e-4), every v5 block passes after the draw, and the v4 constants of the same keys
/// never move.
/// Class v5's mixer-draw rule (AP-F4-1), the known-failed case first: a block of cheap multipliers (NAF sum under
/// 163) or repeated rotations is inadmissible; a scan of day keys finds days the rule redraws (the census's 6.1e-4),
/// every v5 block passes after the draw, and the v4 constants of the same keys never move.
#[test]
fn class_v5_mixer_draw_rule() {
assert_eq!(naf_weight(0), 0);
@ -859,34 +828,16 @@ mod tests {
assert_eq!(naf_weight(3), 2, "11 = 100 - 1");
assert_eq!(naf_weight(7), 2, "111 = 1000 - 1");
assert_eq!(naf_weight(0xffff_ffff), 2);
assert_eq!(naf32_weight(0xffff_ffff), 1, "the carry digit at position 32 is not paid");
assert_eq!(naf32_weight(0xc000_0001), 2, "2^32 - 2^30 + 1: the digit at 32 dropped, -2^30 and +1 kept");
assert_eq!(naf_weight(0b1010_1010), 4);
assert_eq!(naf32_weight(0b1010_1010), 4);
let good_rot = [1u32, 5, 9, 13, 17, 21, 25, 29];
let cheap = [0x8000_0001u32; 16];
assert_eq!(mixer_cost(&cheap), 64 + 16, "16 two-adder multipliers");
assert!(!mixer_block_admissible(&good_rot, &cheap), "the known-failed case: 16 two-adder multipliers");
let dense = [0xaaaa_aaabu32; 16];
assert!(mixer_cost(&dense) > MIXER_COST_REJECT_MAX);
assert!(mixer_block_admissible(&good_rot, &dense));
assert!(!mixer_block_admissible(&[7u32; 8], &dense), "one rotation amount");
assert!(mixer_block_admissible(&[1u32, 2, 3, 3, 3, 3, 3, 3], &dense), "three distinct amounts pass the agreed form");
// one two-adder multiplier among dense ones is rejected on its own (k >= 1)
let mut one_cheap = dense;
one_cheap[5] = 0x0000_0401;
assert!(mixer_cost(&one_cheap) > MIXER_COST_REJECT_MAX);
assert!(!mixer_block_admissible(&good_rot, &one_cheap), "a two-adder multiplier");
// the known-failed day
assert!(!mixer_block_admissible(&[1u32, 2, 3, 3, 3, 3, 3, 3], &dense), "three distinct amounts");
let v5 = Shape { mixer_mult: 8, cache_log2_words: 26, derive_len: 0, state: true };
let v4 = Shape { mixer_mult: 8, cache_log2_words: 26, derive_len: 0, state: false };
let worst = crate::seed::seed_words_from_bytes(&crate::bind::day_bytes(29_337));
let b = MixParams::with_shape(worst, v4);
println!("day 29,337 (2050-04-28) under class v4: cost {}, w32 min {}, distinct rot {}", mixer_cost(&b.mul), b.mul.iter().map(|&m| naf32_weight(m)).min().unwrap(), { let mut r = b.rot.to_vec(); r.sort_unstable(); r.dedup(); r.len() });
assert!(!mixer_block_admissible(&b.rot, &b.mul), "the known-failed day: the sub-version 3 block of day 29,337 is one the rule rejects");
let a = MixParams::with_shape(worst, v5);
assert!(a.redraws >= 1 && mixer_block_admissible(&a.rot, &a.mul), "class v5 redraws day 29,337");
assert_ne!((a.rot, a.mul), (b.rot, b.mul));
let mut redrawn = 0;
let mut scanned = 0;
for d in 0..60_000u64 {
@ -908,7 +859,7 @@ mod tests {
break;
}
}
assert!(redrawn >= 1, "no redraw in {scanned} days (the census says about 5.69e-4 per day)");
assert!(redrawn >= 1, "no redraw in {scanned} days (the census says about 6.1e-4 per day)");
}
#[test]

279
igneum-pow/src/verify.rs
View file

@ -9,38 +9,18 @@ use crate::seed::day_key;
/// window of the load site, `k = min(win, D - 26)` (0 when `D <= 26`), `idx = ((y & (MASK >> k)) | ((off &
/// (2^k - 1)) << (D - k))) & MASK`. For every other class `idx = x & MASK`, the lottery hash's address. `mask` is
/// `2^D - 1`. The acceptance mirror calls this at the rule's constant `D = 28`.
///
/// Class v6 lane 1, the index fold (`docs/design/class-v6-rotating-family.md` section 2; `EraParams::fold`): the
/// product's low bits are folded before the rotation, `y = x * M; y ^= y >> 16; y = rotl(y, R)` ([`stride`]), so no
/// era's R lands a biased product bit (bit 0 of `x * M` for odd M is bit 0 of x; bit 1 is set at 3/8) on an address
/// bit. Every era of every other class keeps the plain product.
#[inline(always)]
pub fn load_index(era: Option<&EraParams>, ins: &Instr, x: u32, mask: u32, log2: u32) -> u32 {
match era {
None => x & mask,
Some(e) => {
let (wm, off) = window(ins, mask, log2);
let y = stride(e, x);
let y = x.wrapping_mul(e.stride_mul).rotate_left(e.stride_rot);
((y & wm) | off) & mask
}
}
}
/// The era's stride of a source value: `rotl(x * M, R)`, with the index fold of class v6 lane 1 between the product
/// and the rotation when the era carries it (`y ^= y >> 16`). The one place the form lives on the CPU side; the three
/// emitters write the same text ([`crate::emit`]).
#[inline(always)]
pub fn stride(e: &EraParams, x: u32) -> u32 {
let mut y = x.wrapping_mul(e.stride_mul);
if e.fold {
y ^= y >> 16;
}
y.rotate_left(e.stride_rot)
}
/// The fold's shift (`y ^= y >> INDEX_FOLD_SHIFT`), named for the pack texts and the tests.
pub const INDEX_FOLD_SHIFT: u32 = 16;
/// The window of a load site at a dataset of `2^log2` words: `(window mask, offset)` such that
/// `idx = (y & window mask) | offset` lies in the site's aligned window of `2^(log2 - k)` words.
#[inline(always)]
@ -51,122 +31,6 @@ pub fn window(ins: &Instr, mask: u32, log2: u32) -> (u32, u32) {
(wm, off)
}
/// The window of a load site in the 32-bit SOURCE space (the multiply-shift mapping, research class ds55,
/// 8 October 2026): `k = min(win, log2 - 26)` as [`window`] with `log2 = floor(log2(N))`, and `(window mask,
/// offset)` such that `v = (y & window mask) | offset` lies in the site's aligned window of `2^(32 - k)` source
/// values; `idx = (v * N) >> 32` then lands in a contiguous run of about `N / 2^k` words, the site's window of the
/// dataset.
#[inline(always)]
pub fn window32(ins: &Instr, log2: u32) -> (u32, u32) {
let k = (ins.win as u32).min(log2.saturating_sub(26));
let wm = u32::MAX >> k;
let off = ((((ins.off as u32) & ((1u32 << k) - 1)) as u64) << (32 - k)) as u32;
(wm, off)
}
/// The dataset's geometry: `2^log2` words under the lottery hash's `src AND MASK` (`mulshift` false), or `words`
/// words, any multiple of 2^16 in `2^28 ..= 2^31`, under the multiply-shift range reduction of spec 01 section
/// 1.13.3 (`mulshift` true: `idx = (src * words) >> 32` in 64 bits; research class ds55, 8 October 2026, no
/// consensus object moves). A power of two given as a word count takes the mask path, so `--dataset-words 2^28`
/// is `--dataset-log2 28` byte for byte. `log2` is `floor(log2(words))` under the multiply-shift (the era
/// window's floor rule reads it); the item index `words / 16 - 1` stays 32-bit.
#[derive(Clone, Copy, Debug, PartialEq, Eq, Hash)]
pub struct DatasetGeom {
pub log2: u32,
pub words: u64,
pub mulshift: bool,
}
impl DatasetGeom {
/// The smallest word count `--dataset-words` takes: 2^28 (1 GiB).
pub const MIN_WORDS: u64 = 1 << 28;
/// The largest: 2^31 (8 GiB; the item index stays 32-bit far beyond, the kernel's word index is 32-bit).
pub const MAX_WORDS: u64 = 1 << 31;
/// The step: 2^16 words, so the era layout's interleave (positions below 16) stays a bijection of the dataset
/// and every wide or warp-coalesced load stays inside it.
pub const WORDS_STEP: u64 = 1 << 16;
/// A dataset of `2^log2` words (4 to 32), the lottery hash's mask path.
pub fn pow2(log2: u32) -> Self {
assert!((4..=32).contains(&log2), "dataset log2 must be in 4..=32");
Self { log2, words: 1u64 << log2, mulshift: false }
}
/// A dataset of `words` words: the mask path for a power of two, the multiply-shift otherwise. Refused
/// outside `MIN_WORDS ..= MAX_WORDS` or off the `WORDS_STEP` grid, with the reason.
pub fn words(words: u64) -> Result<Self, String> {
if !(Self::MIN_WORDS..=Self::MAX_WORDS).contains(&words) {
return Err(format!(
"--dataset-words {words}: the word count must be in 2^28 ..= 2^31 ({} ..= {})",
Self::MIN_WORDS,
Self::MAX_WORDS
));
}
if words % Self::WORDS_STEP != 0 {
return Err(format!("--dataset-words {words}: the word count must be a multiple of 2^16 words (65,536; the era interleave and the wide loads)"));
}
if words.is_power_of_two() {
return Ok(Self::pow2(words.trailing_zeros()));
}
Ok(Self { log2: words.ilog2(), words, mulshift: true })
}
/// The last word index (`words - 1`). Under the mask path it is the AND mask.
pub fn last_index(&self) -> u32 {
(self.words - 1) as u32
}
/// The AND mask of the mask path; under the multiply-shift the last index (recorded in packs as `IGNEUM_MASK`
/// for the host's self-test, never ANDed).
pub fn mask(&self) -> u32 {
self.last_index()
}
pub fn items(&self) -> u64 {
self.words >> 4
}
pub fn bytes(&self) -> u64 {
self.words << 2
}
/// The range reduction of a 32-bit source value to a word index.
#[inline(always)]
pub fn reduce(&self, x: u32) -> u32 {
if self.mulshift {
((x as u64 * self.words) >> 32) as u32
} else {
x & self.last_index()
}
}
/// One line for logs: `2^28 words` or `1476395008 words (5.50 GiB, not a power of two, multiply-shift)`.
pub fn describe(&self) -> String {
if self.mulshift {
format!("{} words ({:.2} GiB, not a power of two, multiply-shift)", self.words, self.bytes() as f64 / (1u64 << 30) as f64)
} else {
format!("2^{} words", self.log2)
}
}
}
/// [`load_index`] at a dataset geometry: the mask path unchanged; under the multiply-shift the plain load is
/// `(x * N) >> 32` and the era form windows the source first ([`window32`]) then reduces.
#[inline(always)]
pub fn load_index_geom(era: Option<&EraParams>, ins: &Instr, x: u32, geom: DatasetGeom) -> u32 {
if !geom.mulshift {
return load_index(era, ins, x, geom.mask(), geom.log2);
}
match era {
None => geom.reduce(x),
Some(e) => {
let (wm, off) = window32(ins, geom.log2);
let y = stride(e, x);
geom.reduce((y & wm) | off)
}
}
}
/// Read-width experiment (5 October 2026): a `load` of `W` words folds every word into `dst`:
/// `x = dst XOR w[0]; for j in 1..W: x = (rotl(x, FOLD_ROT) * FOLD_MUL) XOR w[j]; dst = x`. For `W = 1` this is the
/// lottery hash's `dst XOR dataset[...]`. The fold is state-dependent (the rotate-multiply sits between the words),
@ -320,12 +184,8 @@ pub enum Dataset {
/// A dataset of `2^log2` words plus the construction that fills it.
pub struct DatasetSource {
/// `floor(log2(words))`: the size under the mask path, the era window's floor under the multiply-shift.
pub log2_words: u32,
/// The last word index: the AND mask under the mask path (see [`DatasetGeom::mask`]).
pub mask: u32,
/// The geometry (size and range reduction); `log2_words` and `mask` are its `log2` and `mask()`.
pub geom: DatasetGeom,
/// The day key `K`; `d0, d1 = K[0], K[1]`.
pub key: [u32; 8],
/// The bytes `K` was derived from (`"day/<day>"` for a string day, `bind::day_bytes` on the chain), recorded
@ -359,25 +219,12 @@ impl DatasetSource {
/// `2^shape.cache_log2_words` words on the calling thread.
pub fn from_key_shape(key: [u32; 8], mode: DatasetMode, log2_words: u32, shape: Shape) -> Self {
assert!((4..=32).contains(&log2_words), "dataset log2 must be in 4..=32");
Self::from_key_geom(key, mode, DatasetGeom::pow2(log2_words), shape)
}
/// [`DatasetSource::from_key_shape`] at a dataset geometry (the multiply-shift sizes of `--dataset-words`).
pub fn from_key_geom(key: [u32; 8], mode: DatasetMode, geom: DatasetGeom, shape: Shape) -> Self {
let mask = if log2_words == 32 { u32::MAX } else { (1u32 << log2_words) - 1 };
let dataset = match mode {
DatasetMode::ClosedForm => Dataset::ClosedForm { d0: key[0], d1: key[1] },
DatasetMode::MemoryHard => Dataset::MemoryHard(MemhardCpu::with_shape(key, shape)),
};
Self { log2_words: geom.log2, mask: geom.mask(), geom, key, key_bytes: Vec::new(), dataset, hot: None }
}
/// This source at another geometry: the cache, key and construction unchanged (an item has the same value at
/// every size), only the size and the range reduction move.
pub fn with_geom(mut self, geom: DatasetGeom) -> Self {
self.log2_words = geom.log2;
self.mask = geom.mask();
self.geom = geom;
self
Self { log2_words, mask, key, key_bytes: Vec::new(), dataset, hot: None }
}
/// This source with the window's state leaves (class v5, `docs/design/class-v5-stored-state.md`): memory-hard mode
@ -399,7 +246,7 @@ impl DatasetSource {
Dataset::MemoryHard(m) => Dataset::MemoryHard(m.refreshed(leaves)),
Dataset::ClosedForm { .. } => panic!("state leaves on a closed-form dataset"),
};
Self { log2_words: self.log2_words, mask: self.mask, geom: self.geom, key: self.key, key_bytes: self.key_bytes.clone(), dataset, hot: None }
Self { log2_words: self.log2_words, mask: self.mask, key: self.key, key_bytes: self.key_bytes.clone(), dataset, hot: None }
}
/// A copy of this source sharing its cache (and leaves), for a caller that needs an owned source from a shared one.
@ -408,7 +255,7 @@ impl DatasetSource {
Dataset::MemoryHard(m) => Dataset::MemoryHard(crate::memhard::MemhardCpu { params: m.params.clone(), cache: m.cache.clone(), leaves: m.leaves.clone() }),
Dataset::ClosedForm { d0, d1 } => Dataset::ClosedForm { d0: *d0, d1: *d1 },
};
Self { log2_words: self.log2_words, mask: self.mask, geom: self.geom, key: self.key, key_bytes: self.key_bytes.clone(), dataset, hot: None }
Self { log2_words: self.log2_words, mask: self.mask, key: self.key, key_bytes: self.key_bytes.clone(), dataset, hot: None }
}
/// The window's state leaves, when the source carries them.
@ -452,14 +299,8 @@ impl DatasetSource {
}
/// `dataset[w & mask]` under a program's layout (era layout). The closed form has no items and ignores it.
/// Under the multiply-shift geometry `w` must already be a word index (below `words`).
pub fn word_at(&self, layout: Layout, w: u32) -> u32 {
let w = if self.geom.mulshift {
assert!((w as u64) < self.geom.words, "word {w} outside a dataset of {} words", self.geom.words);
w
} else {
w & self.mask
};
let w = w & self.mask;
match &self.dataset {
Dataset::ClosedForm { d0, d1 } => dataset_elem(w, *d0, *d1),
Dataset::MemoryHard(m) => m.word_at(layout, w),
@ -534,38 +375,11 @@ pub fn interpret_warp_scratch(
ds: &DatasetSource,
trace: bool,
) -> (WarpResult, Vec<ScratchEvent>) {
interpret_warp_core(program, seed, base_nonce, ds, trace, None)
}
/// The liveness probe of the reg64 window (`accept::check_window_liveness`): `flip` xors register `k` of every
/// lane with `x` at the start of iteration 0, after the init; `first_load_idx` receives the word index each
/// lane read at iteration 0's first load.
#[derive(Clone, Debug, Default)]
pub struct Probe {
pub flip: Option<(usize, u32)>,
pub first_load_idx: [u32; LANES],
pub seen_first_load: bool,
}
/// [`interpret_warp_init`] under a [`Probe`] (the reg64 liveness rule): the same interpreter, one flip, one read.
pub fn interpret_warp_probe(program: &Program, seed: &[u32; 8], base_nonce: u32, ds: &DatasetSource, probe: &mut Probe) -> WarpResult {
interpret_warp_core(program, seed, base_nonce, ds, false, Some(probe)).0
}
fn interpret_warp_core(
program: &Program,
seed: &[u32; 8],
base_nonce: u32,
ds: &DatasetSource,
trace: bool,
mut probe: Option<&mut Probe>,
) -> (WarpResult, Vec<ScratchEvent>) {
let geom = ds.geom;
let mask = ds.mask;
let log2 = ds.log2_words;
let era = program.class.era;
let layout = program.class.layout();
// 8 registers per lane, or 64 under the reg64 flag (r8..r63 derived from r0..r7 as the kernels do it)
let nregs = program.registers();
let mut r = vec![[0u32; LANES]; nregs];
let mut r = [[0u32; LANES]; 8];
for lane in 0..LANES {
let nonce = base_nonce.wrapping_add(lane as u32);
for i in 0..8 {
@ -574,13 +388,7 @@ fn interpret_warp_core(
x = splitmix32(x);
r[i][lane] = x ^ seed[(i + 1) & 7];
}
for k in 8..nregs {
r[k][lane] = r[k & 7][lane].wrapping_mul(0x9e3779b9u32).wrapping_add(k as u32);
}
}
// the iteration's statements: the drawn program, or its interleaved two-window form under reg64
let scheduled = program.scheduled();
let address_mix = program.address_mix();
let mut items_derived = 0usize;
let mut idx = [0u32; LANES];
let mut val = [0u32; LANES];
@ -595,28 +403,10 @@ fn interpret_warp_core(
let h = ds.hot.as_ref().expect("a hot-table program needs the epoch's hot table on the dataset source");
assert_eq!(h.n_words(), program.hot_words(), "the hot table's size is the class's");
}
for it in 0..ITERATIONS {
if it == 0 {
// the liveness probe's flip: one register of every lane, complemented at the start of iteration 0
if let Some(pr) = probe.as_deref_mut() {
if let Some((k, x)) = pr.flip {
for lane in 0..LANES {
r[k][lane] ^= x;
}
}
}
}
for _ in 0..ITERATIONS {
let sel = r[0];
for ins in &scheduled {
step(ins, &mut r, &sel, geom, era.as_ref(), layout, ds, &mut idx, &mut val, &mut items_derived, address_mix);
if it == 0 && ins.op == Op::Load {
if let Some(pr) = probe.as_deref_mut() {
if !pr.seen_first_load {
pr.first_load_idx = idx;
pr.seen_first_load = true;
}
}
}
for ins in &program.instrs {
step(ins, &mut r, &sel, mask, log2, era.as_ref(), layout, ds, &mut idx, &mut val, &mut items_derived);
if ins.op == Op::Scratch {
let m = scratch.as_mut().expect("a scratch op needs a scratch class");
let (d, a) = (ins.dst as usize, ins.src as usize);
@ -630,17 +420,7 @@ fn interpret_warp_core(
// iteration's `sel`; it is empty on every class without a shadow, so version 2 and class v3 run nothing here.
for _ in 0..program.shadow_reps() {
for ins in &program.shadow {
step(ins, &mut r, &sel, geom, era.as_ref(), layout, ds, &mut idx, &mut val, &mut items_derived, address_mix);
}
}
}
if nregs > 8 {
// reg64: both windows fold into the eight output registers by xor before the hash fold
for k in 0..8 {
for j in (k + 8..nregs).step_by(8) {
for lane in 0..LANES {
r[k][lane] ^= r[j][lane];
}
step(ins, &mut r, &sel, mask, log2, era.as_ref(), layout, ds, &mut idx, &mut val, &mut items_derived);
}
}
}
@ -658,38 +438,19 @@ fn interpret_warp_core(
#[allow(clippy::too_many_arguments)]
fn step(
ins: &Instr,
r: &mut [[u32; LANES]],
r: &mut [[u32; LANES]; 8],
sel: &[u32; LANES],
geom: DatasetGeom,
mask: u32,
log2: u32,
era: Option<&EraParams>,
layout: Layout,
ds: &DatasetSource,
idx: &mut [u32; LANES],
val: &mut [u32; LANES],
items_derived: &mut usize,
address_mix: bool,
) {
let d = ins.dst as usize;
let a = ins.src as usize;
// reg64 full chain: a load's address source is src ^ m, m the rotate-xor chain over the 63 other registers in
// index order (m = first; m = rotl(m, 1) ^ next). The source itself stays out of the chain: with it inside, a
// register whose term lands at rotation 0 mod 32 (r31, r63) cancels its own direct term, a dead register the
// liveness rule found on the pinned draw (first load src r31, 8 October 2026, 15:5x UK).
let addr_src = |r: &[[u32; LANES]], lane: usize| -> u32 {
if !address_mix {
return r[a][lane];
}
let mut m = 0u32;
let mut started = false;
for k in 0..r.len() {
if k == a {
continue;
}
m = if started { m.rotate_left(1) ^ r[k][lane] } else { r[k][lane] };
started = true;
}
r[a][lane] ^ m
};
match ins.op {
Op::Add => {
let (imm, imm2, bit) = (ins.imm, ins.imm2, ins.bit as u32);
@ -758,7 +519,7 @@ fn step(
}
Op::Load if ins.width == 1 => {
for lane in 0..LANES {
idx[lane] = load_index_geom(era, ins, addr_src(r, lane), geom);
idx[lane] = load_index(era, ins, r[a][lane], mask, log2);
}
*items_derived += ds.fetch(idx, val, layout);
for lane in 0..LANES {
@ -770,7 +531,7 @@ fn step(
let width = ins.width as usize;
let align = !(ins.width as u32 - 1);
for lane in 0..LANES {
idx[lane] = load_index_geom(era, ins, addr_src(r, lane), geom) & align;
idx[lane] = load_index(era, ins, r[a][lane], mask, log2) & align;
}
let mut vals = [[0u32; 16]; LANES];
*items_derived += ds.fetch_wide(idx, width, &mut vals, layout);
@ -790,8 +551,8 @@ fn step(
}
}
Op::WLoad => {
// Lane 0's register, range-reduced, aligned down to 32 words; lane l reads word base + l.
let base = geom.reduce(r[a][0]) & !31;
// Lane 0's register, masked, aligned down to 32 words; lane l reads word base + l.
let base = (r[a][0] & mask) & !31;
for lane in 0..LANES {
idx[lane] = base + lane as u32;
}

226
igneum-pow/tests/packs.rs
View file

@ -997,229 +997,3 @@ fn v5_pack_is_the_v4_program_over_the_state_leaves() {
let mh5 = v5_read("v5-genesis", "memhard.h");
assert!(mh5.contains("s[i] ^= leaf[i]"), "the leaf XOR before the first mixer");
}
// ---------------------------------------------------------------------------------------------------------------
// The 64-register window (the hash lane's measurement, 8 October 2026, a research class behind `+reg64`).
/// The `source` line the pinned devnet pack was exported with (vectors.json and vectors.h carry it).
const PINNED_SOURCE: &str = "igneum-pow (Rust) CPU interpreter, generator v3, memory-hard dataset";
/// The pinned devnet pack's epoch again, with the reg64 flag stamped on the program after the draw (as `--reg64`
/// does): the same instructions, the same dataset, 64 registers per lane.
fn epoch_reg64_of_devnet() -> Epoch {
let pack = "mx8-devnet-epoch0";
let j = json(pack, "program.json");
let seed = j["seed"].as_str().unwrap();
let seed_bytes = unhex(&j["seed_bytes"]);
let day_bytes = unhex(&j["dataset"]["day_bytes"]);
let log2 = j["dataset"]["log2_words"].as_u64().unwrap() as u32;
let era = j.get("era_seed_bytes").map(unhex);
let mut program = generate_from_seed_bytes_program_class(seed, &seed_bytes, ProgramClass::V3, era.as_deref());
program.class = program.class.with_reg64();
let shape = Shape::for_class(&program.class);
let mut dataset = DatasetSource::from_key_shape(igneum_pow::seed::seed_words_from_bytes(&day_bytes), DatasetMode::MemoryHard, log2, shape);
dataset.key_bytes = day_bytes;
Epoch { program, dataset }
}
/// The plain path re-exports the pinned devnet pack unchanged: every file `export` writes, byte for byte (the
/// reg64 code sits behind the flag; a pack without it does not move).
#[test]
fn reg64_plain_path_reexports_the_pinned_devnet_pack_unchanged() {
let pack = "mx8-devnet-epoch0";
let e = epoch(pack);
assert!(!e.program.class.reg64);
assert_eq!(e.program.registers(), 8);
assert_eq!(e.program.scheduled(), e.program.instrs, "without the flag the schedule is the drawn program");
let out = export_pack(e, &day_label(pack), PINNED_SOURCE);
assert_eq!(out.files.len(), 12);
for (name, text) in &out.files {
assert_same_text(pack, name, text);
}
}
/// The reg64 variant: the CPU verifier and the emitted CUDA text read one schedule, so the kernel's statements name
/// the registers the verifier wrote, in the verifier's order; the vectors are the verifier's; the plain pack's
/// instructions are untouched.
#[test]
fn reg64_cpu_verifier_and_cuda_text_agree() {
let plain = epoch("mx8-devnet-epoch0");
let e = epoch_reg64_of_devnet();
let p = &e.program;
assert!(p.class.reg64);
assert_eq!(p.registers(), 64);
assert_eq!(p.class.name(), "mx8-erad810f22d+reg64", "the pinned pack's era class, the window outermost");
assert_eq!(LoadClass::parse("mx8+reg64"), Some(V3_CLASS.with_reg64()), "the class name parses");
assert_eq!(LoadClass::MX8.with_reg64().name(), "mx8+reg64", "and round-trips");
assert_eq!(p.instrs, plain.program.instrs, "the draw is the class's without the flag");
assert_ne!(p.program_id(), plain.program.program_id(), "the id carries the flag");
// the schedule: 128 statements, instruction i of the draw on window 0 (+8 * (i % 4)) then on window 1 (+32)
let sched = p.scheduled();
assert_eq!(sched.len(), 128);
let mut touched = [false; 64];
for (i, ins) in p.instrs.iter().enumerate() {
let off = (8 * (i % 4)) as u8;
let a = sched[2 * i];
let b = sched[2 * i + 1];
assert_eq!((a.op, a.dst, a.src, a.src2), (ins.op, ins.dst + off, ins.src + off, ins.src2 + off), "instruction {i} on window 0");
assert_eq!((b.op, b.dst, b.src, b.src2), (ins.op, a.dst + 32, a.src + 32, a.src2 + 32), "instruction {i} on window 1");
touched[a.dst as usize] = true;
touched[b.dst as usize] = true;
}
// the fixed extension 8 * (i % 4) gives each octet 16 of the 64 instructions, so a draw need not write every
// register: the pinned draw writes 28 of 32 per window (r11, r17, r24, r29 unwritten, read or folded only)
let written = touched.iter().filter(|&&t| t).count();
assert!(written >= 48 && written % 2 == 0, "the schedule writes {written} of 64 registers (both windows alike)");
// the CUDA texts: 64 registers declared, 56 derived, every statement of the schedule in order, the fold, 32 masked loads
let mp = e.dataset.memhard().map(|m| &m.params);
for (name, text) in [("kernel.cu", cuda_kernel(p, mp)), ("kernel_bound.cu", cuda_kernel_bound(p, mp))] {
assert!(text.contains("uint32_t r0, r1, r2, r3, r4, r5, r6, r7, r8, r9,") && text.contains(", r62, r63;"), "{name}: 64 registers");
for k in 8..64 {
assert!(text.contains(&format!(" r{k} = r{} * 0x9e3779b9u + {k}u;\n", k & 7)), "{name}: r{k} derived");
}
for k in 0..8 {
let terms: Vec<String> = (k + 8..64).step_by(8).map(|j| format!("r{j}")).collect();
assert!(text.contains(&format!(" r{k} = r{k} ^ {};\n", terms.join(" ^ "))), "{name}: the fold of r{k}");
}
assert_eq!(text.matches("ds[((rotl_imm(").count(), 2 * LOAD_SLOTS, "{name}: 32 loads (the era form)");
assert_eq!(text.matches(" & mask]").count(), 2 * LOAD_SLOTS, "{name}: 32 masked loads");
// every statement line names the schedule's registers: dst first, then src on every op that reads one
let lines: Vec<&str> = text.lines().filter(|l| l.starts_with(" r") && l.contains(" // ")).collect();
assert_eq!(lines.len(), 128, "{name}: 128 statements per iteration");
for (k, (line, ins)) in lines.iter().zip(sched.iter()).enumerate() {
let stmt = line.trim_start();
assert!(stmt.starts_with(&format!("r{} = r{} ", ins.dst, ins.dst)) || stmt.starts_with(&format!("r{} = ", ins.dst)), "{name} statement {k}: dst r{}: {stmt}", ins.dst);
assert!(line.ends_with(&format!("// {k} {}", ins.op.name())), "{name} statement {k}: numbered: {line}");
if ins.op != Op::Rotl && ins.op != Op::Load {
assert!(stmt.contains(&format!("r{}", ins.src)) || stmt.contains(&format!("r{},", ins.src)), "{name} statement {k}: src r{}: {stmt}", ins.src);
}
}
}
// Metal and OpenCL carry the same window (class-v6, 8 October 2026): 64 registers declared, 56 derived, the fold,
// 128 statements per iteration in the schedule's order, 32 loads of the era form
for (name, text) in [
("program.metal", metal_program(p, e.dataset.log2_words, LoadSource::Stored)),
("program_bound.metal", metal_program_bound(p, e.dataset.log2_words)),
("kernel.cl", opencl_kernel(p, mp)),
("kernel_bound.cl", opencl_kernel_bound(p, mp)),
] {
assert!(!text.starts_with("// reg64: no "), "{name}: the window text, not a stub");
assert!(text.contains("uint r0, r1, r2, r3, r4, r5, r6, r7, r8, r9,") && text.contains(", r62, r63;"), "{name}: 64 registers");
for k in 8..64 {
assert!(text.contains(&format!(" r{k} = r{} * 0x9e3779b9u + {k}u;\n", k & 7)), "{name}: r{k} derived");
}
for k in 0..8 {
let terms: Vec<String> = (k + 8..64).step_by(8).map(|j| format!("r{j}")).collect();
assert!(text.contains(&format!(" r{k} = r{k} ^ {};\n", terms.join(" ^ "))), "{name}: the fold of r{k}");
}
// a statement line opens with its register, or with the shuffle's block in OpenCL ("{ uint t_; ...")
let lines: Vec<&str> = text.lines().filter(|l| (l.starts_with(" r") || l.starts_with(" { uint t_;")) && l.contains(" // ")).collect();
assert!(lines.len() % 128 == 0 && !lines.is_empty(), "{name}: {} statement lines, a multiple of 128", lines.len());
for (k, (line, ins)) in lines.iter().take(128).zip(sched.iter()).enumerate() {
let stmt = line.trim_start();
assert!(stmt.starts_with(&format!("r{} = ", ins.dst)) || (stmt.starts_with("{ uint t_;") && stmt.contains(&format!("r{} = r{} ^ t_", ins.dst, ins.dst))), "{name} statement {k}: dst r{}: {stmt}", ins.dst);
}
assert_eq!(text.matches("rotl_imm(").count() >= 2 * LOAD_SLOTS, true, "{name}: the era loads");
}
// the pack: the vectors are the reg64 verifier's and differ from the plain pack's; program.h and program.json carry the flag
let out = export_pack(&e, &day_label("mx8-devnet-epoch0"), "test");
for (i, &base) in out.bases.iter().enumerate() {
assert_eq!(out.outs[i], e.hash_warp(base));
assert_ne!(out.outs[i], plain.hash_warp(base), "base {base}: the reg64 hashes differ from the plain pack's");
}
let file = |n: &str| out.files.iter().find(|(f, _)| f == n).map(|(_, t)| t.clone()).unwrap();
assert!(file("program.h").contains("#define IGNEUM_REG64 1\n#define IGNEUM_REGISTERS 64\n"));
let j: Value = serde_json::from_str(&file("program.json")).unwrap();
assert_eq!(j["load_class"].as_str().unwrap(), "mx8-erad810f22d+reg64");
assert_eq!(j["reg64"]["registers"].as_u64().unwrap(), 64);
assert_eq!(j["reg64"]["statements_per_iteration"].as_u64().unwrap(), 128);
assert_eq!(j["program_class"].as_str().unwrap(), "v3");
let v: Value = serde_json::from_str(&file("vectors.json")).unwrap();
assert_eq!(v["warps"].as_array().unwrap().len(), 3);
}
/// The full-chain variant (`+reg64c`): every load's address source is `src ^ m` with `m` the rotate-xor mix of all
/// 64 registers; the CUDA text carries the mix statement before each of the 32 loads, the verifier the same chain,
/// and the hashes differ from the arithmetic-only window's.
#[test]
fn reg64_full_chain_cpu_verifier_and_cuda_text_agree() {
let mut e = epoch_reg64_of_devnet();
let arithmetic_only: Vec<[u64; 32]> = [0u32, 4096].iter().map(|&b| e.hash_warp(b)).collect();
e.program.class = e.program.class.with_reg64_chain();
let p = &e.program;
assert!(p.address_mix());
assert_eq!(p.class.name(), "mx8-erad810f22d+reg64c");
assert_eq!(LoadClass::parse("mx8+reg64c"), Some(V3_CLASS.with_reg64().with_reg64_chain()));
assert_ne!(LoadClass::parse("mx8+reg64c"), LoadClass::parse("mx8+reg64"));
assert_ne!(p.program_id(), epoch_reg64_of_devnet().program.program_id(), "the id carries the chain");
assert_ne!(p.program_id(), 0x3deee2320e70e1bf, "the sound fold's id differs from the benched text's (the source inside the chain)");
assert_eq!(p.scheduled().len(), 128, "the schedule is the window's");
let mp = e.dataset.memhard().map(|m| &m.params);
let sched = p.scheduled();
let load_srcs: Vec<u8> = sched.iter().filter(|i| i.op == Op::Load).map(|i| i.src).collect();
for (name, text) in [("kernel.cu", cuda_kernel(p, mp)), ("kernel_bound.cu", cuda_kernel_bound(p, mp))] {
assert!(text.contains(" uint32_t m; // reg64 full chain"), "{name}: m declared");
assert_eq!(text.matches(" ^ m)").count(), 2 * LOAD_SLOTS, "{name}: 32 mixed address sources");
assert_eq!(text.matches(" & mask]").count(), 2 * LOAD_SLOTS, "{name}: 32 masked loads");
// the mix line sits right before its load: 63 terms in index order, the load's own source left out, and the
// load reads (rS ^ m) with the schedule's src
let lines: Vec<&str> = text.lines().collect();
let mut loads = 0;
for (i, l) in lines.iter().enumerate() {
let t = l.trim_start();
if t.starts_with("m = r") && t.contains("rotl_imm(m, 1u)") {
let src = load_srcs[loads];
let next = lines[i + 1].trim_start();
assert!(next.contains("ds[") && next.contains(&format!("(r{src} ^ m)")), "{name}: load {loads} follows its mix with src r{src}: {next}");
let mut want = String::new();
for k in 0..64u8 {
if k == src {
continue;
}
if want.is_empty() {
want.push_str(&format!("m = r{k};"));
} else {
want.push_str(&format!(" m = rotl_imm(m, 1u) ^ r{k};"));
}
}
assert_eq!(t, want, "{name}: the mix of load {loads} (src r{src})");
loads += 1;
}
}
assert_eq!(loads, 2 * LOAD_SLOTS, "{name}: the mix before each of the 32 loads");
}
for (i, &b) in [0u32, 4096].iter().enumerate() {
assert_ne!(e.hash_warp(b), arithmetic_only[i], "base {b}: the chain's hashes differ from the window's");
}
let out = export_pack(&e, &day_label("mx8-devnet-epoch0"), "test");
let file = |n: &str| out.files.iter().find(|(f, _)| f == n).map(|(_, t)| t.clone()).unwrap();
assert!(file("program.h").contains("#define IGNEUM_REG64_ADDRESS_MIX 1"));
let j: Value = serde_json::from_str(&file("program.json")).unwrap();
assert_eq!(j["load_class"].as_str().unwrap(), "mx8-erad810f22d+reg64c");
assert_eq!(j["reg64"]["variant"].as_str().unwrap(), "window, full chain");
assert_eq!(j["reg64"]["address_mix"].as_u64().unwrap(), 1);
assert!(j["reg64"]["liveness"].as_str().unwrap().contains("64 independently necessary values"));
assert_eq!(j["program_class"].as_str().unwrap(), "v3");
}
/// The window's liveness rule (`accept::check_window_liveness`): the arithmetic-only window is the known-failed
/// fixture (a load's address reads its own window's register, so complementing a register of another octet leaves
/// iteration 0's first load address where it was); the full-chain form passes; a program without the window is
/// `NotAWindow`.
#[test]
fn reg64_liveness_rule_refuses_the_subset_fold_and_passes_the_full_chain() {
use igneum_pow::accept::{check_window_liveness, Reject};
let plain = epoch("mx8-devnet-epoch0");
assert_eq!(check_window_liveness(&plain.program), Err(Reject::NotAWindow));
let window = epoch_reg64_of_devnet();
match check_window_liveness(&window.program) {
Err(Reject::DeadWindowRegister { address_changed, result_changed, .. }) => {
assert!(!address_changed, "the subset fold leaves the first load address where it was");
assert!(result_changed, "the end fold still reads every register");
}
other => panic!("the arithmetic-only window must be refused as a dead register: {other:?}"),
}
let mut chain = epoch_reg64_of_devnet();
chain.program.class = chain.program.class.with_reg64_chain();
assert_eq!(check_window_liveness(&chain.program), Ok(()), "the full chain keeps all 64 registers live across the chain");
}

253
infra/fast-time/proving-cache-history.mjs Normal file
View file

@ -0,0 +1,253 @@
#!/usr/bin/env node
// The two-node cache-history case (review B F01, Phase 1 item (c), 8 October 2026): on a node with the verdict cache
// of verdict-cache-fix-node 3f672661, a proof's bytes refused once for CONTEXT (another statement) must still pay
// when the same bytes arrive later as their own honest record, answered from the cached facts with no second SP1
// verify; bytes refused once for being INVALID must be refused again at a later carrier, from the cache, with no
// second verify. Three local nodes at fast time as in proving-enforcement.mjs: H1 and H2 honest (the verifier in
// consensus), A the carrier (skip rule, trust verify) whose templates carry what its pool holds.
//
// The real proof is one of this chain's own blocks: after the chain has run, block n is exported from H1
// (igneum_exportSegments), cut by igneum-prove-export and proven compressed by igneum-prove-host on the box's CPU
// (about 70 to 110 s), so its statement is H1's native statement for (n, shard 0, PAYOUT). The node binaries and the
// proving binaries must carry the SAME pair (the node's elf/ overlay and the host's embedded pin: pin C tonight), and
// the node's object names that pair (proving_shard_program_id / proving_aggregator_id from the manifest).
//
// node infra/fast-time/proving-cache-history.mjs --prove-bin <dir with igneum-prove-host and igneum-prove-export>
// [--before 90] [--after 120] [--slot 0] [--out <json>]
//
// Verdict: (1) block n shard 0 paid exactly once, to PAYOUT; (2) H1's verifies.run grew by exactly 2 over the run (the
// real proof once, the invalid bytes once) and verifies.cacheAnswers by at least 2; (3) the context refusal of phase 1
// and the invalid refusal of phase 3 both read in H1's log; (4) the second use of the invalid bytes produced a refusal
// with no new verify.
import { spawn, spawnSync } from 'node:child_process';
import { mkdirSync, rmSync, writeFileSync, readFileSync, openSync, existsSync, appendFileSync } from 'node:fs';
import { createHash, randomBytes } from 'node:crypto';
const ROOT = new URL('../../', import.meta.url).pathname;
const FILE = `${ROOT}infra/fast-time/override-60x.json`;
const MANIFEST = `${ROOT}proving/igneum-prove/elf/manifest.json`;
const IGNEUMD = process.env.IGNEUMD || `${ROOT}vendor/igneum-node/target/release/igneumd`;
const CPU_MINER = process.env.IGNEUM_MINER || `${ROOT}vendor/igneum-node/target/release/igneum-miner`;
const args = process.argv.slice(2);
const flag = (name, dflt) => { const i = args.indexOf(`--${name}`); return i >= 0 ? Number(args[i + 1]) : dflt; };
const sflag = (name, dflt) => { const i = args.indexOf(`--${name}`); return i >= 0 ? args[i + 1] : dflt; };
const BEFORE = flag('before', 90), AFTER = flag('after', 120), SLOT = flag('slot', 0);
const PROVE_BIN = sflag('prove-bin', `${ROOT}proving/igneum-prove/target/release`);
const GENESIS_BITS = flag('genesis-bits', 0x1f010000);
const OUT = sflag('out') || `${ROOT}docs/plans/proving-enforcement/cache-history.json`;
const BASE = 30890 + SLOT * 40, SUFFIX = 985 + SLOT;
const CHAIN = `igneum-devnet-${SUFFIX}`;
const TMP = `/tmp/igneum-fast-time-ch${SLOT}`;
const PIDS = `${TMP}/pids`;
const NEVER = '18446744073709551615';
for (const b of [IGNEUMD, CPU_MINER, `${PROVE_BIN}/igneum-prove-host`, `${PROVE_BIN}/igneum-prove-export`]) if (!existsSync(b)) { console.error(`missing ${b}`); process.exit(2); }
const t0 = Date.now();
const log = (s) => { const t = new Date().toISOString().slice(11, 23); console.log(`${t} ch${SLOT} ${s}`); };
const sleep = (ms) => new Promise(r => setTimeout(r, ms));
const hex = (b) => '0x' + Buffer.from(b).toString('hex');
const sha256 = (b) => createHash('sha256').update(b).digest('hex');
// leftovers of an earlier run on this slot: by the pid file only
rmSync(TMP, { recursive: true, force: true }); mkdirSync(TMP, { recursive: true });
const started = [];
const track = (proc) => { started.push(proc); try { appendFileSync(PIDS, `${proc.pid}\n`); } catch { } };
// the override: the 60x file with the pair named from the manifest, the verifier in consensus from genesis, no succession
let baseText = readFileSync(FILE, 'utf8');
for (const k of ['proving_payment_activation_daa']) baseText = baseText.replace(new RegExp(`\\s*"${k}":\\s*[^,}\\n]+,?`), '');
const manifest = JSON.parse(readFileSync(MANIFEST, 'utf8'));
function mergeOverrideText(text, fields) {
let out = text;
for (const k of Object.keys(fields)) out = out.replace(new RegExp(`\\s*"${k}":\\s*[^,}\\n]+,?`), '');
const extra = Object.entries(fields).map(([k, v]) => `"${k}": ${typeof v === 'string' && !/^\d+$/.test(v) && v !== 'true' && v !== 'false' ? JSON.stringify(v) : v}`).join(', ');
return out.replace(/,?\s*}\s*$/, `,\n ${extra}\n}\n`);
}
const override = `${TMP}/override.json`;
writeFileSync(override, mergeOverrideText(baseText, {
genesis_bits: GENESIS_BITS, skip_proof_of_work: false, proving_v0_activation_daa: '0', proving_consensus_verify_daa: '0',
proving_shard_program_id: manifest.shard.program_id, proving_aggregator_id: manifest.aggregator.program_id,
verifier_in_consensus: 'true', proving_key_succession_daa: NEVER, proving_key_succession_window_daa: '120',
proving_next_shard_program_id: '', proving_next_aggregator_id: '',
program_class_v3_activation_daa: NEVER, program_class_v4_activation_daa: NEVER,
}));
const DAY_MS = (() => { const m = /"pow_day_ms":\s*([0-9]+)/.exec(baseText); return m ? +m[1] : 1440000; })();
class Node {
constructor(name, i, peers = [], env = {}) {
this.name = name; this.i = i; this.grpcPort = BASE + i * 10; this.p2pPort = BASE + i * 10 + 1; this.jsonPort = BASE + i * 10 + 2; this.evmPort = BASE + i * 10 + 3;
this.peers = peers; this.env = env; this.dir = `${TMP}/${name}`; this.logFile = `${this.dir}/node.log`;
}
get grpc() { return `grpc://127.0.0.1:${this.grpcPort}`; }
async start() {
mkdirSync(this.dir, { recursive: true });
const out = openSync(this.logFile, 'a');
const a = ['--devnet', `--devnet-suffix=${SUFFIX}`, '--nodnsseed', '--disable-upnp', '--nologfiles', '--enable-unsynced-mining', '--utxoindex',
`--appdir=${this.dir}`, `--rpclisten=127.0.0.1:${this.grpcPort}`, `--rpclisten-json=127.0.0.1:${this.jsonPort}`, `--evm-rpclisten=127.0.0.1:${this.evmPort}`,
`--listen=127.0.0.1:${this.p2pPort}`, `--override-params-file=${override}`, '--loglevel=info', '--yes'];
if (this.peers.length) for (const p of this.peers) a.push(`--addpeer=127.0.0.1:${p}`); else a.push('--outpeers=0');
this.proc = spawn(IGNEUMD, a, { stdio: ['ignore', out, out], env: { ...process.env, ...this.env } });
track(this.proc);
for (let k = 0; k < 60; k++) {
await sleep(1000);
try { await this.exec('igneum_getExecStatus', []); log(`${this.name} up (pid ${this.proc.pid}) after ${k + 1} s`); return; } catch { }
if (this.proc.exitCode !== null) throw new Error(`${this.name} exited ${this.proc.exitCode}: ${readFileSync(this.logFile, 'utf8').split('\n').slice(-5).join(' | ')}`);
}
throw new Error(`${this.name} did not answer in 60 s`);
}
async exec(method, params) {
const r = await fetch(`http://127.0.0.1:${this.evmPort}`, { method: 'POST', headers: { 'content-type': 'application/json' }, body: JSON.stringify({ jsonrpc: '2.0', id: 1, method, params }) });
const j = await r.json();
if (j.error) throw new Error(`${method}: ${j.error.message || JSON.stringify(j.error)}`);
return j.result;
}
logText() { try { return readFileSync(this.logFile, 'utf8'); } catch { return ''; } }
stop() { try { this.proc.kill('SIGINT'); } catch { } }
}
function startMiner(node, label, secs) {
const out = openSync(`${TMP}/${label}.log`, 'a');
const p = spawn(CPU_MINER, ['mine', node.grpc, '1', String(secs), label, '--engine', 'igneum-pow', '--payout-label', label, '--status-secs', '30', '--no-vote'], { stdio: ['ignore', out, out], env: { ...process.env, IGNEUM_POW_DAY_MS: String(DAY_MS) } });
track(p);
return p;
}
function signRecord(label, chain, block, number, shard, payout, statement, proofHash) {
const r = spawnSync(CPU_MINER, ['sign-record', label, chain, block, String(number), String(shard), payout, statement, proofHash], { encoding: 'utf8' });
if (r.status !== 0) throw new Error(`sign-record failed: ${r.stderr || r.stdout}`);
return JSON.parse(r.stdout.trim().split('\n').pop());
}
const H1 = new Node('H1', 0);
const H2 = new Node('H2', 1, [H1.p2pPort]);
const A = new Node('A', 2, [H1.p2pPort, H2.p2pPort], { IGNEUM_TEST_SKIP_PROOF_RULE: '1', IGNEUM_PROOF_VERIFY: 'trust' });
const nodes = [H1, H2, A];
let miners = [];
async function stopAll() {
for (const m of miners) { try { m.kill('SIGTERM'); } catch { } }
for (const n of nodes) n.stop();
await sleep(2000);
for (const p of started) { try { p.kill('SIGKILL'); } catch { } }
}
process.on('SIGINT', async () => { await stopAll(); process.exit(130); });
const PAYOUT = '0x' + 'c3'.repeat(20);
async function tipNumber(node) { const s = await node.exec('igneum_getExecStatus', []); return Number(s.executedTip ?? 0); }
async function sameTip(a, h) {
try { const x = await a.exec('igneum_getExecStatus', []); const y = await h.exec('igneum_getExecStatus', []); return x.executedTipHash && x.executedTipHash === y.executedTipHash; } catch { return false; }
}
async function rejoin(label) {
for (let k = 0; k < 120; k++) { if (await sameTip(A, H1)) { log(`${label}: A is on H1's tip (${k * 5} s)`); return true; } await sleep(5000); }
log(`${label}: A did not rejoin H1's tip in 600 s`); return false;
}
async function verifies(node) { const s = await node.exec('igneum_getProvingStatus', []); return s.verifies || { run: 0, cacheAnswers: 0 }; }
const refusals = (node) => (node.logText().match(/a carried proof record's proof does not verify: [^\n]*/g) || []);
async function submitVia(node, name, record, proof) {
let out;
try { out = await node.exec('igneum_submitProofRecord', [{ record, proof }]); } catch (e) { out = { accepted: false, reason: String(e.message) }; }
log(`${name}: A's pool ${out.accepted ? 'accepted' : 'refused'} (${out.reason || ''})`);
return out;
}
async function observe(n) {
try {
const r = await H1.exec('igneum_getProofRecords', ['0x' + n.toString(16)]);
return { paid: (r.paid || []).filter(Boolean), carried: (r.carried || []).map(c => ({ key: c.keyHash, carrier: c.carrierNumber, rejected: c.rejected, paidWei: c.paidWei })) };
} catch (e) { return { error: e.message }; }
}
/// A real proof of block n of this chain: exported from H1, cut and proven by the host (CPU), returning the proof
/// bytes and the statement the host printed (which must equal H1's native statement for (n, 0, PAYOUT)).
function proveBlock(n) {
const dir = `${TMP}/prove-${n}`; mkdirSync(dir, { recursive: true });
const body = JSON.stringify({ jsonrpc: '2.0', id: 1, method: 'igneum_exportSegments', params: ['0x0', '0x' + n.toString(16)] });
const r = spawnSync('curl', ['-s', '-m', '600', '-X', 'POST', `http://127.0.0.1:${H1.evmPort}`, '-H', 'Content-Type: application/json', '--data-binary', body], { encoding: 'utf8', maxBuffer: 1 << 30 });
if (r.status !== 0) throw new Error(`export failed: ${r.stderr}`);
writeFileSync(`${dir}/export.json`, JSON.stringify(JSON.parse(r.stdout).result));
const e = spawnSync(`${PROVE_BIN}/igneum-prove-export`, [`${dir}/export.json`, String(n), `${dir}/block-${n}.json`, '--source', `fast-time cache-history ${CHAIN}`], { encoding: 'utf8' });
if (e.status !== 0) throw new Error(`igneum-prove-export failed: ${(e.stdout + e.stderr).slice(-400)}`);
log(`block ${n} cut: ${(e.stdout.trim().split('\n').pop() || '').slice(0, 160)}`);
const h = spawnSync(`${PROVE_BIN}/igneum-prove-host`, [`${dir}/block-${n}.json`, '--mode', 'compressed', '--shard', '0', '--prover', PAYOUT, '--out', `${dir}/results.json`], { encoding: 'utf8', env: { ...process.env, SP1_PROVER: 'cpu' } });
const line = (h.stdout.match(/RESULT compressed shard 0: [^\n]*/) || [''])[0];
log(`host: ${line.slice(0, 200)}${h.status !== 0 ? ` (exit ${h.status}: ${(h.stderr || '').slice(-300)})` : ''}`);
if (h.status !== 0 || !/VERIFIED/.test(line)) throw new Error(`the host did not prove block ${n}`);
const file = `${dir}/block-${n}-shard-0-compressed.bin`;
if (!existsSync(file)) throw new Error(`no proof file ${file}`);
const statement = (line.match(/statement (0x[0-9a-f]{64})/) || [])[1];
return { bytes: readFileSync(file), statement, line };
}
const result = { case: 'cache-history', node: IGNEUMD, proveBin: PROVE_BIN, pair: { shard: manifest.shard.program_id, aggregator: manifest.aggregator.program_id }, startedAt: new Date().toISOString(), phases: {} };
try {
await H1.start(); await H2.start(); await A.start();
miners = [startMiner(H1, 'h1', BEFORE + 3 * AFTER + 900), startMiner(H2, 'h2', BEFORE + 3 * AFTER + 900), startMiner(A, 'attacker', BEFORE + 3 * AFTER + 900)];
log(`phase 0: mining ${BEFORE} s`);
await sleep(BEFORE * 1000);
const v0 = await verifies(H1);
// the block to prove: outside the exclusive window, executed on every node
let tip = await tipNumber(H1);
const n = Math.max(1, tip - 15);
const plan = await H1.exec('igneum_getShardPlan', ['0x' + n.toString(16), PAYOUT]);
const native = plan.shards[0].statement;
log(`block ${n} ${plan.hash}: H1's native statement for shard 0 and ${PAYOUT}: ${native}`);
const proof = proveBlock(n);
result.proof = { block: n, statement: proof.statement, native, bytes: proof.bytes.length, sha256: '0x' + sha256(proof.bytes), hostLine: proof.line };
if (proof.statement !== native) throw new Error(`the host's statement ${proof.statement} is not H1's ${native} (host and node layouts or pairs differ)`);
// phase 1: the real bytes under a record for another block (the statement of block m): refused for context
await rejoin('before phase 1');
tip = await tipNumber(A);
const m = Math.max(1, tip - 15);
const planM = await A.exec('igneum_getShardPlan', ['0x' + m.toString(16), PAYOUT]);
const recWrong = signRecord('ch-wrong', CHAIN, planM.hash, m, 0, PAYOUT, planM.shards[0].statement, '0x' + sha256(proof.bytes)).record;
const r1 = await submitVia(A, 'phase 1 (real bytes, another block\'s statement)', recWrong, hex(proof.bytes));
await sleep(AFTER * 1000);
const v1 = await verifies(H1);
const o1 = await observe(m);
result.phases.context = { block: m, submit: r1, observed: o1, verifies: v1, refusals: refusals(H1) };
log(`phase 1: H1 verifies ${JSON.stringify(v1)}; block ${m} paid ${JSON.stringify(o1.paid)}; refusals ${refusals(H1).length}`);
// phase 2: the same bytes as their own honest record: answered from the cached facts, paid once
await rejoin('before phase 2');
const recHonest = signRecord('ch-honest', CHAIN, plan.hash, n, 0, PAYOUT, native, '0x' + sha256(proof.bytes)).record;
const r2 = await submitVia(A, 'phase 2 (the same bytes, the honest record)', recHonest, hex(proof.bytes));
await sleep(AFTER * 1000);
const v2 = await verifies(H1);
const o2 = await observe(n);
result.phases.honest = { block: n, submit: r2, observed: o2, verifies: v2 };
log(`phase 2: H1 verifies ${JSON.stringify(v2)}; block ${n} paid ${JSON.stringify(o2.paid)}; carried ${JSON.stringify(o2.carried)}`);
// phase 3: invalid bytes twice, under two carriers: verified once, refused from the cache the second time
await rejoin('before phase 3');
const bad = randomBytes(2048);
tip = await tipNumber(A);
const p = Math.max(1, tip - 15);
const planP = await A.exec('igneum_getShardPlan', ['0x' + p.toString(16), PAYOUT]);
const recBad1 = signRecord('ch-bad1', CHAIN, planP.hash, p, 0, PAYOUT, planP.shards[0].statement, '0x' + sha256(bad)).record;
const r3 = await submitVia(A, 'phase 3a (invalid bytes)', recBad1, hex(bad));
await sleep(AFTER * 1000);
const v3a = await verifies(H1);
const ref3a = refusals(H1).length;
await rejoin('before phase 3b');
tip = await tipNumber(A);
const q = Math.max(1, tip - 15);
const planQ = await A.exec('igneum_getShardPlan', ['0x' + q.toString(16), PAYOUT]);
const recBad2 = signRecord('ch-bad2', CHAIN, planQ.hash, q, 0, PAYOUT, planQ.shards[0].statement, '0x' + sha256(bad)).record;
const r4 = await submitVia(A, 'phase 3b (the same invalid bytes, a later carrier)', recBad2, hex(bad));
await sleep(AFTER * 1000);
const v3b = await verifies(H1);
const ref3b = refusals(H1).length;
result.phases.invalid = { blocks: [p, q], submits: [r3, r4], verifies: { after3a: v3a, after3b: v3b }, refusals: { after3a: ref3a, after3b: ref3b }, reasons: refusals(H1) };
log(`phase 3: H1 verifies after 3a ${JSON.stringify(v3a)}, after 3b ${JSON.stringify(v3b)}; refusals ${ref3a} then ${ref3b}`);
// the verdict
const paidOnce = (o2.paid || []).length === 1 && (o1.paid || []).length === 0;
const runs = Number(v3b.run) - Number(v0.run);
const cache = Number(v3b.cacheAnswers) - Number(v0.cacheAnswers);
const contextRefused = result.phases.context.refusals.length >= 1;
const invalidRefusedTwiceNoVerify = ref3b > ref3a && Number(v3b.run) === Number(v3a.run);
const pass = paidOnce && runs === 2 && cache >= 2 && contextRefused && invalidRefusedTwiceNoVerify;
result.verdict = { paidOnce, verifiesRun: runs, cacheAnswers: cache, contextRefused, invalidRefusedTwiceNoVerify, pass };
result.endedAt = new Date().toISOString();
writeFileSync(OUT, JSON.stringify(result, null, 2));
log(`RESULT ${pass ? 'PASS' : 'FAIL'}: paid once=${paidOnce}; SP1 verifies run=${runs} (expect 2); cache answers=${cache} (expect >= 2); context refusal read=${contextRefused}; invalid refused again without a verify=${invalidRefusedTwiceNoVerify}; ${OUT}`);
await stopAll();
process.exit(pass ? 0 : 1);
} catch (e) {
log(`ERROR: ${e.message}`);
result.error = e.message; result.endedAt = new Date().toISOString();
try { writeFileSync(OUT, JSON.stringify(result, null, 2)); } catch { }
await stopAll();
process.exit(1);
}