diff --git a/igneum-pow/src/emit.rs b/igneum-pow/src/emit.rs index 2b5299e60..d198a0c29 100644 --- a/igneum-pow/src/emit.rs +++ b/igneum-pow/src/emit.rs @@ -579,6 +579,76 @@ pub fn cuda_kernel_bound(p: &Program, memhard: Option<&MixParams>) -> String { s } +/// The instruction lines of the OpenCL hash kernel body (shared by `igneum_hash` and `igneum_hash_bound`). +fn opencl_instr_lines(p: &Program) -> String { + let mut s = String::with_capacity(6000); + 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 +} + +/// `kernel_bound.cl`: `kernel.cl` plus the header-bound kernel `igneum_hash_bound`, whose init words come from a +/// fifth argument (`__global const uint* initw`, 8 words, `bind::block_init_words`). One source file so the serve +/// mode of proto-opencl/host.c builds cache fill, dataset build and the bound hash from it at runtime. +pub fn opencl_kernel_bound(p: &Program, memhard: Option<&MixParams>) -> String { + let mut s = opencl_kernel(p, memhard); + s.push('\n'); + s.push_str("// Header-bound variant (bind.rs): the init words come from initw, not SEEDW. Same body as igneum_hash.\n"); + s.push_str("IGNEUM_KERNEL_HASH void igneum_hash_bound(__global const uint* ds, __global ulong* out, uint baseNonce, uint mask, __global const uint* initw) {\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(" uint iw0 = initw[0], iw1 = initw[1], iw2 = initw[2], iw3 = initw[3], iw4 = initw[4], iw5 = initw[5], iw6 = initw[6], iw7 = initw[7];\n"); + s.push_str("#if IGNEUM_EXCHANGE == 0\n"); + s.push_str(" IGNEUM_LOCAL_WORDS(xch, 2 * IGNEUM_GROUP);\n"); + 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(&format!( + " {{ uint x = nonce ^ iw{i}; x += 0x9e3779b9u * {}u; x = splitmix32(x); r{i} = x ^ iw{}; }}\n", + i + 1, + (i + 1) & 7 + )); + } + s.push_str(&format!("\n for (uint it = 0u; it < {ITERATIONS}u; ++it) {{\n uint sel = r0;\n")); + s.push_str(&opencl_instr_lines(p)); + 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 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); @@ -685,31 +755,7 @@ pub fn opencl_kernel(p: &Program, memhard: Option<&MixParams>) -> String { 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(&opencl_instr_lines(p)); 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"); @@ -1126,6 +1172,7 @@ pub fn export_pack(epoch: &Epoch, day: &str, source: &str) -> Pack { // Header-bound kernels (3 October 2026, bind.rs): new files, the seven above are unchanged. ("program_bound.metal".to_string(), metal_program_bound(p, ds.log2_words)), ("kernel_bound.cu".to_string(), cuda_kernel_bound(p, memhard)), + ("kernel_bound.cl".to_string(), opencl_kernel_bound(p, memhard)), ]; if let Some(mp) = memhard { files.push(("memhard.h".to_string(), cuda_memhard_header(p, mp))); diff --git a/igneum-pow/src/main.rs b/igneum-pow/src/main.rs index 0c4b5e932..3fc03eee3 100644 --- a/igneum-pow/src/main.rs +++ b/igneum-pow/src/main.rs @@ -1,9 +1,10 @@ //! 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 export --seed --out [--day ] [--closed-form] [--dataset-log2 28] [--epoch-hex <64 hex> --day-hex ] //! igneum-pow hash --seed --nonce [--day ] [--closed-form] [--dataset-log2 28] //! igneum-pow hash-bound --seed --prehash <64 hex> --nonce [--day ] [--closed-form] [--dataset-log2 28] +//! [--epoch-hex <64 hex> --day-hex ] byte seeds instead of strings (Epoch::from_seed_bytes) use igneum_pow::emit::export_pack; use igneum_pow::memhard::Cache; @@ -21,6 +22,8 @@ struct Args { warps: usize, nonce: u64, prehash: String, + epoch_hex: Option, + day_hex: Option, } fn usage() -> ! { @@ -45,6 +48,8 @@ fn parse() -> Args { warps: 20, nonce: 0, prehash: "00".repeat(32), + epoch_hex: None, + day_hex: None, }; let mut it = std::env::args().skip(1); a.cmd = it.next().unwrap_or_else(|| usage()); @@ -59,6 +64,8 @@ fn parse() -> Args { "--warps" => a.warps = val().parse().unwrap_or_else(|_| usage()), "--nonce" => a.nonce = val().parse().unwrap_or_else(|_| usage()), "--prehash" => a.prehash = val(), + "--epoch-hex" => a.epoch_hex = Some(val()), + "--day-hex" => a.day_hex = Some(val()), _ => usage(), } } @@ -78,7 +85,15 @@ fn main() { "hash-bound" => { let bytes = igneum_pow::bind::unhex(&a.prehash).unwrap_or_else(|| usage()); let prehash: [u8; 32] = bytes.as_slice().try_into().unwrap_or_else(|_| usage()); - let e = Epoch::new(&a.seed, &a.day, mode, a.dataset_log2); + // --epoch-hex / --day-hex: the chain's byte seeds (Epoch::from_seed_bytes), as the worker protocol carries them + let e = match (&a.epoch_hex, &a.day_hex) { + (Some(eh), Some(dh)) => { + let eb = igneum_pow::bind::unhex(eh).unwrap_or_else(|| usage()); + let db = igneum_pow::bind::unhex(dh).unwrap_or_else(|| usage()); + Epoch::from_seed_bytes(&eb, &db, "cli") + } + _ => Epoch::new(&a.seed, &a.day, mode, a.dataset_log2), + }; let init = igneum_pow::bind::block_init_words(&prehash, a.nonce); println!( "init words {}", @@ -139,7 +154,14 @@ fn bench(a: &Args, mode: DatasetMode) { 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 e = match (&a.epoch_hex, &a.day_hex) { + (Some(eh), Some(dh)) => { + let eb = igneum_pow::bind::unhex(eh).unwrap_or_else(|| usage()); + let db = igneum_pow::bind::unhex(dh).unwrap_or_else(|| usage()); + Epoch::from_seed_bytes(&eb, &db, &format!("igneum-epoch/{eh}/day/{dh}")) + } + _ => 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!( diff --git a/igneum-pow/tests/packs.rs b/igneum-pow/tests/packs.rs index 7583a1c38..658156306 100644 --- a/igneum-pow/tests/packs.rs +++ b/igneum-pow/tests/packs.rs @@ -257,6 +257,7 @@ fn check_export(pack: &str, e: &Epoch) { "program.metal", "program_bound.metal", "kernel_bound.cu", + "kernel_bound.cl", "memhard.h", "memhard.metal", ] @@ -271,6 +272,7 @@ fn check_export(pack: &str, e: &Epoch) { "program.metal", "program_bound.metal", "kernel_bound.cu", + "kernel_bound.cl", ] }; assert_eq!(out.files.iter().map(|(n, _)| n.as_str()).collect::>(), expected);