diff --git a/igneum-pow/.gitignore b/igneum-pow/.gitignore new file mode 100644 index 000000000..2f7896d1d --- /dev/null +++ b/igneum-pow/.gitignore @@ -0,0 +1 @@ +target/ diff --git a/igneum-pow/Cargo.lock b/igneum-pow/Cargo.lock new file mode 100644 index 000000000..4cf2d0537 --- /dev/null +++ b/igneum-pow/Cargo.lock @@ -0,0 +1,105 @@ +# This file is automatically @generated by Cargo. +# It is not intended for manual editing. +version = 4 + +[[package]] +name = "igneum-pow" +version = "0.1.0" +dependencies = [ + "serde_json", +] + +[[package]] +name = "itoa" +version = "1.0.18" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8f42a60cbdf9a97f5d2305f08a87dc4e09308d1276d28c869c684d7777685682" + +[[package]] +name = "memchr" +version = "2.8.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "cf8baf1c55e62ffcace7a9f06f4bd9cd3f0c4beb022d3b367256b91b87513d98" + +[[package]] +name = "proc-macro2" +version = "1.0.107" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "985e7ec9bb745e6ce6535b544d84d6cd6f7ad8bd711c398938ae983b91a766d9" +dependencies = [ + "unicode-ident", +] + +[[package]] +name = "quote" +version = "1.0.47" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1fbf4db142a473a8d80c26bbf18454ed458bf8d26c8219c331daecfdbd079001" +dependencies = [ + "proc-macro2", +] + +[[package]] +name = "serde" +version = "1.0.229" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4148590afebada386688f18773da617792bf2ef03ffc1e4cbd2b1d45b023e0ba" +dependencies = [ + "serde_core", +] + +[[package]] +name = "serde_core" +version = "1.0.229" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "67dca2c9c51e58a4791a4b1ed58308b39c64224d349a935ab5039aa360942a48" +dependencies = [ + "serde_derive", +] + +[[package]] +name = "serde_derive" +version = "1.0.229" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e7a5d71263a5a7d47b41f6b3f06ba276f10cc18b0931f1799f710578e2309348" +dependencies = [ + "proc-macro2", + "quote", + "syn", +] + +[[package]] +name = "serde_json" +version = "1.0.151" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c841b55ecdae098c80dcae9cf767f6f8a0c2cdb3416bbef72181df4d0fe73f14" +dependencies = [ + "itoa", + "memchr", + "serde", + "serde_core", + "zmij", +] + +[[package]] +name = "syn" +version = "3.0.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8593e8e72159ed2257d083c7a454a85cbf854f37a0966d8d483aff8c8a3ebcee" +dependencies = [ + "proc-macro2", + "quote", + "unicode-ident", +] + +[[package]] +name = "unicode-ident" +version = "1.0.26" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d245f478577f809a851594d02313b640fb437e0bb33866753cff937863096954" + +[[package]] +name = "zmij" +version = "1.0.23" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "29666d0abbfad1e3dc4dcf6144730dd3a3ab225bbbdac83319345b1b44ccfc1b" diff --git a/igneum-pow/Cargo.toml b/igneum-pow/Cargo.toml new file mode 100644 index 000000000..5405b4f96 --- /dev/null +++ b/igneum-pow/Cargo.toml @@ -0,0 +1,30 @@ +[package] +name = "igneum-pow" +version = "0.1.0" +edition = "2021" +description = "Igneum random-program GPU proof-of-work: seed, program generator, memory-hard dataset, CPU warp verifier and kernel emitters, bit-exact with proto-metal" +license = "MIT" +publish = false + +[lib] +name = "igneum_pow" +path = "src/lib.rs" + +[[bin]] +name = "igneum-pow" +path = "src/main.rs" + +[dependencies] + +[dev-dependencies] +serde_json = "1" + +# The cache fill is 2^22 ChaCha12 blocks and the vector tests derive thousands of items. +# Unoptimised builds would make `cargo test` take minutes, so the dev profile is optimised too. +[profile.dev] +opt-level = 3 + +[profile.release] +opt-level = 3 +lto = true +codegen-units = 1 diff --git a/igneum-pow/README.md b/igneum-pow/README.md new file mode 100644 index 000000000..e4d760ccd --- /dev/null +++ b/igneum-pow/README.md @@ -0,0 +1,97 @@ +# igneum-pow + +The Igneum lottery hash in Rust, bit-exact with the Swift prototype in `proto-metal/main.swift`. This is the +crate the rusty-kaspa fork will call (`docs/fork-map.md`, rows a1 to a3) so a node written in Rust can verify any +block and hand miners the kernel source for the epoch. No dependency outside the standard library; `serde_json` +is a dev-dependency for reading the packs in the tests. + +Date: 3 October 2026. Toolchain: rustc 1.99.0 via rustup (the Homebrew 1.69 on PATH is too old; use +`~/.cargo/bin/cargo`). + +## Modules + +| Module | What it is | Swift namesake | +|---|---|---| +| `seed` | 32-byte seed words from a string (FNV-1a 64, four salts, finalised); `seed_words_from_bytes` is the boundary where the chain will feed the VDF output; SplitMix64 | `seedWords`, `SplitMix64` | +| `generator` | the 64-instruction program for a seed (op, dst, src, src2, imm, imm2, rot, bit, mask); levers `load_weight` and `wide_frac` | `generateProgram`, `GeneratorConfig` | +| `memhard` | 256 MiB cache (2^16 chains of 64 ChaCha12 blocks), mixer parameters, 8-round item derivation with the 32 lanes interleaved, `MemhardCpu::fetch` | `cpuFillCache`, `MixParams`, `deriveItems`, `MemhardCPU` | +| `verify` | the 32-lane warp interpreter, `DatasetMode::{ClosedForm, MemoryHard}`, `Epoch`, `hash_warp`, `verify_block` | `cpuWarp`, `DatasetSource` | +| `emit` | Metal, CUDA and OpenCL source, program.h, memhard.h, vectors.h, program.json, vectors.json, `export_pack` | `generateMSL`, `memhardMSL`, `emitMemhardCore`, `generateCUDA`, `generateOpenCL`, `exportPack` | + +## The API the fork calls + +```rust +use igneum_pow::{Epoch, DatasetMode}; + +// Once per epoch and day: generates the program and fills the 256 MiB cache (about 0.18 s on one core). +let epoch = Epoch::memory_hard("igneum-genesis", "2026-10-03"); + +let h: u64 = epoch.hash(nonce); // one nonce (computes its aligned 32-nonce warp) +let w: [u64; 32] = epoch.hash_warp(base_nonce); // one warp +let ok: bool = epoch.verify_block(nonce, target_u64); + +// Miner programs for the epoch, byte-identical to the Swift exporter. +let pack = igneum_pow::emit::export_pack(&epoch, "2026-10-03", "igneum node"); +pack.write_to(std::path::Path::new("out"))?; // kernel.cu, kernel.cl, program.metal, memhard.h, ... +``` + +`Epoch` is `Send + Sync`; build one and share it. `DatasetMode::ClosedForm` reproduces the two old packs +(`igneum-genesis`, `igneum-hourly`) and is not memory-hard. The hash is 64 bits; the fork maps it into its +256-bit target space in `consensus/pow/src/lib.rs`. + +## CLI + +``` +cargo build --release +./target/release/igneum-pow bench --seed igneum-genesis [--warps 20] [--closed-form] [--day 2026-10-03] +./target/release/igneum-pow export --seed igneum-genesis --out [--closed-form] +./target/release/igneum-pow hash --seed igneum-genesis --nonce 4103 +``` + +## Tests + +`cargo test` (23 tests, 0.7 s after compile; the dev profile is optimised so the cache fill is quick): + +| Check | Pack | Result | +|---|---|---| +| program.json instruction by instruction, op mix, loads per hash | igneum-genesis, igneum-genesis-mh, igneum-hourly | 3 x 64 match | +| Mixer parameters (key, rot, mul, rc) | igneum-genesis-mh | match | +| Cache head, last line, FNV-1a 64 `48c4f5bf24166b2e` | igneum-genesis-mh | match | +| Dataset head (16), `[MASK]`, 64 sampled words | all three | match | +| 96 hash vectors (3 warps x 32 lanes) | igneum-genesis-mh | 96/96 | +| 96 hash vectors | igneum-genesis, igneum-hourly | 96/96 each | +| kernel.cu, program.metal, kernel.cl, program.h byte-identical | all three | identical | +| memhard.h, memhard.metal byte-identical | igneum-genesis-mh | identical | +| program.json byte-identical (after the fix below) | all three | identical | +| vectors.json, vectors.h byte-identical apart from the provenance string | all three | identical | + +An independent `diff -r` of `igneum-pow export` output against the checked-in packs shows the same two lines +only: the provenance string and the `"item"` line. + +One deliberate difference: `proto-cuda/packs/igneum-genesis-mh/program.json` as written by the Swift is not +valid JSON (main.swift line 1291 uses `jhex` inside the `"item"` string, so the cache line mask is quoted inside a +quoted string). The Rust emitter writes `0x003fffff` bare; the test normalises that one line before comparing. +A node must hand miners valid JSON, so the Rust side does not reproduce the defect. + +## Measured, 3 October 2026, Apple M5 Max, one core, release build + +| Step | Rust | Swift (MEMHARD.md) | +|---|---|---| +| Cache fill, 256 MiB, 65,536 chains x 64 ChaCha12 blocks | 175 to 181 ms (5 quiet runs; 200 ms once with another build running) | 184.5 to 190.6 ms (C++ host reference 161.5) | +| CPU verify per warp, igneum-genesis, 104 loads, 3,328 items, avg of 20 | 0.441 ms | 0.649 ms | +| igneum-genesis/epoch1, 104 loads | 0.411 ms | 0.631 ms | +| igneum-genesis/epoch2, 112 loads | 0.488 ms | 0.701 ms | +| igneum-second-seed, 104 loads | 0.482 ms | 0.801 ms | +| igneum-second-seed/epoch1, 144 loads, 4,608 items | 0.579 ms | 1.205 ms | +| Cold single warps across the five seeds | 0.41 to 0.87 ms | 1.16 to 2.11 ms | +| Closed form, igneum-genesis | 0.002 ms | 0.017 ms | + +The Rust verifier is 1.4x to 2.1x faster than the Swift one per warp; the registers are kept register-major +(`r[reg][lane]`) so the lane loops vectorise, and the item derivation interleaves the 32 lanes round by round as +the Swift does. The 10 ms gate holds with a margin of about 17x on the steady figure and 11x on the worst cold warp. + +## Not done here + +- No GPU. The vectors tie this crate to the Metal and CUDA results through the packs; nothing here runs a kernel. +- The epoch seed is still a string. `seed::seed_words_from_bytes` is where the VDF output will enter. +- The 256-bit target mapping and the `kaspa_pow::State` shape belong to the fork, not to this crate. diff --git a/igneum-pow/rustfmt.toml b/igneum-pow/rustfmt.toml new file mode 100644 index 000000000..c775577ee --- /dev/null +++ b/igneum-pow/rustfmt.toml @@ -0,0 +1,2 @@ +max_width = 120 +use_small_heuristics = "Max" diff --git a/igneum-pow/src/emit.rs b/igneum-pow/src/emit.rs new file mode 100644 index 000000000..c8350e170 --- /dev/null +++ b/igneum-pow/src/emit.rs @@ -0,0 +1,1044 @@ +//! Kernel source emitters. Each function here writes the same bytes as its namesake in +//! `proto-metal/main.swift` (`generateMSL`, `memhardMSL`, `emitMemhardCore`, `generateCUDA`, `generateOpenCL`, +//! `generateProgramHeader`, `generateMemhardHeader`, `generateVectorsHeader`, `generateProgramJSON`, +//! `generateVectorsJSON`). The pack tests diff them against `proto-cuda/packs/*`. +//! +//! One deliberate difference: `program_json` writes the cache line mask inside the `"item"` string as a bare +//! `0x003fffff`. The Swift writes it quoted (`jhex`), which makes `igneum-genesis-mh/program.json` invalid +//! JSON. A node must hand miners valid JSON, so the Rust side does not reproduce that defect. + +use crate::generator::{Op, Program, INSTR_COUNT, ITERATIONS}; +use crate::memhard::{ + MixParams, CACHE_LINES_PER_SEGMENT, CACHE_LINE_MASK, CACHE_LOG2_WORDS, CACHE_SEGMENTS, CACHE_SEGMENT_LOG2_LINES, + CACHE_TAG, CACHE_WORDS, CHACHA_ROUNDS, CHACHA_SIGMA, ITEM_ROUNDS, +}; +use crate::seed::SplitMix64; +use crate::verify::{DatasetMode, DatasetSource, Epoch}; + +pub fn hex(v: u32) -> String { + format!("0x{v:08x}u") +} +pub fn hex64(v: u64) -> String { + format!("0x{v:016x}ull") +} +fn jhex(v: u32) -> String { + format!("\"0x{v:08x}\"") +} +fn jhex64(v: u64) -> String { + format!("\"0x{v:016x}\"") +} +/// JSON string with the three escapes the Swift applies (quote, backslash, newline). +fn jstr(s: &str) -> String { + let mut o = String::with_capacity(s.len() + 2); + o.push('"'); + for c in s.chars() { + match c { + '"' => o.push_str("\\\""), + '\\' => o.push_str("\\\\"), + '\n' => o.push_str("\\n"), + _ => o.push(c), + } + } + o.push('"'); + o +} +fn join_hex(v: &[u32]) -> String { + v.iter().map(|&x| hex(x)).collect::>().join(", ") +} +fn join_jhex(v: &[u32]) -> String { + v.iter().map(|&x| jhex(x)).collect::>().join(", ") +} + +/// The three dialects of the memory-hard core. +#[derive(Clone, Copy, Debug, PartialEq, Eq)] +pub enum CoreDialect { + Metal, + Cuda, + OpenCl, +} + +/// How the hash kernel gets dataset words. `Stored` is the honest kernel; the inline variants are the +/// shortcut measurements of MEMHARD.md section 2.2. +pub enum LoadSource<'a> { + Stored, + InlineClosed(u32, u32), + InlineMemhard(&'a MixParams), +} + +fn log2_segments() -> usize { + CACHE_SEGMENTS.trailing_zeros() as usize +} + +/// The memory-hard core as source text (`emitMemhardCore`). Every parameter is a literal. +pub fn emit_memhard_core(mp: &MixParams, dialect: CoreDialect) -> String { + let (u, fn_, cptr, wptr, lptr, lcptr) = match dialect { + CoreDialect::Metal => { + ("uint", "inline", "device const uint*", "device uint*", "thread uint*", "const thread uint*") + } + CoreDialect::Cuda => ("uint32_t", "IGNEUM_HD", "const uint32_t*", "uint32_t*", "uint32_t*", "const uint32_t*"), + CoreDialect::OpenCl => { + ("uint", "static inline", "__global const uint*", "__global uint*", "uint*", "const uint*") + } + }; + let k = &mp.key; + let r = &mp.rot; + let m = &mp.mul; + let c = &mp.rc; + let mut s = String::with_capacity(6000); + s.push_str(&format!( + "// Memory-hard dataset core (MEMHARD.md). Cache: 2^{} words in 2^{} segments of {} chained ChaCha{} lines.\n", + CACHE_LOG2_WORDS, + log2_segments(), + CACHE_LINES_PER_SEGMENT, + CHACHA_ROUNDS + )); + s.push_str("// Item: 8 rounds of seed-parameterised mixer + one 64-byte cache read, then a final mixer. All parameters are literals.\n"); + s.push_str(&format!("#define MH_CACHE_LINE_MASK {}\n", hex(CACHE_LINE_MASK))); + s.push_str(&format!("#define MH_SEGMENT_LINES {}u\n", CACHE_LINES_PER_SEGMENT)); + s.push_str("#define MH_QR(a, b, c, d, r1, r2, r3, r4) { a += b; d ^= a; d = mh_rotl(d, r1); c += d; b ^= c; b = mh_rotl(b, r2); a += b; d ^= a; d = mh_rotl(d, r3); c += d; b ^= c; b = mh_rotl(b, r4); }\n"); + s.push_str(&format!( + "{fn_} {u} mh_rotl({u} x, {u} n) {{ return (x << n) | (x >> (32u - n)); }} // n in 1..31 at every call site\n" + )); + s.push('\n'); + s.push_str(&format!("// y = ChaCha{CHACHA_ROUNDS} core(x) + x\n")); + s.push_str(&format!("{fn_} void mh_chacha_block({lcptr} x, {lptr} y) {{\n")); + s.push_str(&format!(" for ({u} i = 0u; i < 16u; ++i) y[i] = x[i];\n")); + s.push_str(&format!(" for ({u} r = 0u; r < {}u; ++r) {{\n", CHACHA_ROUNDS / 2)); + s.push_str( + " MH_QR(y[0], y[4], y[8], y[12], 16u, 12u, 8u, 7u) MH_QR(y[1], y[5], y[9], y[13], 16u, 12u, 8u, 7u)\n", + ); + s.push_str( + " MH_QR(y[2], y[6], y[10], y[14], 16u, 12u, 8u, 7u) MH_QR(y[3], y[7], y[11], y[15], 16u, 12u, 8u, 7u)\n", + ); + s.push_str( + " MH_QR(y[0], y[5], y[10], y[15], 16u, 12u, 8u, 7u) MH_QR(y[1], y[6], y[11], y[12], 16u, 12u, 8u, 7u)\n", + ); + s.push_str( + " MH_QR(y[2], y[7], y[8], y[13], 16u, 12u, 8u, 7u) MH_QR(y[3], y[4], y[9], y[14], 16u, 12u, 8u, 7u)\n", + ); + s.push_str(" }\n"); + s.push_str(&format!(" for ({u} i = 0u; i < 16u; ++i) y[i] += x[i];\n")); + s.push_str("}\n"); + s.push('\n'); + s.push_str(&format!( + "// One cache segment: {} chained lines written at cache[seg * {}]. in_j = prev ^ (sigma || K || seg || j || tag), prev_0 = 0.\n", + CACHE_LINES_PER_SEGMENT, + CACHE_LINES_PER_SEGMENT * 16 + )); + s.push_str(&format!("{fn_} void mh_cache_segment({wptr} cache, {u} seg) {{\n")); + s.push_str(&format!(" {u} prev[16]; {u} x[16]; {u} y[16];\n")); + s.push_str(&format!(" for ({u} i = 0u; i < 16u; ++i) prev[i] = 0u;\n")); + s.push_str(&format!(" for ({u} j = 0u; j < MH_SEGMENT_LINES; ++j) {{\n")); + s.push_str(&format!( + " x[0] = {} ^ prev[0]; x[1] = {} ^ prev[1]; x[2] = {} ^ prev[2]; x[3] = {} ^ prev[3];\n", + hex(CHACHA_SIGMA[0]), + hex(CHACHA_SIGMA[1]), + hex(CHACHA_SIGMA[2]), + hex(CHACHA_SIGMA[3]) + )); + for i in 0..8 { + s.push_str(&format!(" x[{}] = {} ^ prev[{}];\n", 4 + i, hex(k[i]), 4 + i)); + } + s.push_str(&format!( + " x[12] = seg ^ prev[12]; x[13] = j ^ prev[13]; x[14] = {} ^ prev[14]; x[15] = {} ^ prev[15];\n", + hex(CACHE_TAG[0]), + hex(CACHE_TAG[1]) + )); + s.push_str(" mh_chacha_block(x, y);\n"); + s.push_str(&format!(" {wptr} line = cache + ((seg * MH_SEGMENT_LINES + j) * 16u);\n")); + s.push_str(&format!(" for ({u} i = 0u; i < 16u; ++i) {{ line[i] = y[i]; prev[i] = y[i]; }}\n")); + s.push_str(" }\n"); + s.push_str("}\n"); + s.push('\n'); + s.push_str("// M_r: per word (s ^ (RC + rk)) * MUL, then a column round and a diagonal round with the seed-drawn rotations.\n"); + s.push_str(&format!("{fn_} void mh_mixer({lptr} s, {u} rk) {{\n")); + for i in 0..16 { + s.push_str(&format!(" s[{i}] = (s[{i}] ^ ({} + rk)) * {};\n", hex(c[i]), hex(m[i]))); + } + let col = (0..4).map(|i| format!("{}u", r[i])).collect::>().join(", "); + let dia = (4..8).map(|i| format!("{}u", r[i])).collect::>().join(", "); + s.push_str(&format!(" MH_QR(s[0], s[4], s[8], s[12], {col}) MH_QR(s[1], s[5], s[9], s[13], {col})\n")); + s.push_str(&format!(" MH_QR(s[2], s[6], s[10], s[14], {col}) MH_QR(s[3], s[7], s[11], s[15], {col})\n")); + s.push_str(&format!(" MH_QR(s[0], s[5], s[10], s[15], {dia}) MH_QR(s[1], s[6], s[11], s[12], {dia})\n")); + s.push_str(&format!(" MH_QR(s[2], s[7], s[8], s[13], {dia}) MH_QR(s[3], s[4], s[9], s[14], {dia})\n")); + s.push_str("}\n"); + s.push('\n'); + s.push_str(&format!( + "// Item t: 16 words. s = (K, t * MUL[i] + RC[i]); {ITEM_ROUNDS} rounds of mixer + cache line s[0] & mask; final mixer.\n" + )); + s.push_str(&format!("{fn_} void mh_item({cptr} cache, {u} t, {lptr} s) {{\n")); + for i in 0..8 { + s.push_str(&format!(" s[{i}] = {};\n", hex(k[i]))); + } + for i in 0..8 { + s.push_str(&format!(" s[{}] = t * {} + {};\n", 8 + i, hex(m[i]), hex(c[i]))); + } + s.push_str(&format!(" for ({u} r = 0u; r < {ITEM_ROUNDS}u; ++r) {{\n")); + s.push_str(" mh_mixer(s, 0x9E3779B9u * (r + 1u));\n"); + s.push_str(&format!(" {cptr} line = cache + ((s[0] & MH_CACHE_LINE_MASK) * 16u);\n")); + s.push_str(&format!(" for ({u} i = 0u; i < 16u; ++i) s[i] ^= line[i];\n")); + s.push_str(" }\n"); + s.push_str(&format!(" mh_mixer(s, 0x9E3779B9u * {}u);\n", ITEM_ROUNDS + 1)); + s.push_str("}\n"); + s.push_str("// dataset[w] without the dataset: derive item w >> 4 and take word w & 15.\n"); + s.push_str(&format!( + "{fn_} {u} mh_word({cptr} cache, {u} w) {{ {u} s[16]; mh_item(cache, w >> 4u, s); return s[w & 15u]; }}\n" + )); + s +} + +/// Metal library with the cache fill and dataset build kernels for one day key (`memhardMSL`, memhard.metal). +pub fn metal_memhard(mp: &MixParams) -> String { + let mut s = String::new(); + s.push_str("#include \n"); + s.push_str("using namespace metal;\n"); + s.push_str(&emit_memhard_core(mp, CoreDialect::Metal)); + s.push('\n'); + s.push_str(&format!("// One thread per segment (2^{} threads).\n", log2_segments())); + s.push_str( + "kernel void igneum_cache_fill(device uint* cache [[buffer(0)]], uint gid [[thread_position_in_grid]]) {\n", + ); + s.push_str(" mh_cache_segment(cache, gid);\n"); + s.push_str("}\n"); + s.push_str("// One thread per 64-byte item (dataset words / 16 threads).\n"); + s.push_str( + "kernel void igneum_build(device const uint* cache [[buffer(0)]], device uint* dataset [[buffer(1)]],\n", + ); + s.push_str(" uint gid [[thread_position_in_grid]]) {\n"); + s.push_str(" uint s[16];\n"); + s.push_str(" mh_item(cache, gid, s);\n"); + s.push_str(" device uint* d = dataset + gid * 16u;\n"); + s.push_str(" for (uint i = 0u; i < 16u; ++i) d[i] = s[i];\n"); + s.push_str("}\n"); + s +} + +const DS_ELEM_BODY: &str = " x *= 0x9E3779B1u; x ^= x >> 15;\n x += d1;\n x *= 0x85EBCA77u; x ^= x >> 13;\n x *= 0xC2B2AE3Du; x ^= x >> 16;\n return x;\n}\n"; + +/// The Metal hash kernel (`generateMSL`, program.metal). +pub fn metal_program(p: &Program, dataset_log2: u32, source: LoadSource) -> String { + let mask = mask_for(dataset_log2); + let mut s = String::with_capacity(5000); + s.push_str("#include \n"); + s.push_str("using namespace metal;\n"); + s.push('\n'); + s.push_str(&format!("#define MASK {}\n", hex(mask))); + s.push_str(&format!("constant uint SEEDW[8] = {{ {} }};\n", join_hex(&p.seed))); + s.push('\n'); + s.push_str("inline uint splitmix32(uint x) {\n"); + s.push_str(" x ^= x >> 16; x *= 0x7feb352du;\n"); + s.push_str(" x ^= x >> 15; x *= 0x846ca68bu;\n"); + s.push_str(" x ^= x >> 16;\n"); + 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("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() { + s.push_str("#define WMASK (MASK & ~31u)\n\n"); + } + let mut buffer0 = "device const uint* dataset [[buffer(0)]]"; + if let LoadSource::InlineMemhard(mp) = &source { + s.push_str(&emit_memhard_core(mp, CoreDialect::Metal)); + s.push('\n'); + buffer0 = "device const uint* cache [[buffer(0)]]"; + } + s.push_str(&format!("kernel void igneum_hash({buffer0},\n")); + s.push_str(" device ulong* out [[buffer(1)]],\n"); + s.push_str(" constant uint& baseNonce [[buffer(2)]],\n"); + s.push_str(" uint gid [[thread_position_in_grid]]) {\n"); + s.push_str(" uint nonce = baseNonce + gid;\n"); + 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"); + } + for i in 0..8 { + s.push_str(&format!( + " {{ uint x = nonce ^ SEEDW[{i}]; x += 0x9e3779b9u * {}u; x = splitmix32(x); r{i} = x ^ SEEDW[{}]; }}\n", + i + 1, + (i + 1) & 7 + )); + } + s.push_str(&format!("\n for (uint it = 0u; it < {ITERATIONS}u; ++it) {{\n uint sel = r0;\n")); + let word_index = |a: &str, wide: bool| -> String { + if wide { + format!("(simd_broadcast({a}, 0) & WMASK) + lane") + } else { + format!("{a} & MASK") + } + }; + let fetch = |idx: String| -> String { + match &source { + LoadSource::Stored => format!("dataset[{idx}]"), + LoadSource::InlineClosed(d0, d1) => format!("ds_elem({idx}, {}, {})", hex(*d0), hex(*d1)), + LoadSource::InlineMemhard(_) => format!("mh_word(cache, {idx})"), + } + }; + 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); + let line = match ins.op { + Op::Add => format!( + "{d} = {d} + {a} + select({}, {}, ((sel >> {}u) & 1u) != 0u);", + hex(ins.imm), + hex(ins.imm2), + ins.bit + ), + Op::Sub => format!("{d} = {d} - {a};"), + Op::Mul => format!("{d} = {d} * {a};"), + Op::MulHi => format!("{d} = mulhi({d}, {a});"), + Op::Xor => format!("{d} = {d} ^ {a};"), + Op::Or => format!("{d} = {d} | {a};"), + Op::Rotl => format!("{d} = rotl_imm({d}, {}u);", ins.rot), + Op::Rotr => format!("{d} = rotr_var({d}, {a});"), + Op::Mad => format!("{d} = {a} * {b} + {d};"), + Op::Shfl => format!("{d} = {d} ^ simd_shuffle_xor({a}, (ushort){});", ins.mask), + Op::Load => format!("{d} = {d} ^ {};", fetch(word_index(&a, false))), + Op::WLoad => format!("{d} = {d} ^ {};", fetch(word_index(&a, true))), + }; + s.push_str(&format!(" {line} // {k}\n")); + } + s.push_str(" }\n"); + 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"); + s.push_str("}\n"); + s +} + +/// The Metal closed-form fill kernel (`fillMSL`). +pub const METAL_FILL: &str = "#include \nusing namespace metal;\ninline uint ds_elem(uint i, uint d0, uint d1) {\n uint x = i ^ d0;\n x *= 0x9E3779B1u; x ^= x >> 15;\n x += d1;\n x *= 0x85EBCA77u; x ^= x >> 13;\n x *= 0xC2B2AE3Du; x ^= x >> 16;\n return x;\n}\nkernel void igneum_fill(device uint* dataset [[buffer(0)]],\n constant uint2& day [[buffer(1)]],\n uint gid [[thread_position_in_grid]]) {\n dataset[gid] = ds_elem(gid, day.x, day.y);\n}"; + +fn generated_by(seed: &str) -> String { + format!("// Generated by proto-metal/igneum-bench --export-pack for seed \"{seed}\". Do not edit by hand.\n") +} + +fn init_line(p: &Program, u: &str, i: usize) -> String { + let addc = 0x9e3779b9u32.wrapping_mul(i as u32 + 1); + format!( + " {{ {u} x = nonce ^ {}; x += {}; x = splitmix32(x); r{i} = x ^ {}; }} // SEEDW[{i}], 0x9e3779b9u * {}u, SEEDW[{}]\n", + hex(p.seed[i]), + hex(addc), + hex(p.seed[(i + 1) & 7]), + i + 1, + (i + 1) & 7 + ) +} + +/// The CUDA kernel (`generateCUDA`, kernel.cu). `memhard` is `None` for a closed-form pack. +pub fn cuda_kernel(p: &Program, memhard: Option<&MixParams>) -> String { + let mut s = String::with_capacity(9000); + s.push_str(&generated_by(&p.seed_string)); + s.push_str( + "// Bit-exact twin of the Metal kernel for the same seed (see proto-cuda/CHECKLIST.md and program.metal).\n", + ); + s.push_str("// Compiled ahead of time by nvcc together with proto-cuda/host.cu. No NVRTC.\n"); + s.push_str("#include \n"); + s.push_str("#include \n"); + s.push_str("#include \"program.h\"\n"); + if memhard.is_some() { + s.push_str("#include \"memhard.h\"\n"); + } + s.push('\n'); + s.push_str("__device__ __forceinline__ uint32_t splitmix32(uint32_t x) {\n"); + s.push_str(" x ^= x >> 16; x *= 0x7feb352du;\n"); + s.push_str(" x ^= x >> 15; x *= 0x846ca68bu;\n"); + s.push_str(" x ^= x >> 16;\n"); + s.push_str(" return x;\n"); + 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("// 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"); + s.push_str(" uint32_t x = i ^ d0;\n"); + s.push_str(DS_ELEM_BODY); + s.push('\n'); + if memhard.is_none() { + s.push_str("// dataset[i] = ds_elem(i, d0, d1) for i < n. Same closed form as the Metal igneum_fill kernel.\n"); + s.push_str("__global__ void igneum_fill(uint32_t* ds, uint32_t n, uint32_t d0, uint32_t d1) {\n"); + s.push_str(" uint32_t i = blockIdx.x * blockDim.x + threadIdx.x;\n"); + s.push_str(" if (i < n) ds[i] = ds_elem(i, d0, d1);\n"); + s.push_str("}\n"); + s.push('\n'); + } else { + s.push_str( + "// Memory-hard dataset (MEMHARD.md). One thread per cache segment; one thread per 64-byte dataset item.\n", + ); + s.push_str( + "// The core functions (mh_cache_segment, mh_item) are in memhard.h and are also compiled for the host.\n", + ); + s.push_str("__global__ void igneum_cache_fill(uint32_t* cache, uint32_t nSegments) {\n"); + s.push_str(" uint32_t seg = blockIdx.x * blockDim.x + threadIdx.x;\n"); + s.push_str(" if (seg < nSegments) mh_cache_segment(cache, seg);\n"); + s.push_str("}\n"); + s.push_str("__global__ void igneum_build(uint32_t* ds, const uint32_t* cache, uint32_t nItems) {\n"); + s.push_str(" uint32_t t = blockIdx.x * blockDim.x + threadIdx.x;\n"); + s.push_str(" if (t < nItems) {\n"); + s.push_str(" uint32_t s[16];\n"); + s.push_str(" mh_item(cache, t, s);\n"); + s.push_str(" uint32_t* d = ds + (size_t)t * 16u;\n"); + s.push_str(" for (uint32_t i = 0u; i < 16u; ++i) d[i] = s[i];\n"); + s.push_str(" }\n"); + s.push_str("}\n"); + s.push('\n'); + } + s.push_str("// One hash per thread. blockDim.x is a multiple of 32; lane = threadIdx.x & 31 and every\n"); + s.push_str("// __shfl_xor_sync stays inside the lane's own warp, exactly like simd_shuffle_xor inside a\n"); + s.push_str("// 32-wide Metal SIMD group. Control flow is uniform, so the full 0xffffffff member mask is valid.\n"); + s.push_str("__global__ void igneum_hash(const uint32_t* ds, uint64_t* out, uint32_t baseNonce, uint32_t mask) {\n"); + 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"); + 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(&format!("\n for (uint32_t it = 0u; it < {ITERATIONS}u; ++it) {{\n uint32_t sel = r0;\n")); + 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); + 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!( + "{d} = {d} + {a} + ((((sel >> {}u) & 1u) != 0u) ? {} : {});", + ins.bit, + hex(ins.imm2), + hex(ins.imm) + ), + Op::Sub => format!("{d} = {d} - {a};"), + Op::Mul => format!("{d} = {d} * {a};"), + Op::MulHi => format!("{d} = __umulhi({d}, {a});"), + Op::Xor => format!("{d} = {d} ^ {a};"), + Op::Or => format!("{d} = {d} | {a};"), + Op::Rotl => format!("{d} = rotl_imm({d}, {}u);", ins.rot), + Op::Rotr => format!("{d} = rotr_var({d}, {a});"), + Op::Mad => format!("{d} = {a} * {b} + {d};"), + Op::Shfl => format!("{d} = {d} ^ __shfl_xor_sync(0xffffffffu, {a}, {});", ins.mask), + Op::Load => format!("{d} = {d} ^ ds[{a} & mask];"), + Op::WLoad => format!("{d} = {d} ^ ds[(__shfl_sync(0xffffffffu, {a}, 0) & wmask) + lane];"), + }; + s.push_str(&format!(" {line} // {k} {}\n", ins.op.name())); + } + s.push_str(" }\n"); + 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"); + s.push_str("}\n"); + s.push('\n'); + s.push_str("// Host-side launch wrappers. Declared in program.h, called from host.cu.\n"); + if memhard.is_none() { + s.push_str("cudaError_t igneum_launch_fill(uint32_t* ds, uint32_t nWords, uint32_t d0, uint32_t d1) {\n"); + s.push_str(" if (nWords == 0u) return cudaErrorInvalidValue;\n"); + s.push_str(" uint32_t block = 256u;\n"); + s.push_str(" uint32_t grid = (nWords + block - 1u) / block;\n"); + s.push_str(" igneum_fill<<>>(ds, nWords, d0, d1);\n"); + s.push_str(" return cudaGetLastError();\n"); + s.push_str("}\n"); + s.push('\n'); + } else { + s.push_str("cudaError_t igneum_launch_cache_fill(uint32_t* cache, uint32_t nSegments) {\n"); + s.push_str(" if (nSegments == 0u) return cudaErrorInvalidValue;\n"); + s.push_str(" uint32_t block = 256u;\n"); + s.push_str(" uint32_t grid = (nSegments + block - 1u) / block;\n"); + s.push_str(" igneum_cache_fill<<>>(cache, nSegments);\n"); + s.push_str(" return cudaGetLastError();\n"); + s.push_str("}\n"); + s.push('\n'); + s.push_str("cudaError_t igneum_launch_build(uint32_t* ds, const uint32_t* cache, uint32_t nItems) {\n"); + s.push_str(" if (nItems == 0u) return cudaErrorInvalidValue;\n"); + s.push_str(" uint32_t block = 256u;\n"); + s.push_str(" uint32_t grid = (nItems + block - 1u) / block;\n"); + s.push_str(" igneum_build<<>>(ds, cache, nItems);\n"); + s.push_str(" return cudaGetLastError();\n"); + s.push_str("}\n"); + s.push('\n'); + } + s.push_str( + "cudaError_t igneum_launch_hash(const uint32_t* ds, uint64_t* out, uint32_t baseNonce, uint32_t mask,\n", + ); + s.push_str(" uint32_t nonces, uint32_t blockWarps) {\n"); + s.push_str(" if (blockWarps == 0u || blockWarps > 32u) return cudaErrorInvalidValue;\n"); + s.push_str(" uint32_t block = 32u * blockWarps;\n"); + s.push_str(" if (nonces == 0u || (nonces % block) != 0u) return cudaErrorInvalidValue;\n"); + s.push_str(" igneum_hash<<>>(ds, out, baseNonce, mask);\n"); + s.push_str(" return cudaGetLastError();\n"); + s.push_str("}\n"); + s.push('\n'); + s.push_str("cudaError_t igneum_hash_info(int* numRegs, int* blocksPerSM, uint32_t blockWarps) {\n"); + s.push_str(" cudaFuncAttributes attr;\n"); + s.push_str(" cudaError_t e = cudaFuncGetAttributes(&attr, igneum_hash);\n"); + s.push_str(" if (e != cudaSuccess) return e;\n"); + s.push_str(" *numRegs = attr.numRegs;\n"); + s.push_str(" return cudaOccupancyMaxActiveBlocksPerMultiprocessor(blocksPerSM, igneum_hash, (int)(32u * blockWarps), 0);\n"); + s.push_str("}\n"); + s +} + +/// The OpenCL C 1.2 kernel (`generateOpenCL`, kernel.cl). +pub fn opencl_kernel(p: &Program, memhard: Option<&MixParams>) -> String { + let mut s = String::with_capacity(14000); + s.push_str(&generated_by(&p.seed_string)); + s.push_str("// OpenCL C twin of the Metal kernel for the same seed (see proto-opencl/README.md, WAVEFRONT.md and program.metal).\n"); + s.push_str("// Built from source at runtime by proto-opencl/host.c, which passes these defines:\n"); + s.push_str("// IGNEUM_GROUP work-group size of igneum_hash, a multiple of 32 (default 32: one work-group = one 32-lane unit)\n"); + s.push_str( + "// IGNEUM_EXCHANGE 0 = local-memory exchange with a barrier (any device, any wave width; the default)\n", + ); + s.push_str("// 1 = sub_group_shuffle_xor (cl_khr_subgroup_shuffle), only with IGNEUM_GROUP 32 and a sub-group size of exactly 32\n"); + s.push_str("// 2 = intel_sub_group_shuffle_xor (cl_intel_subgroups), same condition\n"); + s.push_str("// The verification unit is always 32 lanes. A 64-wide hardware wave (AMD GCN/CDNA, RDNA in wave64) runs two units;\n"); + 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("#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"); + s.push_str("#if IGNEUM_EXCHANGE == 1\n"); + s.push_str("#ifdef cl_khr_subgroups\n#pragma OPENCL EXTENSION cl_khr_subgroups : enable\n#endif\n"); + s.push_str("#ifdef cl_khr_subgroup_shuffle\n#pragma OPENCL EXTENSION cl_khr_subgroup_shuffle : enable\n#endif\n"); + s.push_str("#elif IGNEUM_EXCHANGE == 2\n#pragma OPENCL EXTENSION cl_intel_subgroups : enable\n#endif\n"); + s.push_str("#else\n"); + s.push_str("// Not an OpenCL compiler: proto-opencl/emu compiles this file as C++ and supplies the built-ins and these two macros.\n"); + s.push_str("#include \"emu_opencl.h\"\n#endif\n"); + s.push('\n'); + s.push_str("#if IGNEUM_EXCHANGE == 1\n"); + s.push_str("#define IGNEUM_SHFL_XOR(dst, a, m) dst = sub_group_shuffle_xor((a), (uint)(m))\n"); + s.push_str("#define IGNEUM_BCAST0(dst, a) dst = sub_group_broadcast((a), 0u)\n"); + s.push_str("#elif IGNEUM_EXCHANGE == 2\n"); + s.push_str("#define IGNEUM_SHFL_XOR(dst, a, m) dst = intel_sub_group_shuffle_xor((a), (uint)(m))\n"); + s.push_str("#define IGNEUM_BCAST0(dst, a) dst = sub_group_broadcast((a), 0u)\n"); + s.push_str("#else\n"); + s.push_str("// Local-memory exchange. Two buffers of IGNEUM_GROUP words alternate (xk counts exchanges), so one barrier per\n"); + s.push_str("// exchange is enough: a lane can only overwrite buffer b at exchange k+2 after passing barrier k+1, and every lane\n"); + s.push_str("// reaches barrier k+1 only after its read of buffer b at exchange k. The partner lid ^ m stays inside the lane's\n"); + s.push_str( + "// aligned run of 32 because m < 32. Control flow is uniform, so every work-item reaches every barrier.\n", + ); + s.push_str("#define IGNEUM_SHFL_XOR(dst, a, m) { xch[(xk & 1u) * IGNEUM_GROUP + lid] = (a); barrier(CLK_LOCAL_MEM_FENCE); dst = xch[(xk & 1u) * IGNEUM_GROUP + (lid ^ (uint)(m))]; xk += 1u; }\n"); + s.push_str("#define IGNEUM_BCAST0(dst, a) { xch[(xk & 1u) * IGNEUM_GROUP + lid] = (a); barrier(CLK_LOCAL_MEM_FENCE); dst = xch[(xk & 1u) * IGNEUM_GROUP + (lid & ~31u)]; xk += 1u; }\n"); + s.push_str("#endif\n"); + s.push('\n'); + s.push_str("static inline uint splitmix32(uint x) {\n"); + s.push_str(" x ^= x >> 16; x *= 0x7feb352du;\n"); + s.push_str(" x ^= x >> 15; x *= 0x846ca68bu;\n"); + s.push_str(" x ^= x >> 16;\n"); + s.push_str(" return x;\n"); + 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("// 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"); + s.push_str(" uint x = i ^ d0;\n"); + s.push_str(DS_ELEM_BODY); + s.push('\n'); + if let Some(mp) = memhard { + s.push_str(&emit_memhard_core(mp, CoreDialect::OpenCl)); + s.push('\n'); + s.push_str("// Memory-hard dataset (MEMHARD.md). One work-item per cache segment; one work-item per 64-byte dataset item.\n"); + s.push_str("// The same constants as memhard.h in this pack (one emitter, three dialects).\n"); + s.push_str("__kernel void igneum_cache_fill(__global uint* cache, uint nSegments) {\n"); + s.push_str(" uint seg = (uint)get_global_id(0);\n"); + s.push_str(" if (seg < nSegments) mh_cache_segment(cache, seg);\n"); + s.push_str("}\n"); + s.push_str("__kernel void igneum_build(__global uint* ds, __global const uint* cache, uint nItems) {\n"); + s.push_str(" uint t = (uint)get_global_id(0);\n"); + s.push_str(" if (t < nItems) {\n"); + s.push_str(" uint s[16];\n"); + s.push_str(" mh_item(cache, t, s);\n"); + s.push_str(" __global uint* d = ds + ((ulong)t * 16u);\n"); + s.push_str(" for (uint i = 0u; i < 16u; ++i) d[i] = s[i];\n"); + s.push_str(" }\n"); + s.push_str("}\n"); + s.push('\n'); + } else { + s.push_str("// dataset[i] = ds_elem(i, d0, d1) for i < n. Same closed form as the Metal igneum_fill kernel.\n"); + s.push_str("__kernel void igneum_fill(__global uint* ds, uint n, uint d0, uint d1) {\n"); + s.push_str(" uint i = (uint)get_global_id(0);\n"); + s.push_str(" if (i < n) ds[i] = ds_elem(i, d0, d1);\n"); + s.push_str("}\n"); + s.push('\n'); + } + s.push_str("// One hash per work-item. IGNEUM_GROUP is a multiple of 32; lane = lid & 31 and every exchange stays inside the\n"); + s.push_str("// lane's own aligned run of 32 work-items, exactly like simd_shuffle_xor inside a 32-wide Metal SIMD group and\n"); + s.push_str("// __shfl_xor_sync inside a CUDA warp. Control flow is uniform (no branches at all).\n"); + s.push_str("IGNEUM_KERNEL_HASH void igneum_hash(__global const uint* ds, __global ulong* out, uint baseNonce, uint mask) {\n"); + s.push_str(" uint gid = (uint)get_global_id(0);\n"); + s.push_str(" uint lid = (uint)get_local_id(0);\n"); + s.push_str(" uint nonce = baseNonce + gid;\n"); + 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"); + s.push_str("#else\n"); + s.push_str(" (void)lid;\n"); + s.push_str("#endif\n"); + if p.has_wide() { + 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(&format!("\n for (uint it = 0u; it < {ITERATIONS}u; ++it) {{\n uint sel = r0;\n")); + 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); + let line = match ins.op { + Op::Add => format!( + "{d} = {d} + {a} + ((((sel >> {}u) & 1u) != 0u) ? {} : {});", + ins.bit, + hex(ins.imm2), + hex(ins.imm) + ), + Op::Sub => format!("{d} = {d} - {a};"), + Op::Mul => format!("{d} = {d} * {a};"), + Op::MulHi => format!("{d} = mul_hi({d}, {a});"), + Op::Xor => format!("{d} = {d} ^ {a};"), + Op::Or => format!("{d} = {d} | {a};"), + Op::Rotl => format!("{d} = rotl_imm({d}, {}u);", ins.rot), + Op::Rotr => format!("{d} = rotr_var({d}, {a});"), + 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 => format!("{d} = {d} ^ ds[{a} & mask];"), + Op::WLoad => format!("{{ uint t_; IGNEUM_BCAST0(t_, {a}); {d} = {d} ^ ds[(t_ & wmask) + lane]; }}"), + }; + s.push_str(&format!(" {line} // {k} {}\n", ins.op.name())); + } + s.push_str(" }\n"); + 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"); + s.push_str("}\n"); + s.push('\n'); + s.push_str("#if IGNEUM_EXCHANGE != 0\n"); + s.push_str("// Reports the sub-group size this device uses for a work-group of IGNEUM_GROUP items. host.c runs it only when the\n"); + s.push_str("// per-kernel query (clGetKernelSubGroupInfoKHR on igneum_hash) is unavailable; that query is preferred because a\n"); + s.push_str("// compiler may pick a different wave width per kernel (RDNA: wave32 or wave64). See WAVEFRONT.md.\n"); + s.push_str("IGNEUM_KERNEL_HASH void igneum_probe_subgroup(__global uint* out) {\n"); + s.push_str(" if (get_local_id(0) == 0u) { out[0] = get_sub_group_size(); out[1] = get_num_sub_groups(); }\n"); + s.push_str("}\n"); + s.push_str("#endif\n"); + s +} + +fn mask_for(dataset_log2: u32) -> u32 { + if dataset_log2 >= 32 { + u32::MAX + } else { + (1u32 << dataset_log2) - 1 + } +} + +const STDINT_BLOCK: &str = "#ifdef __cplusplus\n#include \n#else\n#include \n#endif\n"; + +/// program.h (`generateProgramHeader`). +pub fn program_header( + p: &Program, + day: &str, + key: &[u32; 8], + dataset_log2: u32, + memhard: Option<&MixParams>, +) -> String { + 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"); + s.push_str("// Also included by proto-opencl/host.c (C99), which defines IGNEUM_NO_CUDA first and reads only the macros.\n"); + s.push_str("#pragma once\n"); + s.push_str(STDINT_BLOCK); + s.push_str("#ifndef IGNEUM_NO_CUDA\n#include \n#endif\n"); + s.push('\n'); + s.push_str(&format!("#define IGNEUM_SEED_STRING {}\n", jstr(&p.seed_string))); + s.push_str(&format!("#define IGNEUM_DAY_STRING {}\n", jstr(day))); + s.push_str(&format!("#define IGNEUM_DAY0 {}\n", hex(key[0]))); + s.push_str(&format!("#define IGNEUM_DAY1 {}\n", hex(key[1]))); + 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")); + s.push_str(&format!("#define IGNEUM_LOADS_PER_HASH {}\n", p.loads_per_hash())); + s.push_str(&format!("#define IGNEUM_WIDE_LOADS_PER_HASH {}\n", p.wide_loads_per_hash())); + s.push_str(&format!("#define IGNEUM_OP_MIX {}\n", jstr(&p.op_mix()))); + s.push_str("// 0 = closed-form dataset (ds_elem), 1 = memory-hard cache construction (MEMHARD.md, memhard.h)\n"); + s.push_str(&format!("#define IGNEUM_DATASET_MODE {}\n", if memhard.is_some() { 1 } else { 0 })); + s.push('\n'); + s.push_str(&format!("#define IGNEUM_SEEDW_INIT {{ {} }}\n", join_hex(&p.seed))); + if let Some(mp) = memhard { + s.push_str(&format!("#define IGNEUM_KEY_INIT {{ {} }}\n", join_hex(&mp.key))); + s.push_str(&format!("#define IGNEUM_CACHE_LOG2_WORDS {CACHE_LOG2_WORDS}\n")); + s.push_str(&format!("#define IGNEUM_CACHE_SEGMENT_LOG2_LINES {CACHE_SEGMENT_LOG2_LINES}\n")); + s.push_str(&format!("#define IGNEUM_CACHE_SEGMENTS {CACHE_SEGMENTS}u\n")); + s.push_str(&format!("#define IGNEUM_ITEM_ROUNDS {ITEM_ROUNDS}\n")); + s.push_str(&format!( + "#define IGNEUM_MIX_ROT_INIT {{ {} }}\n", + mp.rot.iter().map(|r| format!("{r}u")).collect::>().join(", ") + )); + s.push_str(&format!("#define IGNEUM_MIX_MUL_INIT {{ {} }}\n", join_hex(&mp.mul))); + s.push_str(&format!("#define IGNEUM_MIX_RC_INIT {{ {} }}\n", join_hex(&mp.rc))); + s.push('\n'); + s.push_str("#ifndef IGNEUM_NO_CUDA\n"); + s.push_str("// Defined in kernel.cu. All launch on the default stream and return cudaGetLastError().\n"); + s.push_str("cudaError_t igneum_launch_cache_fill(uint32_t* cache, uint32_t nSegments);\n"); + s.push_str("cudaError_t igneum_launch_build(uint32_t* ds, const uint32_t* cache, uint32_t nItems);\n"); + } else { + s.push_str("#ifndef IGNEUM_NO_CUDA\n"); + s.push_str("// Defined in kernel.cu. Both launch on the default stream and return cudaGetLastError().\n"); + s.push_str("cudaError_t igneum_launch_fill(uint32_t* ds, uint32_t nWords, uint32_t d0, uint32_t d1);\n"); + } + s.push_str( + "cudaError_t igneum_launch_hash(const uint32_t* ds, uint64_t* out, uint32_t baseNonce, uint32_t mask,\n", + ); + s.push_str(" uint32_t nonces, uint32_t blockWarps);\n"); + s.push_str("cudaError_t igneum_hash_info(int* numRegs, int* blocksPerSM, uint32_t blockWarps);\n"); + s.push_str("#endif\n"); + s +} + +/// memhard.h (`generateMemhardHeader`): the core in CUDA C++, compiled for host and device. +pub fn cuda_memhard_header(p: &Program, mp: &MixParams) -> String { + let mut s = String::with_capacity(6000); + s.push_str(&generated_by(&p.seed_string)); + s.push_str("// Memory-hard dataset core, the same text that the Mac's Metal kernels and CPU verifier were checked against.\n"); + s.push_str( + "// Included by kernel.cu (device), host.cu (host reference) and proto-opencl/host.c (C99 host reference).\n", + ); + s.push_str("// See proto-metal/MEMHARD.md for the construction. kernel.cl carries the same text in OpenCL C.\n"); + s.push_str("#pragma once\n"); + s.push_str(STDINT_BLOCK); + s.push_str("#if defined(__CUDACC__)\n"); + s.push_str("#define IGNEUM_HD __host__ __device__ __forceinline__\n"); + s.push_str("#elif defined(_MSC_VER) && !defined(__cplusplus)\n"); + s.push_str("#define IGNEUM_HD static __inline\n"); + s.push_str("#else\n"); + s.push_str("#define IGNEUM_HD static inline\n"); + s.push_str("#endif\n"); + s.push_str(&emit_memhard_core(mp, CoreDialect::Cuda)); + s +} + +/// The self-test values a pack carries beside the 96 hashes. +#[derive(Clone, Debug, Default, PartialEq, Eq)] +pub struct PackVectors { + /// dataset[0..15] + pub head: Vec, + /// dataset[MASK] + pub last: u32, + /// 64 sampled dataset indices and their values + pub sample_idx: Vec, + pub sample_val: Vec, + /// cache[0..15] (memory-hard only) + pub cache_head: Vec, + /// the last cache line (memory-hard only) + pub cache_last: Vec, + /// FNV-1a 64 over the whole cache (memory-hard only) + pub cache_fnv: u64, +} + +/// The base nonces of the three vector warps every pack carries. +pub const PACK_VECTOR_BASES: [u32; 3] = [0, 4096, 1_000_000]; + +/// The 64 sampled dataset indices: SplitMix64 seeded with "mhsample", low 32 bits masked. +pub fn sample_indices(mask: u32) -> Vec { + let mut sr = SplitMix64::new(0x6d68_7361_6d70_6c65); + (0..64).map(|_| (sr.next() as u32) & mask).collect() +} + +/// vectors.h (`generateVectorsHeader`). +pub fn vectors_header( + p: &Program, + bases: &[u32], + outs: &[[u64; 32]], + v: &PackVectors, + mask: u32, + source: &str, + memhard: bool, +) -> String { + let mut s = String::with_capacity(6000); + s.push_str(&generated_by(&p.seed_string)); + s.push_str(&format!("// Expected outputs: {source}\n")); + s.push_str("#pragma once\n"); + s.push_str(STDINT_BLOCK); + s.push('\n'); + s.push_str(&format!("#define IGNEUM_VEC_WARPS {}\n", bases.len())); + s.push_str(&format!( + "static const uint32_t IGNEUM_VEC_BASE[IGNEUM_VEC_WARPS] = {{ {} }};\n", + bases.iter().map(|b| format!("{b}u")).collect::>().join(", ") + )); + s.push_str("static const uint64_t IGNEUM_VEC_OUT[IGNEUM_VEC_WARPS][32] = {\n"); + for (i, o) in outs.iter().enumerate() { + s.push_str(&format!(" {{ // base nonce {}\n", bases[i])); + for row in 0..4 { + s.push_str(" "); + s.push_str(&(0..8).map(|c| hex64(o[row * 8 + c])).collect::>().join(", ")); + s.push_str(if row == 3 { "\n" } else { ",\n" }); + } + s.push_str(if i == outs.len() - 1 { " }\n" } else { " },\n" }); + } + s.push_str("};\n"); + s.push('\n'); + s.push_str(&format!("// Dataset self-test: dataset[0..15] and dataset[IGNEUM_MASK] ({mask}).\n")); + s.push_str("static const uint32_t IGNEUM_DS_HEAD[16] = {\n"); + s.push_str(&format!(" {},\n", join_hex(&v.head[..8]))); + s.push_str(&format!(" {}\n", join_hex(&v.head[8..16]))); + s.push_str("};\n"); + s.push_str(&format!("static const uint32_t IGNEUM_DS_LAST_INDEX = {mask}u;\n")); + s.push_str(&format!("static const uint32_t IGNEUM_DS_LAST = {};\n", hex(v.last))); + s.push_str("// 64 sampled dataset words (index, value) computed on the Mac.\n"); + s.push_str(&format!("#define IGNEUM_DS_SAMPLES {}\n", v.sample_idx.len())); + s.push_str("static const uint32_t IGNEUM_DS_SAMPLE_INDEX[IGNEUM_DS_SAMPLES] = {\n"); + s.push_str(&format!(" {}\n", v.sample_idx.iter().map(|i| format!("{i}u")).collect::>().join(", "))); + s.push_str("};\n"); + s.push_str("static const uint32_t IGNEUM_DS_SAMPLE_VALUE[IGNEUM_DS_SAMPLES] = {\n"); + s.push_str(&format!(" {}\n", join_hex(&v.sample_val))); + s.push_str("};\n"); + if memhard { + s.push_str(&format!( + "// Cache self-test (memory-hard mode): cache[0..15], the last 16 words, and FNV-1a 64 over all 2^{CACHE_LOG2_WORDS} words.\n" + )); + s.push_str("static const uint32_t IGNEUM_CACHE_HEAD[16] = {\n"); + s.push_str(&format!(" {},\n", join_hex(&v.cache_head[..8]))); + s.push_str(&format!(" {}\n", join_hex(&v.cache_head[8..16]))); + s.push_str("};\n"); + s.push_str("static const uint32_t IGNEUM_CACHE_LAST[16] = {\n"); + s.push_str(&format!(" {},\n", join_hex(&v.cache_last[..8]))); + s.push_str(&format!(" {}\n", join_hex(&v.cache_last[8..16]))); + s.push_str("};\n"); + s.push_str(&format!("static const uint64_t IGNEUM_CACHE_FNV64 = {};\n", hex64(v.cache_fnv))); + } + s +} + +/// program.json (`generateProgramJSON`). Valid JSON (see the module note about the `"item"` line). +pub fn program_json(p: &Program, day: &str, key: &[u32; 8], dataset_log2: u32, memhard: Option<&MixParams>) -> String { + let mask = mask_for(dataset_log2); + let mut s = String::with_capacity(13000); + s.push_str("{\n"); + s.push_str(" \"format\": \"igneum-program-pack-2\",\n"); + s.push_str(&format!( + " \"dataset_mode\": {},\n", + jstr(if memhard.is_some() { "memory-hard" } else { "closed-form" }) + )); + s.push_str(&format!(" \"seed\": {},\n", jstr(&p.seed_string))); + s.push_str(&format!(" \"seed_words\": [{}],\n", join_jhex(&p.seed))); + s.push_str(" \"seed_derivation\": \"FNV-1a 64 over UTF-8 of seed, basis ^ (salt * 0x9E3779B97F4A7C15) for salt 0..3, then h ^= h>>33; h *= 0xff51afd7ed558ccd; h ^= h>>33; words[2*salt] = low 32, words[2*salt+1] = high 32\",\n"); + s.push_str(" \"lanes\": 32,\n"); + s.push_str(" \"registers\": 8,\n"); + s.push_str(&format!(" \"iterations\": {ITERATIONS},\n")); + s.push_str(&format!(" \"instruction_count\": {INSTR_COUNT},\n")); + s.push_str(&format!(" \"loads_per_hash\": {},\n", p.loads_per_hash())); + s.push_str(&format!( + " \"op_mix\": {{{}}},\n", + p.histogram().iter().map(|(n, c)| format!("{}: {c}", jstr(n))).collect::>().join(", ") + )); + s.push_str(" \"register_init\": \"for i in 0..7: x = nonce ^ seed_words[i]; x += 0x9e3779b9 * (i+1) (mod 2^32); x = splitmix32(x); r[i] = x ^ seed_words[(i+1) & 7]\",\n"); + s.push_str(" \"splitmix32\": \"x ^= x>>16; x *= 0x7feb352d; x ^= x>>15; x *= 0x846ca68b; x ^= x>>16\",\n"); + s.push_str( + " \"iteration\": \"sel = r0 sampled once at the top of each iteration, then all instructions in order\",\n", + ); + s.push_str(" \"output\": \"lo = r0 ^ rotl(r1,7) ^ rotl(r2,14) ^ rotl(r3,21); hi = r4 ^ rotl(r5,9) ^ rotl(r6,18) ^ rotl(r7,27); out = (hi << 32) | lo\",\n"); + s.push_str(" \"op_semantics\": {\n"); + s.push_str(" \"add\": \"dst = dst + src + (bit `bit` of sel ? imm2 : imm)\",\n"); + s.push_str(" \"sub\": \"dst = dst - src\",\n"); + s.push_str(" \"mul\": \"dst = dst * src (low 32)\",\n"); + s.push_str(" \"mulhi\": \"dst = high 32 bits of dst * src\",\n"); + s.push_str(" \"xor\": \"dst = dst ^ src\",\n"); + s.push_str(" \"or\": \"dst = dst | src\",\n"); + s.push_str(" \"rotl\": \"dst = rotl(dst, rot), rot in 1..31\",\n"); + s.push_str(" \"rotr\": \"dst = rotr(dst, src & 31)\",\n"); + s.push_str(" \"mad\": \"dst = src * src2 + dst\",\n"); + s.push_str( + " \"shfl\": \"dst = dst ^ (src of lane (lane ^ mask)), mask in {1,2,4,8,16}, within the 32-lane warp\",\n", + ); + 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)\"\n"); + s.push_str(" },\n"); + s.push_str(" \"dataset\": {\n"); + 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_words_from\": {},\n", jstr(&format!("day/{day}")))); + s.push_str(&format!(" \"d0\": {},\n", jhex(key[0]))); + s.push_str(&format!(" \"d1\": {},\n", jhex(key[1]))); + if let Some(mp) = memhard { + s.push_str(" \"mode\": \"memory-hard\",\n"); + s.push_str(" \"spec\": \"proto-metal/MEMHARD.md\",\n"); + s.push_str(&format!(" \"key\": [{}],\n", join_jhex(&mp.key))); + s.push_str( + " \"key_derivation\": \"the 8 words of seedWords(\\\"day/\\\" + day); d0, d1 are key[0], key[1]\",\n", + ); + s.push_str(&format!( + " \"cache\": {{\"log2_words\": {CACHE_LOG2_WORDS}, \"bytes\": {}, \"line_words\": 16, \"segment_lines\": {CACHE_LINES_PER_SEGMENT}, \"segments\": {CACHE_SEGMENTS}, \"block\": \"ChaCha{CHACHA_ROUNDS} core + feed-forward, rotations 16 12 8 7\", \"sigma\": [{}], \"tag\": [{}], \"chain\": \"in_j = prev_line ^ (sigma[0..3] || key[0..7] || seg || j || tag[0..1]); line_j = block(in_j); prev_0 = 0\"}},\n", + CACHE_WORDS as u64 * 4, + join_jhex(&CHACHA_SIGMA), + join_jhex(&CACHE_TAG) + )); + s.push_str(&format!( + " \"mixer\": {{\"draw\": \"SplitMix64 seeded with key[0] | key[1] << 32: rot[0..7] = 1 + next() % 31, mul[0..15] = low32(next()) | 1, rc[0..15] = low32(next())\", \"rot\": [{}], \"mul\": [{}], \"rc\": [{}], \"round\": \"for i in 0..15: s[i] = (s[i] ^ (rc[i] + (r+1) * 0x9E3779B9)) * mul[i]; then quarter rounds on columns (0,4,8,12) (1,5,9,13) (2,6,10,14) (3,7,11,15) with rot[0..3] and diagonals (0,5,10,15) (1,6,11,12) (2,7,8,13) (3,4,9,14) with rot[4..7]\", \"quarter_round\": \"a += b; d ^= a; d = rotl(d, r1); c += d; b ^= c; b = rotl(b, r2); a += b; d ^= a; d = rotl(d, r3); c += d; b ^= c; b = rotl(b, r4)\"}},\n", + mp.rot.iter().map(|r| r.to_string()).collect::>().join(", "), + join_jhex(&mp.mul), + join_jhex(&mp.rc) + )); + // The Swift writes jhex(cacheLineMask) here, which breaks the JSON. We write the bare literal. + s.push_str(&format!( + " \"item\": \"s[0..7] = key; s[8+i] = t * mul[i] + rc[i] for i in 0..7; for r in 0..{}: s = M_r(s); line = s[0] & 0x{:08x}; s[i] ^= cache[line * 16 + i]; then s = M_{ITEM_ROUNDS}(s); item(t) = s\",\n", + ITEM_ROUNDS - 1, + CACHE_LINE_MASK + )); + s.push_str(" \"word\": \"dataset[w] = item(w >> 4)[w & 15]\"\n"); + } else { + s.push_str(" \"mode\": \"closed-form\",\n"); + s.push_str(" \"formula\": \"x = i ^ d0; x *= 0x9E3779B1; x ^= x>>15; x += d1; x *= 0x85EBCA77; x ^= x>>13; x *= 0xC2B2AE3D; x ^= x>>16 (all mod 2^32)\"\n"); + } + s.push_str(" },\n"); + s.push_str(" \"instructions\": [\n"); + let n = p.instrs.len(); + for (k, ins) in p.instrs.iter().enumerate() { + s.push_str(&format!( + " {{\"i\": {k}, \"op\": {}, \"dst\": {}, \"src\": {}, \"src2\": {}, \"imm\": {}, \"imm2\": {}, \"rot\": {}, \"bit\": {}, \"mask\": {}}}", + jstr(ins.op.name()), + ins.dst, + ins.src, + ins.src2, + jhex(ins.imm), + jhex(ins.imm2), + ins.rot, + ins.bit, + ins.mask + )); + s.push_str(if k == n - 1 { "\n" } else { ",\n" }); + } + s.push_str(" ]\n}\n"); + s +} + +/// vectors.json (`generateVectorsJSON`). +pub fn vectors_json( + p: &Program, + day: &str, + dataset_log2: u32, + bases: &[u32], + outs: &[[u64; 32]], + v: &PackVectors, + mask: u32, + source: &str, + memhard: bool, +) -> String { + 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")); + s.push_str(&format!(" \"mask\": {},\n", jhex(mask))); + s.push_str(" \"lanes\": 32,\n"); + s.push_str(&format!(" \"source\": {},\n", jstr(source))); + s.push_str(" \"warps\": [\n"); + for (i, o) in outs.iter().enumerate() { + s.push_str(&format!(" {{\"base_nonce\": {}, \"expected\": [\n", bases[i])); + for row in 0..4 { + s.push_str(" "); + s.push_str(&(0..8).map(|c| jhex64(o[row * 8 + c])).collect::>().join(", ")); + s.push_str(if row == 3 { "\n" } else { ",\n" }); + } + s.push_str(if i == outs.len() - 1 { " ]}\n" } else { " ]},\n" }); + } + s.push_str(" ],\n"); + s.push_str(&format!(" \"dataset_head\": [{}],\n", join_jhex(&v.head))); + s.push_str(&format!(" \"dataset_last_index\": {mask},\n")); + s.push_str(&format!(" \"dataset_last\": {},\n", jhex(v.last))); + s.push_str(&format!( + " \"dataset_samples\": [{}]", + v.sample_idx + .iter() + .zip(v.sample_val.iter()) + .map(|(i, val)| format!("{{\"index\": {i}, \"value\": {}}}", jhex(*val))) + .collect::>() + .join(", ") + )); + if memhard { + s.push_str(&format!(",\n \"cache_head\": [{}],\n", join_jhex(&v.cache_head))); + s.push_str(&format!(" \"cache_last_line\": [{}],\n", join_jhex(&v.cache_last))); + s.push_str(&format!(" \"cache_fnv1a64\": {}\n", jhex64(v.cache_fnv))); + } else { + s.push('\n'); + } + s.push_str("}\n"); + s +} + +/// A program pack: the files `--export-pack` writes, as (name, text). +pub struct Pack { + pub files: Vec<(String, String)>, + pub bases: Vec, + pub outs: Vec<[u64; 32]>, + pub vectors: PackVectors, +} + +impl Pack { + pub fn write_to(&self, dir: &std::path::Path) -> std::io::Result<()> { + std::fs::create_dir_all(dir)?; + for (name, text) in &self.files { + std::fs::write(dir.join(name), text)?; + } + Ok(()) + } +} + +/// Build the whole pack for an epoch: the three vector warps from the CPU interpreter, the self-test words, +/// and every source file. The `source` string says where the vectors came from. +pub fn export_pack(epoch: &Epoch, day: &str, source: &str) -> Pack { + let p = &epoch.program; + let ds: &DatasetSource = &epoch.dataset; + 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(); + let mut v = PackVectors { + head: (0..16).map(|i| ds.word(i)).collect(), + last: ds.word(mask), + sample_idx: sample_indices(mask), + ..Default::default() + }; + v.sample_val = v.sample_idx.iter().map(|&i| ds.word(i)).collect(); + if let Some(m) = ds.memhard() { + let w = m.cache.words(); + v.cache_head = w[..16].to_vec(); + v.cache_last = w[w.len() - 16..].to_vec(); + v.cache_fnv = m.cache.fnv1a64(); + } + let is_mh = memhard.is_some(); + let mut files = vec![ + ("program.json".to_string(), program_json(p, day, &ds.key, ds.log2_words, memhard)), + ("vectors.json".to_string(), vectors_json(p, day, ds.log2_words, &bases, &outs, &v, mask, source, is_mh)), + ("kernel.cu".to_string(), cuda_kernel(p, memhard)), + ("kernel.cl".to_string(), opencl_kernel(p, memhard)), + ("program.h".to_string(), program_header(p, day, &ds.key, ds.log2_words, memhard)), + ("vectors.h".to_string(), vectors_header(p, &bases, &outs, &v, mask, source, is_mh)), + ("program.metal".to_string(), metal_program(p, ds.log2_words, LoadSource::Stored)), + ]; + if let Some(mp) = memhard { + files.push(("memhard.h".to_string(), cuda_memhard_header(p, mp))); + files.push(("memhard.metal".to_string(), metal_memhard(mp))); + } + Pack { files, bases, outs, vectors: v } +} + +/// The dataset mode a pack was written in, from its program.json text (no JSON parser needed). +pub fn pack_mode_from_json(program_json: &str) -> DatasetMode { + if program_json.contains("\"dataset_mode\": \"memory-hard\"") { + DatasetMode::MemoryHard + } else { + DatasetMode::ClosedForm + } +} diff --git a/igneum-pow/src/generator.rs b/igneum-pow/src/generator.rs new file mode 100644 index 000000000..3240984e0 --- /dev/null +++ b/igneum-pow/src/generator.rs @@ -0,0 +1,276 @@ +//! The program generator: 64 integer instructions over 8 x u32 lane registers, run for 8 iterations. +//! Draw order, weights and the lever rules are those of `generateProgram` in `proto-metal/main.swift`. + +use crate::seed::{program_rng, seed_words}; + +/// Iterations of the instruction list per hash. +pub const ITERATIONS: usize = 8; +/// Instructions per program. +pub const INSTR_COUNT: usize = 64; +/// Lanes per verification unit (one SIMD group / warp). +pub const LANES: usize = 32; + +#[derive(Clone, Copy, Debug, PartialEq, Eq, Hash)] +pub enum Op { + Add, + Sub, + Mul, + MulHi, + Xor, + Or, + Rotl, + Rotr, + Mad, + Shfl, + Load, + /// Warp-coalesced load (lever b). Never emitted unless `wide_frac > 0`. + WLoad, +} + +impl Op { + /// The name used in program.json, kernel comments and the op mix. + pub fn name(self) -> &'static str { + match self { + Op::Add => "add", + Op::Sub => "sub", + Op::Mul => "mul", + Op::MulHi => "mulhi", + Op::Xor => "xor", + Op::Or => "or", + Op::Rotl => "rotl", + Op::Rotr => "rotr", + Op::Mad => "mad", + Op::Shfl => "shfl", + Op::Load => "load", + Op::WLoad => "wload", + } + } + + pub fn from_name(s: &str) -> Option { + Some(match s { + "add" => Op::Add, + "sub" => Op::Sub, + "mul" => Op::Mul, + "mulhi" => Op::MulHi, + "xor" => Op::Xor, + "or" => Op::Or, + "rotl" => Op::Rotl, + "rotr" => Op::Rotr, + "mad" => Op::Mad, + "shfl" => Op::Shfl, + "load" => Op::Load, + "wload" => Op::WLoad, + _ => return None, + }) + } +} + +/// One instruction. Every field is drawn for every instruction whether the op uses it or not, so the +/// draw stream is identical for every op. +#[derive(Clone, Copy, Debug, PartialEq, Eq)] +pub struct Instr { + pub op: Op, + /// Destination register 0..7. + pub dst: u8, + /// Source register 0..7, never equal to `dst`. + pub src: u8, + /// Second source (mad only). + pub src2: u8, + /// Add immediate A. + pub imm: u32, + /// Add immediate B. + pub imm2: u32, + /// rotl amount 1..31. + pub rot: u32, + /// Selector bit of r0 for add, 0..31. + pub bit: u8, + /// Shuffle xor mask: 1, 2, 4, 8 or 16. + pub mask: u8, +} + +#[derive(Clone, Debug, PartialEq, Eq)] +pub struct Program { + pub seed_string: String, + pub seed: [u32; 8], + pub instrs: Vec, +} + +impl Program { + pub fn loads_per_hash(&self) -> usize { + self.instrs.iter().filter(|i| i.op == Op::Load || i.op == Op::WLoad).count() * ITERATIONS + } + pub fn wide_loads_per_hash(&self) -> usize { + self.instrs.iter().filter(|i| i.op == Op::WLoad).count() * ITERATIONS + } + pub fn has_wide(&self) -> bool { + self.instrs.iter().any(|i| i.op == Op::WLoad) + } + /// Distinct dataset items a 32-lane warp touches per hash: 32 per plain load, 2 per wide load. + pub fn items_per_warp(&self) -> usize { + (self.loads_per_hash() - self.wide_loads_per_hash()) * 32 + self.wide_loads_per_hash() * 2 + } + /// Op histogram, count descending then name ascending, as the Swift prints it. + pub fn histogram(&self) -> Vec<(&'static str, usize)> { + let mut counts: Vec<(&'static str, usize)> = Vec::new(); + for i in &self.instrs { + let name = i.op.name(); + match counts.iter_mut().find(|(n, _)| *n == name) { + Some(e) => e.1 += 1, + None => counts.push((name, 1)), + } + } + counts.sort_by(|a, b| b.1.cmp(&a.1).then_with(|| a.0.cmp(b.0))); + counts + } + /// "load=13 xor=13 ..." as written into program.h. + pub fn op_mix(&self) -> String { + self.histogram().iter().map(|(n, c)| format!("{n}={c}")).collect::>().join(" ") + } +} + +/// Weights sum to 100. Loads are 25 percent so the kernel leans on memory. +pub const OP_WEIGHTS: [(Op, u64); 11] = [ + (Op::Load, 25), + (Op::Add, 12), + (Op::Xor, 10), + (Op::Mul, 8), + (Op::Mad, 8), + (Op::Shfl, 8), + (Op::Rotl, 7), + (Op::Sub, 6), + (Op::MulHi, 6), + (Op::Rotr, 6), + (Op::Or, 4), +]; + +/// Generator levers (MEMHARD.md section 2.4). The defaults reproduce the original generator exactly. +#[derive(Clone, Copy, Debug, PartialEq, Eq)] +pub struct GeneratorConfig { + /// Percent weight of the load op. + pub load_weight: u64, + /// Percent of load instructions emitted as warp-coalesced wide loads. + pub wide_frac: u64, +} + +impl Default for GeneratorConfig { + fn default() -> Self { + Self { load_weight: 25, wide_frac: 0 } + } +} + +impl GeneratorConfig { + /// Scaled weights: load gets `load_weight`, the other ten ops share the rest in their original + /// proportions, rounded by largest remainder so the table still sums to 100. + pub fn weights(&self) -> Vec<(Op, u64)> { + if self.load_weight == 25 { + return OP_WEIGHTS.to_vec(); + } + let others = &OP_WEIGHTS[1..]; + let total: u64 = others.iter().map(|w| w.1).sum(); // 75 + let budget = 100 - self.load_weight; + let mut scaled: Vec<(Op, u64, u64)> = + others.iter().map(|&(op, w)| (op, (w * budget) / total, (w * budget) % total)).collect(); + let mut sum: u64 = scaled.iter().map(|s| s.1).sum(); + let mut order: Vec = (0..scaled.len()).collect(); + order.sort_by(|&a, &b| scaled[b].2.cmp(&scaled[a].2).then_with(|| a.cmp(&b))); + let mut k = 0; + while sum < budget { + scaled[order[k]].1 += 1; + sum += 1; + k += 1; + } + let mut out = vec![(Op::Load, self.load_weight)]; + out.extend(scaled.iter().map(|s| (s.0, s.1))); + out + } +} + +/// The default generator for a seed string. +pub fn generate(seed_string: &str) -> Program { + generate_with(seed_string, &GeneratorConfig::default()) +} + +/// The generator with levers. `generateProgram` in the Swift, draw for draw. +pub fn generate_with(seed_string: &str, cfg: &GeneratorConfig) -> Program { + let seed = seed_words(seed_string); + generate_from_words(seed_string, seed, cfg) +} + +/// The generator from already-derived seed words (what the chain will call once the VDF output is in). +pub fn generate_from_words(seed_string: &str, seed: [u32; 8], cfg: &GeneratorConfig) -> Program { + let mut rng = program_rng(&seed); + let weights = cfg.weights(); + let mut instrs = Vec::with_capacity(INSTR_COUNT); + for _ in 0..INSTR_COUNT { + let mut roll = rng.below(100); + let mut op = Op::Add; + for &(o, w) in &weights { + if roll < w { + op = o; + break; + } + roll -= w; + } + let dst = rng.below(8); + let mut a = rng.below(7); + if a >= dst { + a += 1; + } + let b = rng.below(8); + let imm = rng.next() as u32; + let imm2 = rng.next() as u32; + let rot = 1 + rng.below(31) as u32; + let bit = rng.below(32); + let mask = 1u8 << rng.below(5); + // Lever (b): the already-drawn selector bit decides whether a load is wide, so the stream is unchanged. + if op == Op::Load && bit * 100 < cfg.wide_frac * 32 { + op = Op::WLoad; + } + instrs.push(Instr { op, dst: dst as u8, src: a as u8, src2: b as u8, imm, imm2, rot, bit: bit as u8, mask }); + } + Program { seed_string: seed_string.to_string(), seed, instrs } +} + +#[cfg(test)] +mod tests { + use super::*; + + #[test] + fn default_weights_unchanged() { + assert_eq!(GeneratorConfig::default().weights(), OP_WEIGHTS.to_vec()); + } + + #[test] + fn load_weight_17_table() { + // MEMHARD.md section 2.4: add=13 xor=11 mul=9 mad=9 shfl=9 rotl=8 sub=7 mulhi=7 rotr=6 or=4. + let w = GeneratorConfig { load_weight: 17, wide_frac: 0 }.weights(); + let expect = [ + (Op::Load, 17), + (Op::Add, 13), + (Op::Xor, 11), + (Op::Mul, 9), + (Op::Mad, 9), + (Op::Shfl, 9), + (Op::Rotl, 8), + (Op::Sub, 7), + (Op::MulHi, 7), + (Op::Rotr, 6), + (Op::Or, 4), + ]; + assert_eq!(w, expect.to_vec()); + assert_eq!(w.iter().map(|x| x.1).sum::(), 100); + } + + #[test] + fn genesis_shape() { + let p = generate("igneum-genesis"); + assert_eq!(p.instrs.len(), 64); + assert_eq!(p.loads_per_hash(), 104); + assert_eq!(p.op_mix(), "load=13 xor=13 sub=7 shfl=6 add=5 mulhi=5 mad=4 rotr=4 mul=3 rotl=3 or=1"); + for i in &p.instrs { + assert_ne!(i.dst, i.src); + assert!((1..=31).contains(&i.rot)); + assert!(i.mask.is_power_of_two() && i.mask <= 16); + } + } +} diff --git a/igneum-pow/src/lib.rs b/igneum-pow/src/lib.rs new file mode 100644 index 000000000..f730605a0 --- /dev/null +++ b/igneum-pow/src/lib.rs @@ -0,0 +1,30 @@ +//! igneum-pow: the Igneum lottery hash, bit-exact with the Swift prototype in `proto-metal/main.swift`. +//! +//! The crate has five parts, each mirroring one section of the prototype: +//! +//! * [`seed`]: the 32-byte seed words from a string (FNV-1a 64, four salts) and the SplitMix64 stream. +//! * [`generator`]: the 64-instruction program drawn from a seed. +//! * [`memhard`]: the 256 MiB ChaCha12 cache and the 8-round dataset item derivation (`proto-metal/MEMHARD.md`). +//! * [`verify`]: the 32-lane warp interpreter that computes the 64-bit hash on the CPU, deriving dataset +//! words on demand from the cache (or from the closed form, for the old packs). +//! * [`emit`]: the Metal, CUDA and OpenCL kernel text for a program, byte-identical to the Swift exporter. +//! +//! Nothing here depends on a crate outside the standard library. The integration points for the +//! rusty-kaspa fork (`docs/fork-map.md`) are [`verify::Epoch`], [`verify::Epoch::verify_block`] and +//! [`verify::Epoch::hash_warp`]; the node hands [`emit::Pack`] files to miners. + +// The lane loops are written index style on purpose so they read like the kernels they mirror, and the +// SplitMix64 `next` keeps the Swift name. +#![allow(clippy::needless_range_loop, clippy::should_implement_trait, clippy::large_enum_variant)] +#![allow(clippy::manual_slice_size_calculation, clippy::too_many_arguments)] + +pub mod emit; +pub mod generator; +pub mod memhard; +pub mod seed; +pub mod verify; + +pub use generator::{generate, Instr, Op, Program}; +pub use memhard::{Cache, MemhardCpu, MixParams}; +pub use seed::{fnv1a64, seed_words, SplitMix64}; +pub use verify::{hash_warp, verify_block, DatasetMode, DatasetSource, Epoch}; diff --git a/igneum-pow/src/main.rs b/igneum-pow/src/main.rs new file mode 100644 index 000000000..dbf32d4f5 --- /dev/null +++ b/igneum-pow/src/main.rs @@ -0,0 +1,156 @@ +//! igneum-pow CLI. +//! +//! igneum-pow bench --seed [--day ] [--closed-form] [--dataset-log2 28] [--warps 20] +//! igneum-pow export --seed --out [--day ] [--closed-form] [--dataset-log2 28] +//! igneum-pow hash --seed --nonce [--day ] [--closed-form] [--dataset-log2 28] + +use igneum_pow::emit::export_pack; +use igneum_pow::memhard::Cache; +use igneum_pow::seed::day_key; +use igneum_pow::verify::{DatasetMode, Epoch, DEFAULT_DATASET_LOG2}; +use std::time::Instant; + +struct Args { + cmd: String, + seed: String, + day: String, + out: Option, + closed_form: bool, + dataset_log2: u32, + warps: usize, + nonce: u32, +} + +fn usage() -> ! { + eprintln!( + "igneum-pow --seed [--day 2026-10-03] [--closed-form] [--dataset-log2 28]\n\ + \x20 bench [--warps 20] fill the cache, then time the CPU verifier per 32-lane warp\n\ + \x20 export --out write the program pack (kernel.cu, kernel.cl, program.metal, memhard.h, ...)\n\ + \x20 hash --nonce print the 64-bit hash of one nonce" + ); + std::process::exit(2) +} + +fn parse() -> Args { + let mut a = Args { + cmd: String::new(), + seed: "igneum-genesis".into(), + day: "2026-10-03".into(), + out: None, + closed_form: false, + dataset_log2: DEFAULT_DATASET_LOG2, + warps: 20, + nonce: 0, + }; + let mut it = std::env::args().skip(1); + a.cmd = it.next().unwrap_or_else(|| usage()); + while let Some(k) = it.next() { + let mut val = || it.next().unwrap_or_else(|| usage()); + match k.as_str() { + "--seed" => a.seed = val(), + "--day" => a.day = val(), + "--out" => a.out = Some(val()), + "--closed-form" => a.closed_form = true, + "--dataset-log2" => a.dataset_log2 = val().parse().unwrap_or_else(|_| usage()), + "--warps" => a.warps = val().parse().unwrap_or_else(|_| usage()), + "--nonce" => a.nonce = val().parse().unwrap_or_else(|_| usage()), + _ => usage(), + } + } + a +} + +fn main() { + let a = parse(); + let mode = if a.closed_form { DatasetMode::ClosedForm } else { DatasetMode::MemoryHard }; + match a.cmd.as_str() { + "bench" => bench(&a, mode), + "export" => export(&a, mode), + "hash" => { + let e = Epoch::new(&a.seed, &a.day, mode, a.dataset_log2); + println!("{:016x}", e.hash(a.nonce)); + } + _ => usage(), + } +} + +fn bench(a: &Args, mode: DatasetMode) { + println!( + "igneum-pow bench: seed \"{}\", day \"{}\", dataset 2^{} words ({})", + a.seed, + a.day, + a.dataset_log2, + mode.name() + ); + if mode == DatasetMode::MemoryHard { + // Time the cache fill on its own first (one core), then build the epoch (which fills it again). + let t0 = Instant::now(); + let c = Cache::fill(day_key(&a.day)); + let fill_ms = t0.elapsed().as_secs_f64() * 1e3; + println!("cache: fill {fill_ms:.1} ms on one core (2^26 words, 65536 chains of 64 ChaCha12 blocks), FNV-1a 64 {:016x}", c.fnv1a64()); + drop(c); + } + let t0 = Instant::now(); + let e = Epoch::new(&a.seed, &a.day, mode, a.dataset_log2); + let build_ms = t0.elapsed().as_secs_f64() * 1e3; + println!( + "program: {} loads/hash, {} items/warp, op mix {}; epoch built in {build_ms:.1} ms", + e.program.loads_per_hash(), + e.program.items_per_warp(), + e.program.op_mix() + ); + let bases = [0u32, 4096, 1_000_000]; + for &b in &bases { + let t = Instant::now(); + let r = e.interpret_warp(b); + let ms = t.elapsed().as_secs_f64() * 1e3; + println!( + "warp base {b}: single cold run {ms:.3} ms, {} items derived, lane0 {:016x} lane31 {:016x}", + r.items_derived, r.hashes[0], r.hashes[31] + ); + } + let n = a.warps.max(1); + let t = Instant::now(); + let mut sink = 0u64; + for i in 0..n { + let w = e.hash_warp((i as u32) * 32 + 65536); + sink ^= w[0]; + } + let avg = t.elapsed().as_secs_f64() * 1e3 / n as f64; + println!("CPU verify: {avg:.3} ms per 32-lane warp, avg of {n} (checksum {sink:016x})"); +} + +fn export(a: &Args, mode: DatasetMode) { + let out = a.out.clone().unwrap_or_else(|| usage()); + let t0 = Instant::now(); + let e = Epoch::new(&a.seed, &a.day, mode, a.dataset_log2); + let build_ms = t0.elapsed().as_secs_f64() * 1e3; + println!("igneum-pow export {out}"); + println!( + "seed \"{}\", day \"{}\", dataset 2^{} words ({}), loads/hash {}, wide loads/hash {}; epoch built in {build_ms:.1} ms", + a.seed, + a.day, + a.dataset_log2, + mode.name(), + e.program.loads_per_hash(), + e.program.wide_loads_per_hash() + ); + println!("op mix: {}", e.program.op_mix()); + let source = format!("igneum-pow (Rust) CPU interpreter, {} dataset", mode.name()); + let pack = export_pack(&e, &a.day, &source); + let dir = std::path::Path::new(&out); + if let Err(err) = pack.write_to(dir) { + eprintln!("FAIL: write error {err}"); + std::process::exit(1); + } + for (name, text) in &pack.files { + println!("wrote {}/{name} ({} bytes)", dir.display(), text.len()); + } + for (i, b) in pack.bases.iter().enumerate() { + println!("vector warp base {b}: lane0 {:016x} lane31 {:016x}", pack.outs[i][0], pack.outs[i][31]); + } + if mode == DatasetMode::MemoryHard { + println!("cache FNV-1a 64 {:016x}", pack.vectors.cache_fnv); + } + println!("OVERALL: PASS (pack written)"); +} diff --git a/igneum-pow/src/memhard.rs b/igneum-pow/src/memhard.rs new file mode 100644 index 000000000..3141f7d04 --- /dev/null +++ b/igneum-pow/src/memhard.rs @@ -0,0 +1,313 @@ +//! The memory-hard dataset of `proto-metal/MEMHARD.md`: a 256 MiB cache of chained ChaCha12 blocks keyed +//! by the day key, and 64-byte dataset items derived by 8 dependent cache reads through a seed-parameterised +//! ARX-multiply mixer. The verifier holds the cache and never the dataset. +//! +//! All arithmetic is on u32 modulo 2^32. Rotations are by 1..31 at every call site. + +use crate::seed::{day_key, fnv1a64_words, SplitMix64}; + +pub const CACHE_LOG2_WORDS: usize = 26; +pub const CACHE_SEGMENT_LOG2_LINES: usize = 6; +/// 2^26 words = 256 MiB. +pub const CACHE_WORDS: usize = 1 << CACHE_LOG2_WORDS; +/// 2^22 lines of 16 words. +pub const CACHE_LINES: usize = CACHE_WORDS >> 4; +/// 64 chained lines per segment. +pub const CACHE_LINES_PER_SEGMENT: usize = 1 << CACHE_SEGMENT_LOG2_LINES; +/// 2^16 independent segments. +pub const CACHE_SEGMENTS: usize = CACHE_LINES >> CACHE_SEGMENT_LOG2_LINES; +pub const CACHE_LINE_MASK: u32 = (CACHE_LINES - 1) as u32; +pub const ITEM_ROUNDS: usize = 8; +pub const CHACHA_ROUNDS: usize = 12; +/// The ChaCha constants "expand 32-byte k". +pub const CHACHA_SIGMA: [u32; 4] = [0x61707865, 0x3320646e, 0x79622d32, 0x6b206574]; +/// "Igne", "umMH". +pub const CACHE_TAG: [u32; 2] = [0x49676e65, 0x756d4d48]; + +#[inline(always)] +fn rotl(x: u32, n: u32) -> u32 { + x.rotate_left(n) +} + +/// The ChaCha quarter round with explicit rotations. +#[inline(always)] +fn qr(s: &mut [u32; 16], a: usize, b: usize, c: usize, d: usize, r1: u32, r2: u32, r3: u32, r4: u32) { + s[a] = s[a].wrapping_add(s[b]); + s[d] ^= s[a]; + s[d] = rotl(s[d], r1); + s[c] = s[c].wrapping_add(s[d]); + s[b] ^= s[c]; + s[b] = rotl(s[b], r2); + s[a] = s[a].wrapping_add(s[b]); + s[d] ^= s[a]; + s[d] = rotl(s[d], r3); + s[c] = s[c].wrapping_add(s[d]); + s[b] ^= s[c]; + s[b] = rotl(s[b], r4); +} + +/// `y = ChaCha12 core(x) + x`. Standard rotations 16, 12, 8, 7; column round then diagonal round, six times. +#[inline] +pub fn chacha_block(x: &[u32; 16]) -> [u32; 16] { + let mut y = *x; + for _ in 0..CHACHA_ROUNDS / 2 { + qr(&mut y, 0, 4, 8, 12, 16, 12, 8, 7); + qr(&mut y, 1, 5, 9, 13, 16, 12, 8, 7); + qr(&mut y, 2, 6, 10, 14, 16, 12, 8, 7); + qr(&mut y, 3, 7, 11, 15, 16, 12, 8, 7); + qr(&mut y, 0, 5, 10, 15, 16, 12, 8, 7); + qr(&mut y, 1, 6, 11, 12, 16, 12, 8, 7); + qr(&mut y, 2, 7, 8, 13, 16, 12, 8, 7); + qr(&mut y, 3, 4, 9, 14, 16, 12, 8, 7); + } + for i in 0..16 { + y[i] = y[i].wrapping_add(x[i]); + } + y +} + +/// Mixer parameters drawn from the day key. Draw order: ROT[0..7] (1..31), MUL[0..15] (odd), RC[0..15]. +#[derive(Clone, Debug, PartialEq, Eq)] +pub struct MixParams { + pub key: [u32; 8], + pub rot: [u32; 8], + pub mul: [u32; 16], + pub rc: [u32; 16], +} + +impl MixParams { + pub fn new(key: [u32; 8]) -> Self { + let mut rng = SplitMix64::new(key[0] as u64 | ((key[1] as u64) << 32)); + let mut rot = [0u32; 8]; + let mut mul = [0u32; 16]; + let mut rc = [0u32; 16]; + for r in rot.iter_mut() { + *r = 1 + rng.below(31) as u32; + } + for m in mul.iter_mut() { + *m = (rng.next() as u32) | 1; + } + for c in rc.iter_mut() { + *c = rng.next() as u32; + } + Self { key, rot, mul, rc } + } + /// Parameters for a day string: the key is `seed_words("day/" + day)`. + pub fn for_day(day: &str) -> Self { + Self::new(day_key(day)) + } +} + +/// Round key `(r + 1) * 0x9E3779B9` mod 2^32. +#[inline(always)] +pub fn round_key(r: usize) -> u32 { + ((r + 1) as u32).wrapping_mul(0x9E3779B9) +} + +/// `M_r` on 16 words in place: per word `(s ^ (RC + rk)) * MUL`, then one ChaCha-shaped double round with +/// the four column rotations `ROT[0..3]` and the four diagonal rotations `ROT[4..7]`. +#[inline(always)] +pub fn mixer(s: &mut [u32; 16], rk: u32, mp: &MixParams) { + for i in 0..16 { + s[i] = (s[i] ^ mp.rc[i].wrapping_add(rk)).wrapping_mul(mp.mul[i]); + } + let r = &mp.rot; + qr(s, 0, 4, 8, 12, r[0], r[1], r[2], r[3]); + qr(s, 1, 5, 9, 13, r[0], r[1], r[2], r[3]); + qr(s, 2, 6, 10, 14, r[0], r[1], r[2], r[3]); + qr(s, 3, 7, 11, 15, r[0], r[1], r[2], r[3]); + qr(s, 0, 5, 10, 15, r[4], r[5], r[6], r[7]); + qr(s, 1, 6, 11, 12, r[4], r[5], r[6], r[7]); + qr(s, 2, 7, 8, 13, r[4], r[5], r[6], r[7]); + qr(s, 3, 4, 9, 14, r[4], r[5], r[6], r[7]); +} + +/// The 256 MiB cache for one day key. +pub struct Cache { + pub key: [u32; 8], + words: Vec, +} + +impl Cache { + /// One segment: 64 chained lines written at `cache[seg * 1024 ..]`. + /// `in_j = prev XOR (sigma || K || seg || j || tag)`, `line_j = B(in_j)`, `prev_0 = 0`. + pub fn fill_segment(words: &mut [u32], seg: usize, key: &[u32; 8]) { + let base = (seg << CACHE_SEGMENT_LOG2_LINES) * 16; + let seg_words = &mut words[base..base + CACHE_LINES_PER_SEGMENT * 16]; + let mut prev = [0u32; 16]; + for (j, line) in seg_words.as_chunks_mut::<16>().0.iter_mut().enumerate() { + let mut x = [0u32; 16]; + x[..4].copy_from_slice(&CHACHA_SIGMA); + x[4..12].copy_from_slice(key); + x[12] = seg as u32; + x[13] = j as u32; + x[14] = CACHE_TAG[0]; + x[15] = CACHE_TAG[1]; + for i in 0..16 { + x[i] ^= prev[i]; + } + let y = chacha_block(&x); + line.copy_from_slice(&y); + prev = y; + } + } + + /// The whole cache on the calling thread: 65,536 chains of 64 ChaCha12 blocks, in segment order. + pub fn fill(key: [u32; 8]) -> Cache { + let mut words = vec![0u32; CACHE_WORDS]; + for seg in 0..CACHE_SEGMENTS { + Self::fill_segment(&mut words, seg, &key); + } + Cache { key, words } + } + + pub fn for_day(day: &str) -> Cache { + Self::fill(day_key(day)) + } + + #[inline(always)] + pub fn words(&self) -> &[u32] { + &self.words + } + + /// Cache line `a` (0 <= a < 2^22) as 16 words. + #[inline(always)] + pub fn line(&self, a: u32) -> &[u32] { + let o = (a & CACHE_LINE_MASK) as usize * 16; + &self.words[o..o + 16] + } + + /// FNV-1a 64 over the cache as little-endian bytes (what `vectors.h` carries as `IGNEUM_CACHE_FNV64`). + pub fn fnv1a64(&self) -> u64 { + fnv1a64_words(&self.words) + } +} + +/// Derive `ts.len()` items into `out`, all chains interleaved round by round so the cache-line misses of +/// independent items overlap in the memory system (`deriveItems` in the Swift). +pub fn derive_items(ts: &[u32], mp: &MixParams, cache: &Cache, out: &mut [[u32; 16]]) { + let n = ts.len(); + debug_assert!(out.len() >= n); + for k in 0..n { + let s = &mut out[k]; + let t = ts[k]; + s[..8].copy_from_slice(&mp.key); + for i in 0..8 { + s[8 + i] = t.wrapping_mul(mp.mul[i]).wrapping_add(mp.rc[i]); + } + } + for r in 0..ITEM_ROUNDS { + let rk = round_key(r); + for s in out[..n].iter_mut() { + mixer(s, rk, mp); + } + for s in out[..n].iter_mut() { + let line = cache.line(s[0]); + for i in 0..16 { + s[i] ^= line[i]; + } + } + } + let rk = round_key(ITEM_ROUNDS); + for s in out[..n].iter_mut() { + mixer(s, rk, mp); + } +} + +/// One dataset item, 16 words. +pub fn derive_item(t: u32, mp: &MixParams, cache: &Cache) -> [u32; 16] { + let mut out = [[0u32; 16]; 1]; + derive_items(&[t], mp, cache, &mut out); + out[0] +} + +/// The CPU verifier's view of the memory-hard dataset: the mixer parameters and the 256 MiB cache. +pub struct MemhardCpu { + pub params: MixParams, + pub cache: Cache, +} + +/// Largest batch `MemhardCpu::fetch` accepts (two warps). +pub const FETCH_MAX: usize = 64; + +impl MemhardCpu { + pub fn new(key: [u32; 8]) -> Self { + Self { params: MixParams::new(key), cache: Cache::fill(key) } + } + pub fn for_day(day: &str) -> Self { + Self::new(day_key(day)) + } + /// `dataset[w] = item(w >> 4)[w & 15]`. + pub fn word(&self, w: u32) -> u32 { + derive_item(w >> 4, &self.params, &self.cache)[(w & 15) as usize] + } + /// `out[k] = dataset[idx[k]]` for every k, `idx.len() <= FETCH_MAX`. Equal items are derived once. + /// Returns the number of distinct items derived. + pub fn fetch(&self, idx: &[u32], out: &mut [u32]) -> usize { + let n = idx.len(); + assert!(n <= FETCH_MAX && out.len() >= n); + let mut uniq = [0u32; FETCH_MAX]; + let mut slot = [0u8; FETCH_MAX]; + let mut u = 0usize; + for k in 0..n { + let t = idx[k] >> 4; + let found = uniq[..u].iter().position(|&x| x == t); + let j = match found { + Some(j) => j, + None => { + uniq[u] = t; + u += 1; + u - 1 + } + }; + slot[k] = j as u8; + } + let mut items = [[0u32; 16]; FETCH_MAX]; + derive_items(&uniq[..u], &self.params, &self.cache, &mut items); + for k in 0..n { + out[k] = items[slot[k] as usize][(idx[k] & 15) as usize]; + } + u + } +} + +#[cfg(test)] +mod tests { + use super::*; + + #[test] + fn mix_params_for_day() { + // MEMHARD.md section 1.4 and the igneum-genesis-mh pack. + let mp = MixParams::for_day("2026-10-03"); + assert_eq!(mp.rot, [20, 20, 19, 4, 26, 3, 3, 27]); + assert_eq!(mp.mul[0], 0x42146205); + assert_eq!(mp.mul[15], 0x99cfb423); + assert_eq!(mp.rc[0], 0xbab68293); + assert_eq!(mp.rc[15], 0x31b49ee2); + assert!(mp.mul.iter().all(|m| m & 1 == 1)); + } + + #[test] + fn chacha_block_is_a_permutation_plus_feedforward() { + let x = [1u32; 16]; + let y = chacha_block(&x); + assert_ne!(x, y); + let z = chacha_block(&x); + assert_eq!(y, z); + } + + #[test] + fn first_cache_line_matches_pack() { + // vectors.json cache_head for day 2026-10-03: segment 0, line 0, with prev = 0. + let key = day_key("2026-10-03"); + let mut words = vec![0u32; CACHE_LINES_PER_SEGMENT * 16]; + Cache::fill_segment(&mut words, 0, &key); + assert_eq!( + &words[..16], + &[ + 0x355a86d2, 0x7957db1c, 0xd21772af, 0x6fc1e09b, 0xd55ce61d, 0x6e6a278b, 0xd3f543ce, 0x223d8e82, + 0x143ab337, 0x2e9f05bd, 0x2eb389bf, 0x0c6e449e, 0x5cfa4222, 0xba6560fe, 0x8e3e1aa4, 0xdbcc1d53 + ] + ); + } +} diff --git a/igneum-pow/src/seed.rs b/igneum-pow/src/seed.rs new file mode 100644 index 000000000..68d40911a --- /dev/null +++ b/igneum-pow/src/seed.rs @@ -0,0 +1,114 @@ +//! Seed derivation and the SplitMix64 stream, exactly as `proto-metal/main.swift` does them. +//! +//! Today a seed is a string ("igneum-genesis", "day/2026-10-03"). On the chain the epoch seed will be the +//! output of a class-group VDF over a certified checkpoint hash. [`seed_words_from_bytes`] is the function +//! boundary for that: whatever bytes the chain settles on go through the same FNV-1a construction, so the +//! generator and the day key never need to know where the bytes came from. + +/// FNV-1a 64 over `bytes` with the standard basis. Used for the cache fingerprint in the packs. +pub fn fnv1a64(bytes: &[u8]) -> u64 { + fnv1a64_with_basis(0xcbf29ce484222325, bytes) +} + +#[inline] +fn fnv1a64_with_basis(basis: u64, bytes: &[u8]) -> u64 { + let mut h = basis; + for &b in bytes { + h ^= b as u64; + h = h.wrapping_mul(0x100000001b3); + } + h +} + +/// FNV-1a 64 over 32-bit words in little-endian byte order (the cache is hashed as raw memory). +pub fn fnv1a64_words(words: &[u32]) -> u64 { + let mut h: u64 = 0xcbf29ce484222325; + for &w in words { + for b in w.to_le_bytes() { + h ^= b as u64; + h = h.wrapping_mul(0x100000001b3); + } + } + h +} + +/// The 32-byte seed (8 x u32) from arbitrary bytes: FNV-1a 64 with four salts, each finalised with the +/// murmur-style mix `h ^= h >> 33; h *= 0xff51afd7ed558ccd; h ^= h >> 33`; low word then high word. +pub fn seed_words_from_bytes(bytes: &[u8]) -> [u32; 8] { + let mut words = [0u32; 8]; + for salt in 0..4u64 { + let basis = 0xcbf29ce484222325u64 ^ salt.wrapping_mul(0x9E3779B97F4A7C15); + let mut h = fnv1a64_with_basis(basis, bytes); + h ^= h >> 33; + h = h.wrapping_mul(0xff51afd7ed558ccd); + h ^= h >> 33; + words[2 * salt as usize] = h as u32; + words[2 * salt as usize + 1] = (h >> 32) as u32; + } + words +} + +/// The 32-byte seed from a string (its UTF-8 bytes). `seedWords` in the Swift. +pub fn seed_words(s: &str) -> [u32; 8] { + seed_words_from_bytes(s.as_bytes()) +} + +/// The day key: the 8 words of `seed_words("day/" + day)`. `K` in MEMHARD.md; `d0, d1` are `K[0], K[1]`. +pub fn day_key(day: &str) -> [u32; 8] { + seed_words(&format!("day/{day}")) +} + +/// SplitMix64, the one deterministic stream every draw in the prototype comes from. +#[derive(Clone, Copy, Debug)] +pub struct SplitMix64 { + pub s: u64, +} + +impl SplitMix64 { + pub fn new(s: u64) -> Self { + Self { s } + } + #[inline] + pub fn next(&mut self) -> u64 { + self.s = self.s.wrapping_add(0x9E3779B97F4A7C15); + let mut z = self.s; + z = (z ^ (z >> 30)).wrapping_mul(0xBF58476D1CE4E5B9); + z = (z ^ (z >> 27)).wrapping_mul(0x94D049BB133111EB); + z ^ (z >> 31) + } + /// `next() % n` as the Swift `below` does it (modulo, not rejection sampling). + #[inline] + pub fn below(&mut self, n: u64) -> u64 { + self.next() % n + } +} + +/// The generator's stream for a seed: `(w0 | w1 << 32) ^ ((w2 | w3 << 32) * 0x9E3779B97F4A7C15)`. +pub fn program_rng(seed: &[u32; 8]) -> SplitMix64 { + let lo = seed[0] as u64 | ((seed[1] as u64) << 32); + let hi = seed[2] as u64 | ((seed[3] as u64) << 32); + SplitMix64::new(lo ^ hi.wrapping_mul(0x9E3779B97F4A7C15)) +} + +#[cfg(test)] +mod tests { + use super::*; + + #[test] + fn genesis_seed_words() { + // From proto-cuda/packs/igneum-genesis/program.json. + assert_eq!( + seed_words("igneum-genesis"), + [0x67a9a7be, 0x1a155b25, 0xfddfb732, 0x4b5af2e8, 0xc55caf33, 0xa27c13b7, 0x06628a48, 0x03852469] + ); + } + + #[test] + fn day_key_2026_10_03() { + // MEMHARD.md section 1.1. + assert_eq!( + day_key("2026-10-03"), + [0x3067619f, 0x3c269176, 0x84a03b03, 0xf8c63294, 0xff977c5b, 0xe60def3e, 0x63630141, 0xb8fbcb58] + ); + } +} diff --git a/igneum-pow/src/verify.rs b/igneum-pow/src/verify.rs new file mode 100644 index 000000000..e3e286add --- /dev/null +++ b/igneum-pow/src/verify.rs @@ -0,0 +1,349 @@ +//! The CPU reference interpreter for one 32-lane warp (`cpuWarpTraced` in the Swift) and the API the node +//! calls. Dataset words come from the memory-hard cache (default) or from the closed form (old packs). + +use crate::generator::{generate, Instr, Op, Program, ITERATIONS, LANES}; +use crate::memhard::MemhardCpu; +use crate::seed::day_key; + +/// Dataset element, closed form of (day words, index). The original prototype's six-operation element. +#[inline(always)] +pub fn dataset_elem(i: u32, d0: u32, d1: u32) -> u32 { + let mut x = i ^ d0; + x = x.wrapping_mul(0x9E3779B1); + x ^= x >> 15; + x = x.wrapping_add(d1); + x = x.wrapping_mul(0x85EBCA77); + x ^= x >> 13; + x = x.wrapping_mul(0xC2B2AE3D); + x ^= x >> 16; + x +} + +/// splitmix32, used for the register init. +#[inline(always)] +pub fn splitmix32(v: u32) -> u32 { + let mut x = v; + x ^= x >> 16; + x = x.wrapping_mul(0x7feb352d); + x ^= x >> 15; + x = x.wrapping_mul(0x846ca68b); + x ^= x >> 16; + x +} + +/// Which construction fills the dataset words. +#[derive(Clone, Copy, Debug, PartialEq, Eq)] +pub enum DatasetMode { + /// The six-operation closed form (packs igneum-genesis and igneum-hourly). Not memory-hard. + ClosedForm, + /// The 256 MiB cache and 8 dependent reads per item (pack igneum-genesis-mh, MEMHARD.md). The default. + MemoryHard, +} + +impl DatasetMode { + pub fn name(self) -> &'static str { + match self { + DatasetMode::ClosedForm => "closed-form", + DatasetMode::MemoryHard => "memory-hard", + } + } +} + +/// Where the interpreter reads dataset words from. +pub enum Dataset { + ClosedForm { d0: u32, d1: u32 }, + MemoryHard(MemhardCpu), +} + +/// A dataset of `2^log2` words plus the construction that fills it. +pub struct DatasetSource { + pub log2_words: u32, + pub mask: u32, + /// The day key `K`; `d0, d1 = K[0], K[1]`. + pub key: [u32; 8], + pub dataset: Dataset, +} + +impl DatasetSource { + /// Build the source for a day. Memory-hard mode fills the 256 MiB cache on the calling thread. + pub fn new(day: &str, mode: DatasetMode, log2_words: u32) -> Self { + Self::from_key(day_key(day), mode, log2_words) + } + + pub fn from_key(key: [u32; 8], mode: DatasetMode, log2_words: u32) -> Self { + assert!((4..=32).contains(&log2_words), "dataset log2 must be in 4..=32"); + 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::new(key)), + }; + Self { log2_words, mask, key, dataset } + } + + pub fn mode(&self) -> DatasetMode { + match self.dataset { + Dataset::ClosedForm { .. } => DatasetMode::ClosedForm, + Dataset::MemoryHard(_) => DatasetMode::MemoryHard, + } + } + + pub fn memhard(&self) -> Option<&MemhardCpu> { + match &self.dataset { + Dataset::MemoryHard(m) => Some(m), + _ => None, + } + } + + /// `dataset[w & mask]`. + pub fn word(&self, w: u32) -> u32 { + let w = w & self.mask; + match &self.dataset { + Dataset::ClosedForm { d0, d1 } => dataset_elem(w, *d0, *d1), + Dataset::MemoryHard(m) => m.word(w), + } + } + + /// `out[k] = dataset[idx[k]]`; indices are already masked. Returns items derived (0 for the closed form). + #[inline] + fn fetch(&self, idx: &[u32; LANES], out: &mut [u32; LANES]) -> usize { + match &self.dataset { + Dataset::ClosedForm { d0, d1 } => { + for k in 0..LANES { + out[k] = dataset_elem(idx[k], *d0, *d1); + } + 0 + } + Dataset::MemoryHard(m) => m.fetch(idx, out), + } + } +} + +/// The result of interpreting one warp. +#[derive(Clone, Copy, Debug, PartialEq, Eq)] +pub struct WarpResult { + pub hashes: [u64; LANES], + /// Distinct dataset items derived from the cache (0 in closed-form mode). + pub items_derived: usize, +} + +#[inline(always)] +fn mulhi32(a: u32, b: u32) -> u32 { + ((a as u64 * b as u64) >> 32) as u32 +} + +/// Interpret `program` for the 32 nonces `base_nonce .. base_nonce + 31` (wrapping). Registers are kept +/// register-major (`r[reg][lane]`) so the per-lane loops vectorise; the semantics are those of the Metal +/// and CUDA kernels instruction for instruction. +pub fn interpret_warp(program: &Program, base_nonce: u32, ds: &DatasetSource) -> WarpResult { + let mask = ds.mask; + let seed = &program.seed; + 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 { + let mut x = nonce ^ seed[i]; + x = x.wrapping_add(0x9e3779b9u32.wrapping_mul(i as u32 + 1)); + x = splitmix32(x); + r[i][lane] = x ^ seed[(i + 1) & 7]; + } + } + let mut items_derived = 0usize; + let mut idx = [0u32; LANES]; + let mut val = [0u32; LANES]; + for _ in 0..ITERATIONS { + let sel = r[0]; + for ins in &program.instrs { + step(ins, &mut r, &sel, mask, ds, &mut idx, &mut val, &mut items_derived); + } + } + let mut hashes = [0u64; LANES]; + for lane in 0..LANES { + let lo = r[0][lane] ^ r[1][lane].rotate_left(7) ^ r[2][lane].rotate_left(14) ^ r[3][lane].rotate_left(21); + let hi = r[4][lane] ^ r[5][lane].rotate_left(9) ^ r[6][lane].rotate_left(18) ^ r[7][lane].rotate_left(27); + hashes[lane] = ((hi as u64) << 32) | lo as u64; + } + WarpResult { hashes, items_derived } +} + +#[inline(always)] +#[allow(clippy::too_many_arguments)] +fn step( + ins: &Instr, + r: &mut [[u32; LANES]; 8], + sel: &[u32; LANES], + mask: u32, + ds: &DatasetSource, + idx: &mut [u32; LANES], + val: &mut [u32; LANES], + items_derived: &mut usize, +) { + let d = ins.dst as usize; + let a = ins.src as usize; + match ins.op { + Op::Add => { + let (imm, imm2, bit) = (ins.imm, ins.imm2, ins.bit as u32); + let src = r[a]; + for lane in 0..LANES { + let s = (sel[lane] >> bit) & 1; + let c = if s != 0 { imm2 } else { imm }; + r[d][lane] = r[d][lane].wrapping_add(src[lane]).wrapping_add(c); + } + } + Op::Sub => { + let src = r[a]; + for lane in 0..LANES { + r[d][lane] = r[d][lane].wrapping_sub(src[lane]); + } + } + Op::Mul => { + let src = r[a]; + for lane in 0..LANES { + r[d][lane] = r[d][lane].wrapping_mul(src[lane]); + } + } + Op::MulHi => { + let src = r[a]; + for lane in 0..LANES { + r[d][lane] = mulhi32(r[d][lane], src[lane]); + } + } + Op::Xor => { + let src = r[a]; + for lane in 0..LANES { + r[d][lane] ^= src[lane]; + } + } + Op::Or => { + let src = r[a]; + for lane in 0..LANES { + r[d][lane] |= src[lane]; + } + } + Op::Rotl => { + let n = ins.rot; + for lane in 0..LANES { + r[d][lane] = r[d][lane].rotate_left(n); + } + } + Op::Rotr => { + let src = r[a]; + for lane in 0..LANES { + r[d][lane] = r[d][lane].rotate_right(src[lane] & 31); + } + } + Op::Mad => { + let src = r[a]; + let src2 = r[ins.src2 as usize]; + for lane in 0..LANES { + r[d][lane] = src[lane].wrapping_mul(src2[lane]).wrapping_add(r[d][lane]); + } + } + Op::Shfl => { + let src = r[a]; + let m = ins.mask as usize; + for lane in 0..LANES { + r[d][lane] ^= src[lane ^ m]; + } + } + Op::Load => { + for lane in 0..LANES { + idx[lane] = r[a][lane] & mask; + } + *items_derived += ds.fetch(idx, val); + for lane in 0..LANES { + r[d][lane] ^= val[lane]; + } + } + Op::WLoad => { + // 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; + } + *items_derived += ds.fetch(idx, val); + for lane in 0..LANES { + r[d][lane] ^= val[lane]; + } + } + } +} + +/// The 32 hashes of one warp (`cpuWarp` in the Swift). +pub fn hash_warp(program: &Program, base_nonce: u32, ds: &DatasetSource) -> [u64; LANES] { + interpret_warp(program, base_nonce, ds).hashes +} + +/// Everything a node needs to verify blocks of one epoch on one day: the program for the epoch seed and the +/// dataset source for the day key. Building one in memory-hard mode fills the 256 MiB cache (about 0.2 s on +/// one core); keep it for the whole epoch and share it between threads (`&Epoch` is `Send + Sync`). +pub struct Epoch { + pub program: Program, + pub dataset: DatasetSource, +} + +/// Default dataset size: 2^28 words = 1 GiB. +pub const DEFAULT_DATASET_LOG2: u32 = 28; + +impl Epoch { + pub fn new(seed: &str, day: &str, mode: DatasetMode, dataset_log2: u32) -> Self { + Self { program: generate(seed), dataset: DatasetSource::new(day, mode, dataset_log2) } + } + + /// The production shape: memory-hard, 1 GiB dataset. + pub fn memory_hard(seed: &str, day: &str) -> Self { + Self::new(seed, day, DatasetMode::MemoryHard, DEFAULT_DATASET_LOG2) + } + + /// The 32 hashes of the warp starting at `base_nonce`. + pub fn hash_warp(&self, base_nonce: u32) -> [u64; LANES] { + hash_warp(&self.program, base_nonce, &self.dataset) + } + + pub fn interpret_warp(&self, base_nonce: u32) -> WarpResult { + interpret_warp(&self.program, base_nonce, &self.dataset) + } + + /// The hash of one nonce. The verification unit is a warp, so the 31 sibling nonces of the aligned + /// 32-nonce group are computed too (the shuffles couple the lanes). + pub fn hash(&self, nonce: u32) -> u64 { + self.hash_warp(nonce & !31)[(nonce & 31) as usize] + } + + /// `hash(nonce) <= target`. The hash is 64 bits; the fork maps it into its 256-bit target space. + pub fn verify_block(&self, nonce: u32, target: u64) -> bool { + self.hash(nonce) <= target + } +} + +/// One-shot `verify_block(seed, nonce, target)`: builds the epoch (cache fill included) and checks. For a +/// node use [`Epoch`] and keep it; this exists for scripts and tests. +pub fn verify_block(seed: &str, day: &str, nonce: u32, target: u64) -> bool { + Epoch::memory_hard(seed, day).verify_block(nonce, target) +} + +#[cfg(test)] +mod tests { + use super::*; + + #[test] + fn closed_form_head_matches_pack() { + // vectors.json dataset_head for igneum-genesis (closed form), day 2026-10-03. + let ds = DatasetSource::new("2026-10-03", DatasetMode::ClosedForm, 28); + assert_eq!(ds.word(0), 0x82174c0f); + assert_eq!(ds.word(1), 0x577bdb9c); + assert_eq!(ds.word(0x0fffffff), 0xf78c84a4); + } + + #[test] + fn closed_form_genesis_vector_lane0() { + let e = Epoch::new("igneum-genesis", "2026-10-03", DatasetMode::ClosedForm, 28); + let w = e.hash_warp(0); + assert_eq!(w[0], 0x2941e93c76cb1910); + assert_eq!(w[31], 0x453388e1be04e25f); + assert_eq!(e.hash(0), w[0]); + assert_eq!(e.hash(31), w[31]); + assert!(e.verify_block(0, u64::MAX)); + assert!(e.verify_block(0, w[0])); + assert!(!e.verify_block(0, w[0] - 1)); + } +} diff --git a/igneum-pow/tests/packs.rs b/igneum-pow/tests/packs.rs new file mode 100644 index 000000000..9d4835dac --- /dev/null +++ b/igneum-pow/tests/packs.rs @@ -0,0 +1,286 @@ +//! Agreement with the Swift prototype through the checked-in packs under proto-cuda/packs/. +//! +//! igneum-genesis-mh: memory-hard dataset (cache FNV, head and last line, 64 sampled words, 96 hashes). +//! igneum-genesis and igneum-hourly: closed-form dataset (head, last, 64 samples, 96 hashes each). +//! All three: program.json instruction by instruction, and every emitted source file byte for byte. + +use igneum_pow::emit::{ + cuda_kernel, cuda_memhard_header, export_pack, metal_memhard, metal_program, opencl_kernel, program_header, + program_json, LoadSource, +}; +use igneum_pow::generator::{generate, Op}; +use igneum_pow::memhard::CACHE_WORDS; +use igneum_pow::verify::{DatasetMode, Epoch}; +use serde_json::Value; +use std::path::PathBuf; +use std::sync::OnceLock; + +const DAY: &str = "2026-10-03"; + +fn packs_dir() -> PathBuf { + PathBuf::from(env!("CARGO_MANIFEST_DIR")).join("../proto-cuda/packs") +} + +fn read(pack: &str, file: &str) -> String { + let p = packs_dir().join(pack).join(file); + std::fs::read_to_string(&p).unwrap_or_else(|e| panic!("read {}: {e}", p.display())) +} + +/// The Swift exporter writes the cache line mask quoted inside the "item" string, which is not valid JSON. +/// Normalise that one defect so the file can be parsed and compared; the Rust emitter writes it bare. +fn fix_swift_item_line(s: &str) -> String { + s.replace("& \"0x003fffff\";", "& 0x003fffff;") +} + +fn json(pack: &str, file: &str) -> Value { + let text = fix_swift_item_line(&read(pack, file)); + serde_json::from_str(&text).unwrap_or_else(|e| panic!("{pack}/{file}: {e}")) +} + +fn hex32(v: &Value) -> u32 { + u32::from_str_radix(v.as_str().unwrap().trim_start_matches("0x"), 16).unwrap() +} +fn hex64(v: &Value) -> u64 { + u64::from_str_radix(v.as_str().unwrap().trim_start_matches("0x"), 16).unwrap() +} + +/// One memory-hard epoch shared by every test (the cache fill is 256 MiB and about 0.2 s). +fn mh_epoch() -> &'static Epoch { + static E: OnceLock = OnceLock::new(); + E.get_or_init(|| Epoch::new("igneum-genesis", DAY, DatasetMode::MemoryHard, 28)) +} + +fn closed_epoch(seed: &str) -> Epoch { + Epoch::new(seed, DAY, DatasetMode::ClosedForm, 28) +} + +fn check_program_json(pack: &str) { + let j = json(pack, "program.json"); + let seed = j["seed"].as_str().unwrap(); + let p = generate(seed); + let sw: Vec = j["seed_words"].as_array().unwrap().iter().map(hex32).collect(); + assert_eq!(p.seed.to_vec(), sw, "{pack}: seed words"); + assert_eq!(p.loads_per_hash() as u64, j["loads_per_hash"].as_u64().unwrap(), "{pack}: loads per hash"); + let instrs = j["instructions"].as_array().unwrap(); + assert_eq!(instrs.len(), p.instrs.len(), "{pack}: instruction count"); + for (k, (ins, ji)) in p.instrs.iter().zip(instrs).enumerate() { + assert_eq!(ji["i"].as_u64().unwrap() as usize, k); + assert_eq!(Op::from_name(ji["op"].as_str().unwrap()).unwrap(), ins.op, "{pack} #{k} op"); + assert_eq!(ji["dst"].as_u64().unwrap(), ins.dst as u64, "{pack} #{k} dst"); + assert_eq!(ji["src"].as_u64().unwrap(), ins.src as u64, "{pack} #{k} src"); + assert_eq!(ji["src2"].as_u64().unwrap(), ins.src2 as u64, "{pack} #{k} src2"); + assert_eq!(hex32(&ji["imm"]), ins.imm, "{pack} #{k} imm"); + assert_eq!(hex32(&ji["imm2"]), ins.imm2, "{pack} #{k} imm2"); + assert_eq!(ji["rot"].as_u64().unwrap(), ins.rot as u64, "{pack} #{k} rot"); + assert_eq!(ji["bit"].as_u64().unwrap(), ins.bit as u64, "{pack} #{k} bit"); + assert_eq!(ji["mask"].as_u64().unwrap(), ins.mask as u64, "{pack} #{k} mask"); + } + // Op mix. + let mix = j["op_mix"].as_object().unwrap(); + for (name, count) in p.histogram() { + assert_eq!(mix[name].as_u64().unwrap() as usize, count, "{pack}: op_mix {name}"); + } + assert_eq!(mix.len(), p.histogram().len()); +} + +#[test] +fn program_json_matches_all_packs() { + for pack in ["igneum-genesis", "igneum-genesis-mh", "igneum-hourly"] { + check_program_json(pack); + } +} + +#[test] +fn mixer_params_match_pack() { + let j = json("igneum-genesis-mh", "program.json"); + let mp = &mh_epoch().dataset.memhard().unwrap().params; + let key: Vec = j["dataset"]["key"].as_array().unwrap().iter().map(hex32).collect(); + assert_eq!(mp.key.to_vec(), key); + let rot: Vec = + j["dataset"]["mixer"]["rot"].as_array().unwrap().iter().map(|v| v.as_u64().unwrap() as u32).collect(); + assert_eq!(mp.rot.to_vec(), rot); + let mul: Vec = j["dataset"]["mixer"]["mul"].as_array().unwrap().iter().map(hex32).collect(); + assert_eq!(mp.mul.to_vec(), mul); + let rc: Vec = j["dataset"]["mixer"]["rc"].as_array().unwrap().iter().map(hex32).collect(); + assert_eq!(mp.rc.to_vec(), rc); + assert_eq!(hex32(&j["dataset"]["d0"]), mp.key[0]); + assert_eq!(hex32(&j["dataset"]["d1"]), mp.key[1]); +} + +#[test] +fn cache_matches_vectors() { + let v = json("igneum-genesis-mh", "vectors.json"); + let m = mh_epoch().dataset.memhard().unwrap(); + let w = m.cache.words(); + assert_eq!(w.len(), CACHE_WORDS); + let head: Vec = v["cache_head"].as_array().unwrap().iter().map(hex32).collect(); + assert_eq!(&w[..16], &head[..]); + let last: Vec = v["cache_last_line"].as_array().unwrap().iter().map(hex32).collect(); + assert_eq!(&w[CACHE_WORDS - 16..], &last[..]); + assert_eq!(m.cache.fnv1a64(), 0x48c4f5bf24166b2e, "cache FNV-1a 64 (MEMHARD.md)"); + assert_eq!(m.cache.fnv1a64(), hex64(&v["cache_fnv1a64"])); +} + +fn check_dataset_words(pack: &str, e: &Epoch) { + let v = json(pack, "vectors.json"); + let ds = &e.dataset; + let head: Vec = v["dataset_head"].as_array().unwrap().iter().map(hex32).collect(); + for (i, h) in head.iter().enumerate() { + assert_eq!(ds.word(i as u32), *h, "{pack}: dataset[{i}]"); + } + let last_index = v["dataset_last_index"].as_u64().unwrap() as u32; + assert_eq!(last_index, ds.mask); + assert_eq!(ds.word(last_index), hex32(&v["dataset_last"]), "{pack}: dataset[MASK]"); + let samples = v["dataset_samples"].as_array().unwrap(); + assert_eq!(samples.len(), 64); + let mut n = 0; + for s in samples { + let idx = s["index"].as_u64().unwrap() as u32; + assert_eq!(ds.word(idx), hex32(&s["value"]), "{pack}: dataset[{idx}]"); + n += 1; + } + assert_eq!(n, 64); +} + +#[test] +fn dataset_words_match_memhard_pack() { + check_dataset_words("igneum-genesis-mh", mh_epoch()); +} + +#[test] +fn dataset_words_match_closed_packs() { + check_dataset_words("igneum-genesis", &closed_epoch("igneum-genesis")); + check_dataset_words("igneum-hourly", &closed_epoch("igneum-hourly")); +} + +/// Returns the number of hashes compared (3 warps x 32 lanes = 96). +fn check_vectors(pack: &str, e: &Epoch) -> usize { + let v = json(pack, "vectors.json"); + assert_eq!(v["dataset_mode"].as_str().unwrap(), e.dataset.mode().name()); + assert_eq!(v["dataset_log2_words"].as_u64().unwrap() as u32, e.dataset.log2_words); + let warps = v["warps"].as_array().unwrap(); + assert_eq!(warps.len(), 3); + let mut n = 0; + for w in warps { + let base = w["base_nonce"].as_u64().unwrap() as u32; + let expected: Vec = w["expected"].as_array().unwrap().iter().map(hex64).collect(); + let got = e.hash_warp(base); + for lane in 0..32 { + assert_eq!(got[lane], expected[lane], "{pack}: base {base} lane {lane}"); + n += 1; + } + // The single-nonce API agrees with the warp. + assert_eq!(e.hash(base + 7), expected[7]); + assert!(e.verify_block(base + 7, expected[7])); + assert!(!e.verify_block(base + 7, expected[7] - 1)); + } + n +} + +#[test] +fn vectors_memhard_96() { + assert_eq!(check_vectors("igneum-genesis-mh", mh_epoch()), 96); +} + +#[test] +fn vectors_closed_form_genesis_96() { + assert_eq!(check_vectors("igneum-genesis", &closed_epoch("igneum-genesis")), 96); +} + +#[test] +fn vectors_closed_form_hourly_96() { + assert_eq!(check_vectors("igneum-hourly", &closed_epoch("igneum-hourly")), 96); +} + +fn assert_same_text(pack: &str, file: &str, got: &str) { + let want = read(pack, file); + if got != want { + // Find the first differing line for a readable failure. + let (gl, wl): (Vec<&str>, Vec<&str>) = (got.lines().collect(), want.lines().collect()); + for i in 0..gl.len().max(wl.len()) { + let g = gl.get(i).copied().unwrap_or(""); + let w = wl.get(i).copied().unwrap_or(""); + if g != w { + panic!("{pack}/{file} differs at line {}:\n pack: {w}\n rust: {g}", i + 1); + } + } + panic!("{pack}/{file} differs only in trailing bytes (len {} vs {})", got.len(), want.len()); + } +} + +fn check_sources(pack: &str, e: &Epoch) { + let p = &e.program; + let mp = e.dataset.memhard().map(|m| &m.params); + assert_same_text(pack, "kernel.cu", &cuda_kernel(p, mp)); + assert_same_text(pack, "program.metal", &metal_program(p, e.dataset.log2_words, LoadSource::Stored)); + assert_same_text(pack, "kernel.cl", &opencl_kernel(p, mp)); + assert_same_text(pack, "program.h", &program_header(p, DAY, &e.dataset.key, e.dataset.log2_words, mp)); + if let Some(mp) = mp { + assert_same_text(pack, "memhard.h", &cuda_memhard_header(p, mp)); + assert_same_text(pack, "memhard.metal", &metal_memhard(mp)); + } + // program.json: byte-identical after normalising the Swift quoting defect in the "item" line. + let want = fix_swift_item_line(&read(pack, "program.json")); + let got = program_json(p, DAY, &e.dataset.key, e.dataset.log2_words, mp); + assert_eq!(got, want, "{pack}/program.json"); + let _: Value = serde_json::from_str(&got).expect("Rust program.json is valid JSON"); +} + +#[test] +fn emitted_sources_match_memhard_pack() { + check_sources("igneum-genesis-mh", mh_epoch()); +} + +#[test] +fn emitted_sources_match_closed_packs() { + check_sources("igneum-genesis", &closed_epoch("igneum-genesis")); + check_sources("igneum-hourly", &closed_epoch("igneum-hourly")); +} + +/// The whole pack as `export` writes it: vectors.json and vectors.h match the Swift ones apart from the +/// provenance string (the Swift adds "Metal GPU cross-check PASS", which the Rust side cannot claim). +fn check_export(pack: &str, e: &Epoch) { + let swift_v = json(pack, "vectors.json"); + let source = swift_v["source"].as_str().unwrap(); + let out = export_pack(e, DAY, source); + let file = |name: &str| -> &str { &out.files.iter().find(|(n, _)| n == name).unwrap().1 }; + assert_same_text(pack, "vectors.json", file("vectors.json")); + assert_same_text(pack, "vectors.h", file("vectors.h")); + let expected: Vec<&str> = if e.dataset.mode() == DatasetMode::MemoryHard { + vec![ + "program.json", + "vectors.json", + "kernel.cu", + "kernel.cl", + "program.h", + "vectors.h", + "program.metal", + "memhard.h", + "memhard.metal", + ] + } else { + vec!["program.json", "vectors.json", "kernel.cu", "kernel.cl", "program.h", "vectors.h", "program.metal"] + }; + assert_eq!(out.files.iter().map(|(n, _)| n.as_str()).collect::>(), expected); +} + +#[test] +fn export_pack_matches_memhard_pack() { + check_export("igneum-genesis-mh", mh_epoch()); +} + +#[test] +fn export_pack_matches_closed_packs() { + check_export("igneum-genesis", &closed_epoch("igneum-genesis")); + check_export("igneum-hourly", &closed_epoch("igneum-hourly")); +} + +#[test] +fn item_is_independent_of_dataset_size() { + // MEMHARD.md 1.7: a smaller dataset is a prefix of items, so dataset[w] is the same at every size. + let big = &mh_epoch().dataset; + let m = big.memhard().unwrap(); + for w in [0u32, 1, 15, 16, 17, 0x00ff_ffff, 0x03ff_ffff] { + assert_eq!(big.word(w), m.word(w)); + } +}