reg64: the 64-register window per lane (a research class, +reg64 and +reg64c), the liveness rule, the measurement packs
The hash lane's measurement of 8 October 2026: each lane holds 64 live 32-bit registers. r0..r7 seeded as every class, r[k] = r[k & 7] * 0x9E3779B9 + k for k in 8..63; the drawn program runs twice per iteration in an interleaved order (instruction i on window 0 with every register field + 8 * (i % 4), then on window 1 at +32); both windows fold into r0..r7 by xor before the hash fold. Program::scheduled is the one list the CPU verifier and the CUDA emitters read. Metal and OpenCL carry a refusal stub. Class suffix +reg64 (window, arithmetic-only) and +reg64c (window, full chain: every load's address source is src ^ m, m the rotate-xor chain over the 63 other registers in index order); igneum-pow --reg64 and --reg64-chain. The id carries reg64/, chain/ and the fold byte (REG64_CHAIN_FOLD = 1, the source out of the chain), so the benched text (the source inside the chain, id 0x3deee2320e70e1bf) never shares an id with the committed one. The liveness rule (accept::check_window_liveness, beside check_distinct_indices_v4, not wired into the acceptance): each register in turn is xored with two seed-derived probe words at the start of iteration 0 in every lane of unit 0; the program passes when both words move the iteration's first load address and the final hash. The arithmetic-only window is the known-failed fixture (a load's address reads its own window's register only); the full chain passes. Two pseudo-random words because the chain is linear: a complement or a single bit that enters it twice through a copy the program makes (the pinned draw's r15 ^= r8) cancels. Tests: the plain path re-exports the pinned devnet pack byte for byte; the reg64 verifier and CUDA text agree on one schedule (128 statements, 32 masked loads, the derivation and fold lines, every statement's registers); the full-chain text carries the 63-term mix before each load with the load's own source left out; the liveness rule refuses the subset fold and passes the full chain. Suite green on box 2 (123 passed, 0 failed, 7 ignored). Measured on RunPod (stock, kit worker d84b1b6cb09cebdc, 250 x 2^24, block-warps 1), the rows in the lane's report: 5090 base 141.7 MH/s 303 W regs 30 blocks 24; window 80.4 MH/s 309 W regs 96 blocks 20; full chain 71.0 MH/s 320 W regs 88 blocks 20. 4090 base 62.7 MH/s 209 W regs 29 blocks 24; window 31.6 MH/s 210 W regs 104 blocks 16; full chain 31.4 MH/s 217 W regs 87 blocks 20. No spill on either card in any form. Packs on build-1 /srv/artefacts/packs: hl-reg64.tgz, hl-reg64c-benched.tgz, hl-reg64c.tgz, hl-v6-win.tgz, each with a .sha256. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
This commit is contained in:
igneum-labs
parent
cf029e26f7
commit
306886db18
6 changed files with 589 additions and 19 deletions
|
|
@ -111,6 +111,13 @@ 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 },
|
||||
}
|
||||
|
|
@ -133,6 +140,13 @@ 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)
|
||||
}
|
||||
|
|
@ -213,6 +227,63 @@ pub fn check_distinct_indices_v4(p: &Program) -> Result<(), Reject> {
|
|||
distinct_ratio_pass(p, ACCEPT_UNITS_DISTINCT_V4, MIN_DISTINCT_RATIO_V4).map(|_| ())
|
||||
}
|
||||
|
||||
/// 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.
|
||||
|
|
@ -972,6 +1043,8 @@ mod tests {
|
|||
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",
|
||||
|
|
|
|||
|
|
@ -238,6 +238,13 @@ 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 {{ {}, {}, {} }}
|
||||
|
|
@ -872,6 +879,9 @@ pub fn metal_program_bound(p: &Program, dataset_log2: u32) -> String {
|
|||
}
|
||||
|
||||
fn metal_program_impl(p: &Program, dataset_log2: u32, source: LoadSource, bound: bool) -> String {
|
||||
if p.class.reg64 {
|
||||
return reg64_refusal("Metal");
|
||||
}
|
||||
let mask = mask_for(dataset_log2);
|
||||
let mut s = String::with_capacity(5000);
|
||||
s.push_str("#include <metal_stdlib>\n");
|
||||
|
|
@ -1027,14 +1037,80 @@ fn init_line(p: &Program, u: &str, i: usize) -> String {
|
|||
)
|
||||
}
|
||||
|
||||
/// The per-lane register declaration of the CUDA kernels: r0..r7, or r0..r63 under reg64.
|
||||
fn cuda_reg_decl(p: &Program) -> String {
|
||||
if !p.class.reg64 {
|
||||
return " uint32_t r0, r1, r2, r3, r4, r5, r6, r7;\n".to_string();
|
||||
}
|
||||
let names: Vec<String> = (0..p.registers()).map(|k| format!("r{k}")).collect();
|
||||
let m = if p.address_mix() { "\n uint32_t m; // reg64 full chain: the address mix of all 64 registers before every load" } else { "" };
|
||||
format!(" uint32_t {}; // 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.
|
||||
fn cuda_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). Empty otherwise.
|
||||
fn cuda_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. Empty otherwise.
|
||||
fn cuda_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 Metal and OpenCL texts have no reg64 form (the pods are NVIDIA; the hash lane's measurement of 8 October
|
||||
/// 2026 is CUDA only). A pack of the flag carries this stub where those kernels would be.
|
||||
fn reg64_refusal(dialect: &str) -> String {
|
||||
format!("// reg64: no {dialect} text. The 64-register window (class suffix +reg64, igneum-pow --reg64) is emitted for CUDA only; export the class without the flag for this backend.\n")
|
||||
}
|
||||
|
||||
/// The instruction lines of the CUDA hash kernel body (shared by `igneum_hash` and `igneum_hash_bound`).
|
||||
fn cuda_instr_lines(p: &Program, dataset_log2: u32) -> String {
|
||||
let mut s = String::with_capacity(6000);
|
||||
let era = p.class.era;
|
||||
for (k, ins) in p.instrs.iter().enumerate() {
|
||||
// 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() {
|
||||
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", cuda_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!(
|
||||
|
|
@ -1053,9 +1129,9 @@ fn cuda_instr_lines(p: &Program, dataset_log2: u32) -> 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, &a, dataset_log2), ins.width, WideSource::Stored, None)
|
||||
wide_load_stmt(CoreDialect::Cuda, &d, &load_index_expr(CoreDialect::Cuda, era.as_ref(), ins, &addr_src(&a), dataset_log2), ins.width, WideSource::Stored, None)
|
||||
}
|
||||
Op::Load => format!("{d} = {d} ^ ds[{}];", load_index_expr(CoreDialect::Cuda, era.as_ref(), ins, &a, dataset_log2)),
|
||||
Op::Load => format!("{d} = {d} ^ ds[{}];", load_index_expr(CoreDialect::Cuda, era.as_ref(), ins, &addr_src(&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),
|
||||
|
|
@ -1162,17 +1238,19 @@ pub fn cuda_kernel_at(p: &Program, memhard: Option<&MixParams>, dataset_log2: u3
|
|||
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(" uint32_t r0, r1, r2, r3, r4, r5, r6, r7;\n");
|
||||
s.push_str(&cuda_reg_decl(p));
|
||||
if p.has_wide() {
|
||||
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(&cuda_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, dataset_log2));
|
||||
s.push_str(&shadow_block(p, CoreDialect::Cuda));
|
||||
s.push_str(" }\n");
|
||||
s.push_str(&cuda_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");
|
||||
|
|
@ -1311,7 +1389,7 @@ pub fn cuda_kernel_bound_at(p: &Program, memhard: Option<&MixParams>, dataset_lo
|
|||
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(" uint32_t r0, r1, r2, r3, r4, r5, r6, r7;\n");
|
||||
s.push_str(&cuda_reg_decl(p));
|
||||
if p.has_wide() {
|
||||
s.push_str(" uint32_t lane = threadIdx.x & 31u;\n uint32_t wmask = mask & ~31u;\n");
|
||||
}
|
||||
|
|
@ -1322,10 +1400,12 @@ pub fn cuda_kernel_bound_at(p: &Program, memhard: Option<&MixParams>, dataset_lo
|
|||
(i + 1) & 7
|
||||
));
|
||||
}
|
||||
s.push_str(&cuda_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, dataset_log2));
|
||||
s.push_str(&shadow_block(p, CoreDialect::Cuda));
|
||||
s.push_str(" }\n");
|
||||
s.push_str(&cuda_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");
|
||||
|
|
@ -1414,6 +1494,9 @@ 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 {
|
||||
if p.class.reg64 {
|
||||
return reg64_refusal("OpenCL");
|
||||
}
|
||||
let mut s = opencl_kernel_at(p, memhard, dataset_log2);
|
||||
s.push('\n');
|
||||
s.push_str(
|
||||
|
|
@ -1483,6 +1566,9 @@ 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 {
|
||||
if p.class.reg64 {
|
||||
return reg64_refusal("OpenCL");
|
||||
}
|
||||
let layout = p.class.layout();
|
||||
let mut s = String::with_capacity(14000);
|
||||
s.push_str(&generated_by(&p.seed_string));
|
||||
|
|
@ -1950,6 +2036,9 @@ 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.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");
|
||||
|
|
|
|||
|
|
@ -246,6 +246,26 @@ pub struct LoadClass {
|
|||
/// 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,
|
||||
/// 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
|
||||
|
|
@ -423,13 +443,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 };
|
||||
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, reg64: false, reg64_chain: 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 };
|
||||
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, reg64: false, reg64_chain: 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
|
||||
|
|
@ -570,6 +590,17 @@ impl LoadClass {
|
|||
LoadClass { era: other.era, ..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 }
|
||||
}
|
||||
|
||||
/// The class with the state leaves of class v5 folded into every item ("mx8+sh256x27+state").
|
||||
pub fn with_state(self) -> LoadClass {
|
||||
LoadClass { state: true, ..self }
|
||||
|
|
@ -606,6 +637,14 @@ 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>+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());
|
||||
|
|
@ -734,6 +773,11 @@ 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 {
|
||||
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());
|
||||
|
|
@ -1024,7 +1068,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, ..*class };
|
||||
let base = LoadClass { era: None, reg64: false, reg64_chain: false, ..*class };
|
||||
if base == LoadClass::V2 {
|
||||
Some(ProgramClass::V2)
|
||||
} else if base == V3_CLASS {
|
||||
|
|
@ -1048,7 +1092,52 @@ 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
|
||||
}
|
||||
|
|
@ -1153,7 +1242,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.is_v2() || self.generator == GENERATOR_VERSION_V3 || v4_rung_0 || v5_rung_0 {
|
||||
if !self.class.reg64 && (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
|
||||
|
|
@ -1170,7 +1259,7 @@ impl Program {
|
|||
// 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.is_v2() || self.generator == GENERATOR_VERSION_V3 || v4_rung_0 || v5_rung_0 {
|
||||
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()
|
||||
|
|
@ -1309,6 +1398,17 @@ pub fn program_id_class_recipe(generator: u32, seed: &[u32; 8], attempt: u32, cl
|
|||
// class v5: the state leaves are part of the construction
|
||||
r.lit(b"state/");
|
||||
}
|
||||
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
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -56,6 +56,12 @@ 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).
|
||||
|
|
@ -109,7 +115,9 @@ 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)"
|
||||
\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 --reg64 the 64-register window per lane over the class (research, CUDA text only; the same as --class <class>+reg64)\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)"
|
||||
);
|
||||
std::process::exit(2)
|
||||
}
|
||||
|
|
@ -136,6 +144,8 @@ 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());
|
||||
|
|
@ -161,6 +171,11 @@ 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(),
|
||||
}
|
||||
}
|
||||
|
|
@ -223,6 +238,13 @@ 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);
|
||||
if a.reg64 {
|
||||
// the 64-register window over the drawn program (CUDA text only; Metal and OpenCL carry a refusal stub)
|
||||
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 {
|
||||
|
|
|
|||
|
|
@ -374,12 +374,40 @@ pub fn interpret_warp_scratch(
|
|||
base_nonce: u32,
|
||||
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 mask = ds.mask;
|
||||
let log2 = ds.log2_words;
|
||||
let era = program.class.era;
|
||||
let layout = program.class.layout();
|
||||
let mut r = [[0u32; LANES]; 8];
|
||||
// 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];
|
||||
for lane in 0..LANES {
|
||||
let nonce = base_nonce.wrapping_add(lane as u32);
|
||||
for i in 0..8 {
|
||||
|
|
@ -388,7 +416,13 @@ pub fn interpret_warp_scratch(
|
|||
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];
|
||||
|
|
@ -403,10 +437,28 @@ pub fn interpret_warp_scratch(
|
|||
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 _ in 0..ITERATIONS {
|
||||
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;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
let sel = r[0];
|
||||
for ins in &program.instrs {
|
||||
step(ins, &mut r, &sel, mask, log2, era.as_ref(), layout, ds, &mut idx, &mut val, &mut items_derived);
|
||||
for ins in &scheduled {
|
||||
step(ins, &mut r, &sel, mask, log2, 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;
|
||||
}
|
||||
}
|
||||
}
|
||||
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);
|
||||
|
|
@ -420,7 +472,17 @@ pub fn interpret_warp_scratch(
|
|||
// 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, mask, log2, era.as_ref(), layout, ds, &mut idx, &mut val, &mut items_derived);
|
||||
step(ins, &mut r, &sel, mask, log2, 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];
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
@ -438,7 +500,7 @@ pub fn interpret_warp_scratch(
|
|||
#[allow(clippy::too_many_arguments)]
|
||||
fn step(
|
||||
ins: &Instr,
|
||||
r: &mut [[u32; LANES]; 8],
|
||||
r: &mut [[u32; LANES]],
|
||||
sel: &[u32; LANES],
|
||||
mask: u32,
|
||||
log2: u32,
|
||||
|
|
@ -448,9 +510,29 @@ fn step(
|
|||
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);
|
||||
|
|
@ -519,7 +601,7 @@ fn step(
|
|||
}
|
||||
Op::Load if ins.width == 1 => {
|
||||
for lane in 0..LANES {
|
||||
idx[lane] = load_index(era, ins, r[a][lane], mask, log2);
|
||||
idx[lane] = load_index(era, ins, addr_src(r, lane), mask, log2);
|
||||
}
|
||||
*items_derived += ds.fetch(idx, val, layout);
|
||||
for lane in 0..LANES {
|
||||
|
|
@ -531,7 +613,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(era, ins, r[a][lane], mask, log2) & align;
|
||||
idx[lane] = load_index(era, ins, addr_src(r, lane), mask, log2) & align;
|
||||
}
|
||||
let mut vals = [[0u32; 16]; LANES];
|
||||
*items_derived += ds.fetch_wide(idx, width, &mut vals, layout);
|
||||
|
|
|
|||
|
|
@ -997,3 +997,207 @@ 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 refuse the flag with the one-line stub
|
||||
for text in [metal_program(p, e.dataset.log2_words, LoadSource::Stored), metal_program_bound(p, e.dataset.log2_words), opencl_kernel(p, mp), opencl_kernel_bound(p, mp)] {
|
||||
assert!(text.starts_with("// reg64: no ") && text.contains("emitted for CUDA only"), "{text}");
|
||||
}
|
||||
// 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");
|
||||
}
|
||||
|
|
|
|||
Loading…
Reference in a new issue