From 4b08124b35ffc3c6075ceacedb2b95e227a58193 Mon Sep 17 00:00:00 2001 From: userInner <68625791+userInner@users.noreply.github.com> Date: Wed, 9 Sep 2026 21:24:48 +0800 Subject: [PATCH 1/5] Optimize Apple Goldilocks reduction and hash rejection Reduce Apple u64 kernel overhead by combining reduction carry/borrow corrections, using explicit u32 middle-limb carries, indexing round constants directly, and delaying full result construction until needed. Keep the existing hash, target comparison, and nonce coverage semantics. Add offline arithmetic edge/reference checks, exhaustive small-range coverage checks, and a trusted-kernel CPU parity/ABBA comparison harness. Measured about 7.8% faster than v4.1.0 on one M4 Pro at identical dispatch parameters; this is not a claim about other GPUs or live mining rewards. Validation: stable full-workspace build/tests (58 passed), all-target Clippy with warnings denied, rustfmt, Taplo, GPU reference examples, and an offline CPU benchmark smoke test passed. Independent review of the modular arithmetic and measurements on other Apple GPUs is welcome. Code: crates/engine-gpu/src/kernels/mining_u64_apple.wgsl Validation: crates/engine-gpu/examples/{arithmetic_edges,dispatch_coverage,kernel_compare}.rs --- .../engine-gpu/examples/arithmetic_edges.rs | 135 +++++++++ .../engine-gpu/examples/dispatch_coverage.rs | 42 +++ crates/engine-gpu/examples/kernel_compare.rs | 262 ++++++++++++++++++ .../src/kernels/mining_u64_apple.wgsl | 123 ++++---- 4 files changed, 503 insertions(+), 59 deletions(-) create mode 100644 crates/engine-gpu/examples/arithmetic_edges.rs create mode 100644 crates/engine-gpu/examples/dispatch_coverage.rs create mode 100644 crates/engine-gpu/examples/kernel_compare.rs diff --git a/crates/engine-gpu/examples/arithmetic_edges.rs b/crates/engine-gpu/examples/arithmetic_edges.rs new file mode 100644 index 0000000..77c4c30 --- /dev/null +++ b/crates/engine-gpu/examples/arithmetic_edges.rs @@ -0,0 +1,135 @@ +//! Compare full-width GPU arithmetic against Rust u128, including lazy residues. +use rand::{RngCore, SeedableRng}; +use wgpu::util::DeviceExt; + +fn main() { + tokio::runtime::Runtime::new().unwrap().block_on(run()); +} + +async fn run() { + let instance = wgpu::Instance::default(); + let adapter = instance + .request_adapter(&wgpu::RequestAdapterOptions::default()) + .await + .unwrap(); + let (device, queue) = adapter + .request_device(&wgpu::DeviceDescriptor { + required_features: wgpu::Features::SHADER_INT64, + ..Default::default() + }) + .await + .unwrap(); + let p = 0xffff_ffff_0000_0001u64; + let edges = [0, 1, 0xffff_ffff, 0x1_0000_0000, p - 1, p, p + 1, u64::MAX]; + let mut inputs = Vec::new(); + for a in edges { + for b in edges { + inputs.extend([a, b]); + } + } + let mut rng = rand::rngs::StdRng::seed_from_u64(0x20260909); + for _ in 0..4096 { + inputs.extend([rng.next_u64(), rng.next_u64()]); + } + let count = inputs.len() / 2; + let source = format!( + r#"{} +@group(0) @binding(5) var pairs: array; +@group(0) @binding(6) var answer: array; +@compute @workgroup_size(64) +fn arithmetic_edges(@builtin(global_invocation_id) id: vec3) {{ + if (id.x >= {}u) {{ return; }} + let a = pairs[id.x * 2u]; + let b = pairs[id.x * 2u + 1u]; + let v = mul_wide(a, b); + answer[id.x * 5u] = v.lo; + answer[id.x * 5u + 1u] = v.hi; + answer[id.x * 5u + 2u] = gf64_canon(gf64_mul(a, b)); + answer[id.x * 5u + 3u] = gf64_canon(gf64_sqr(a)); + answer[id.x * 5u + 4u] = gf64_canon(gf64_reduce(U128(a,b))); +}} +"#, + engine_gpu::Kernel::Apple.source(), + count + ); + let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor { + label: None, + source: wgpu::ShaderSource::Wgsl(source.into()), + }); + let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: None, + layout: None, + module: &shader, + entry_point: Some("arithmetic_edges"), + compilation_options: Default::default(), + cache: None, + }); + let input = device.create_buffer_init(&wgpu::util::BufferInitDescriptor { + label: None, + contents: bytemuck::cast_slice(&inputs), + usage: wgpu::BufferUsages::STORAGE, + }); + let size = (count * 5 * 8) as u64; + let output = device.create_buffer(&wgpu::BufferDescriptor { + label: None, + size, + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC, + mapped_at_creation: false, + }); + let staging = device.create_buffer(&wgpu::BufferDescriptor { + label: None, + size, + usage: wgpu::BufferUsages::COPY_DST | wgpu::BufferUsages::MAP_READ, + mapped_at_creation: false, + }); + let group = device.create_bind_group(&wgpu::BindGroupDescriptor { + label: None, + layout: &pipeline.get_bind_group_layout(0), + entries: &[ + wgpu::BindGroupEntry { + binding: 5, + resource: input.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 6, + resource: output.as_entire_binding(), + }, + ], + }); + let mut encoder = device.create_command_encoder(&Default::default()); + { + let mut pass = encoder.begin_compute_pass(&Default::default()); + pass.set_pipeline(&pipeline); + pass.set_bind_group(0, &group, &[]); + pass.dispatch_workgroups((count as u32).div_ceil(64), 1, 1); + } + encoder.copy_buffer_to_buffer(&output, 0, &staging, 0, size); + queue.submit([encoder.finish()]); + let (tx, rx) = std::sync::mpsc::channel(); + staging + .slice(..) + .map_async(wgpu::MapMode::Read, move |r| tx.send(r).unwrap()); + device.poll(wgpu::PollType::wait_indefinitely()).unwrap(); + rx.recv().unwrap().unwrap(); + let mapped = staging.slice(..).get_mapped_range(); + let answers: &[u64] = bytemuck::cast_slice(&mapped); + for i in 0..count { + let a = u128::from(inputs[i * 2]); + let b = u128::from(inputs[i * 2 + 1]); + let product = a * b; + assert_eq!( + &answers[i * 5..i * 5 + 5], + &[ + product as u64, + (product >> 64) as u64, + (product % u128::from(p)) as u64, + ((a * a) % u128::from(p)) as u64, + (((b << 64) | a) % u128::from(p)) as u64 + ], + "pair {i}" + ); + } + drop(mapped); + staging.unmap(); + println!("ARITHMETIC OK: {count} operand pairs, wide product / modular multiply / square / arbitrary u128 reduction"); +} diff --git a/crates/engine-gpu/examples/dispatch_coverage.rs b/crates/engine-gpu/examples/dispatch_coverage.rs new file mode 100644 index 0000000..d71c259 --- /dev/null +++ b/crates/engine-gpu/examples/dispatch_coverage.rs @@ -0,0 +1,42 @@ +//! Check tail batches and high-half nonce carry against an exhaustive CPU search. +use engine_cpu::{AtomicBoolCancelCheck, EngineStatus, MinerEngine, Range}; +use engine_gpu::GpuEngine; +use primitive_types::U512; +use std::sync::atomic::AtomicBool; + +fn main() { + let flag = AtomicBool::new(false); + let cancel = AtomicBoolCancelCheck(&flag); + for batch in [1, 31, 65, 257, 1000] { + let engine = GpuEngine::try_new(batch, 0, false).expect("GPU required"); + for header in [0u8, 17, 255] { + let mut ctx = engine.prepare_context([header; 32], U512::one()); + let start = (U512::one() << 256) - U512::from(100u32); + let end = start + U512::from(512u32); + let (hash, nonce) = (0..513u32) + .map(|offset| { + let nonce = start + U512::from(offset); + (pow_core::hash_from_nonce(&ctx, nonce), nonce) + }) + .min() + .unwrap(); + // Only the CPU's minimum hash qualifies. Missing it exposes gaps + // in dispatch coverage, unlike accepting any easy random solution. + ctx.target = hash + U512::one(); + match engine.search_range(&ctx, Range { start, end }, &cancel) { + EngineStatus::Found { candidate, .. } => { + assert_eq!(candidate.nonce, nonce); + assert_eq!(candidate.hash, hash); + } + other => panic!("batch {batch}: expected CPU minimum, got {other:?}"), + } + ctx.target = U512::zero(); + match engine.search_range(&ctx, Range { start, end }, &cancel) { + EngineStatus::Exhausted { hash_count } => assert_eq!(hash_count, 513), + other => panic!("batch {batch}: expected exhausted range, got {other:?}"), + } + } + GpuEngine::clear_worker_resources(); + } + println!("COVERAGE OK: 15 minimum-hash cases and 15 exhausted ranges"); +} diff --git a/crates/engine-gpu/examples/kernel_compare.rs b/crates/engine-gpu/examples/kernel_compare.rs new file mode 100644 index 0000000..dc3df90 --- /dev/null +++ b/crates/engine-gpu/examples/kernel_compare.rs @@ -0,0 +1,262 @@ +//! Compare two trusted mining kernels with CPU parity checks and ABBA timing. +//! +//! Usage: kernel_compare baseline.wgsl candidate.wgsl +//! Both inputs must implement the miner's buffer layout and bounds checks; +//! runtime shader checks are disabled to match the production mining pipeline. +use pow_core::{hash_from_nonce, mining_midstate, JobContext}; +use primitive_types::U512; +use rand::{RngCore, SeedableRng}; +use std::time::Instant; + +fn u512_to_u32s_le(v: U512) -> [u32; 16] { + let bytes = v.to_little_endian(); + let mut out = [0u32; 16]; + for i in 0..16 { + out[i] = u32::from_le_bytes(bytes[i * 4..(i + 1) * 4].try_into().unwrap()); + } + out +} + +struct Runner { + device: wgpu::Device, + queue: wgpu::Queue, + pipeline: wgpu::ComputePipeline, + results: wgpu::Buffer, + midstate: wgpu::Buffer, + start_nonce: wgpu::Buffer, + target: wgpu::Buffer, + cfg: wgpu::Buffer, + staging: wgpu::Buffer, +} + +impl Runner { + async fn new(trusted: bool, source: &str) -> Self { + let instance = wgpu::Instance::new(&wgpu::InstanceDescriptor { + backends: wgpu::Backends::PRIMARY, + ..Default::default() + }); + let adapter = instance + .request_adapter(&wgpu::RequestAdapterOptions::default()) + .await + .expect("no adapter"); + assert!( + adapter.features().contains(wgpu::Features::SHADER_INT64), + "SHADER_INT64 required" + ); + let (device, queue) = adapter + .request_device(&wgpu::DeviceDescriptor { + label: None, + required_features: wgpu::Features::SHADER_INT64, + ..Default::default() + }) + .await + .unwrap(); + + let kernel = engine_gpu::Kernel::for_adapter(&adapter); + let desc = wgpu::ShaderModuleDescriptor { + label: Some(kernel.label()), + source: wgpu::ShaderSource::Wgsl(source.into()), + }; + let shader = if trusted { + // SAFETY: This offline developer tool accepts only trusted mining + // shaders with the production layout and bounds-checked accesses. + unsafe { + device.create_shader_module_trusted(desc, wgpu::ShaderRuntimeChecks::unchecked()) + } + } else { + device.create_shader_module(desc) + }; + let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: None, + layout: None, + module: &shader, + entry_point: Some("mining_main"), + compilation_options: Default::default(), + cache: None, + }); + let mk = |size: u64, usage| { + device.create_buffer(&wgpu::BufferDescriptor { + label: None, + size, + usage, + mapped_at_creation: false, + }) + }; + use wgpu::BufferUsages as U; + Runner { + pipeline, + results: mk(132, U::STORAGE | U::COPY_SRC | U::COPY_DST), + midstate: mk(96, U::STORAGE | U::COPY_DST), + start_nonce: mk(64, U::STORAGE | U::COPY_DST), + target: mk(64, U::STORAGE | U::COPY_DST), + cfg: mk(12, U::STORAGE | U::COPY_DST), + staging: mk(132, U::MAP_READ | U::COPY_DST), + device, + queue, + } + } + + fn run_batch(&self, header: [u8; 32], start: U512, batch: u32, target: U512) -> f64 { + let nonce_be = start.to_big_endian(); + let mid = mining_midstate(header, nonce_be[..32].try_into().unwrap()); + let mut mid_u32 = [0u32; 24]; + for (i, f) in mid.iter().enumerate() { + mid_u32[2 * i] = *f as u32; + mid_u32[2 * i + 1] = (*f >> 32) as u32; + } + self.queue + .write_buffer(&self.midstate, 0, bytemuck::cast_slice(&mid_u32)); + self.queue.write_buffer( + &self.start_nonce, + 0, + bytemuck::cast_slice(&u512_to_u32s_le(start)), + ); + self.queue.write_buffer( + &self.target, + 0, + bytemuck::cast_slice(&u512_to_u32s_le(target)), + ); + self.queue + .write_buffer(&self.cfg, 0, bytemuck::cast_slice(&[batch, 1u32, batch])); + self.queue.write_buffer(&self.results, 0, &[0u8; 132]); + + let layout = self.pipeline.get_bind_group_layout(0); + let entries: Vec> = [ + (0, &self.results), + (1, &self.midstate), + (2, &self.start_nonce), + (3, &self.target), + (4, &self.cfg), + ] + .iter() + .map(|(i, b)| wgpu::BindGroupEntry { + binding: *i, + resource: b.as_entire_binding(), + }) + .collect(); + let bind_group = self.device.create_bind_group(&wgpu::BindGroupDescriptor { + label: None, + layout: &layout, + entries: &entries, + }); + + let t0 = Instant::now(); + let mut encoder = self + .device + .create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None }); + { + let mut cpass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { + label: None, + timestamp_writes: None, + }); + cpass.set_pipeline(&self.pipeline); + cpass.set_bind_group(0, &bind_group, &[]); + cpass.dispatch_workgroups(batch.div_ceil(256), 1, 1); + } + encoder.copy_buffer_to_buffer(&self.results, 0, &self.staging, 0, 132); + self.queue.submit(Some(encoder.finish())); + + let slice = self.staging.slice(..); + let (tx, rx) = std::sync::mpsc::channel(); + slice.map_async(wgpu::MapMode::Read, move |r| tx.send(r).unwrap()); + loop { + let _ = self.device.poll(wgpu::PollType::Wait { + submission_index: None, + timeout: None, + }); + if rx.try_recv().is_ok() { + break; + } + } + let elapsed = t0.elapsed().as_secs_f64(); + self.staging.unmap(); + elapsed + } + + fn read_results(&self) -> [u32; 33] { + let slice = self.staging.slice(..); + let (tx, rx) = std::sync::mpsc::channel(); + slice.map_async(wgpu::MapMode::Read, move |r| tx.send(r).unwrap()); + let _ = self.device.poll(wgpu::PollType::Wait { + submission_index: None, + timeout: None, + }); + rx.recv().unwrap().unwrap(); + let data = slice.get_mapped_range(); + let mut out = [0u32; 33]; + out.copy_from_slice(bytemuck::cast_slice(&data)); + drop(data); + self.staging.unmap(); + out + } +} + +fn main() { + let batch = 1_000_000u32; + let rt = tokio::runtime::Runtime::new().unwrap(); + let paths: Vec<_> = std::env::args().skip(1).collect(); + assert_eq!(paths.len(), 2, "baseline.wgsl candidate.wgsl"); + let runners: Vec<_> = paths + .iter() + .map(|path| { + let source = std::fs::read_to_string(path).unwrap(); + rt.block_on(Runner::new(true, &source)) + }) + .collect(); + for runner in &runners { + let header = [9u8; 32]; + let ctx = JobContext::new(header, U512::one()); + let start = (U512::from(7u64) << 300) | U512::from(123456789u64); + runner.run_batch(header, start, 256, ctx.target); + let r = runner.read_results(); + assert_eq!(r[0], 1); + let nonce = U512::from_little_endian(bytemuck::cast_slice(&r[1..17])); + let hash = U512::from_little_endian(bytemuck::cast_slice(&r[17..33])); + assert_eq!(hash, hash_from_nonce(&ctx, nonce)); + // Force equality of every high word, then exercise strict comparison + // of the low half. Easy random targets rarely reach this branch. + let mut rng = rand::rngs::StdRng::seed_from_u64(0x517a); + for _ in 0..64 { + let mut header = [0u8; 32]; + let mut bytes = [0u8; 64]; + rng.fill_bytes(&mut header); + rng.fill_bytes(&mut bytes); + let nonce = U512::from_big_endian(&bytes); + let ctx = JobContext::new(header, U512::one()); + let hash = hash_from_nonce(&ctx, nonce); + assert!(hash > U512::zero() && hash < U512::MAX); + for (target, found) in [ + (hash - U512::one(), false), + (hash, false), + (hash + U512::one(), true), + ] { + runner.run_batch(header, nonce, 1, target); + let result = runner.read_results(); + assert_eq!(result[0] != 0, found, "strict target comparison"); + if found { + assert_eq!( + U512::from_little_endian(bytemuck::cast_slice(&result[1..17])), + nonce + ); + assert_eq!( + U512::from_little_endian(bytemuck::cast_slice(&result[17..33])), + hash + ); + } + } + } + runner.run_batch([42; 32], U512::one() << 200, batch, U512::one()); + } + let mut totals = [0.0; 2]; + for iteration in 0..100u64 { + for (offset, index) in [0, 1, 1, 0].into_iter().enumerate() { + let start = (U512::one() << 200) + + U512::from((iteration * 4 + offset as u64) * u64::from(batch)); + totals[index] += runners[index].run_batch([42; 32], start, batch, U512::one()); + } + } + for index in 0..2 { + println!("{}: {:.4} MH/s", paths[index], 200.0 / totals[index]); + } + println!("speedup: {:.3}%", (totals[0] / totals[1] - 1.0) * 100.0); +} diff --git a/crates/engine-gpu/src/kernels/mining_u64_apple.wgsl b/crates/engine-gpu/src/kernels/mining_u64_apple.wgsl index 7feb7cf..8505e65 100644 --- a/crates/engine-gpu/src/kernels/mining_u64_apple.wgsl +++ b/crates/engine-gpu/src/kernels/mining_u64_apple.wgsl @@ -74,12 +74,15 @@ struct U128 { // 2^64 ≡ EPS64 and 2^96 ≡ -1 (mod P). fn gf64_reduce(v: U128) -> u64 { let hi_hi = v.hi >> 32u; - let hi_lo = v.hi & EPS64; - var t0 = v.lo - hi_hi; - t0 = t0 - select(0lu, EPS64, v.lo < hi_hi); - let t1 = hi_lo * EPS64; - let t2 = t0 + t1; - return t2 + select(0lu, EPS64, t2 < t0); + let folded = (v.hi & EPS64) * EPS64; + let sum = v.lo + folded; + let result = sum - hi_hi; + // Combine the addition carry and subtraction borrow before folding. + // Their signed difference is -1, 0 or 1, so only one EPS64 correction + // is needed; the high-limb bounds prevent a further wrap correction. + let correction = i64(select(0i, 1i, sum < v.lo) - select(0i, 1i, sum < hi_hi)); + let bits = bitcast(correction); + return result + ((bits << 32u) - bits); } fn mul_wide(a: u64, b: u64) -> U128 { @@ -91,8 +94,12 @@ fn mul_wide(a: u64, b: u64) -> U128 { let lh = a_lo * b_hi; let hl = a_hi * b_lo; let hh = a_hi * b_hi; - let mid = (ll >> 32u) + (lh & EPS64) + (hl & EPS64); - return U128((mid << 32u) | (ll & EPS64), hh + (lh >> 32u) + (hl >> 32u) + (mid >> 32u)); + // Accumulate the middle limb in 32 bits; keep both carry bits explicitly. + let mid0 = u32(ll >> 32u) + u32(lh); + let c0 = select(0u, 1u, mid0 < u32(lh)); + let mid = mid0 + u32(hl); + let c = c0 + select(0u, 1u, mid < mid0); + return U128((u64(mid) << 32u) | u64(u32(ll)), hh + (lh >> 32u) + (hl >> 32u) + u64(c)); } // (a*b + addend) mod P for b <= 2^64 - 2^32 (all MDS_DIAG entries): the addend's @@ -113,8 +120,12 @@ fn gf64_sqr(a: u64) -> u64 { let ll = a_lo * a_lo; let lh = a_lo * a_hi; let hh = a_hi * a_hi; - let mid = (ll >> 32u) + ((lh & EPS64) << 1u); - return gf64_reduce(U128((mid << 32u) | (ll & EPS64), hh + ((lh >> 32u) << 1u) + (mid >> 32u))); + // Accumulate the middle limb in 32 bits; keep both carry bits explicitly. + let mid0 = u32(ll >> 32u) + u32(lh); + let c0 = select(0u, 1u, mid0 < u32(lh)); + let mid = mid0 + u32(lh); + let c = c0 + select(0u, 1u, mid < mid0); + return gf64_reduce(U128((u64(mid) << 32u) | u64(u32(ll)), hh + ((lh >> 32u) << 1u) + u64(c))); } fn gf64_sbox(x: u64) -> u64 { @@ -146,7 +157,7 @@ fn mds4(x0: u64, x1: u64, x2: u64, x3: u64) -> array { // External linear layer: 4x4 MDS on each chunk, then circulant sums, plus the // next round's constants. Additions are accumulated unreduced (at most 27 // carries) and folded once per output. -fn ext_layer64(state: ptr>, rc: array) { +fn ext_layer64(state: ptr>, rc_index: u32) { var y: array; for (var chunk = 0u; chunk < 3u; chunk++) { let o = chunk * 4u; @@ -158,9 +169,9 @@ fn ext_layer64(state: ptr>, rc: array) { } for (var k = 0u; k < 4u; k++) { let s = acc_add2(acc_add2(y[k], y[k + 4u]), y[k + 8u]); - (*state)[k] = acc_fold(acc_add(acc_add2(y[k], s), rc[k])); - (*state)[k + 4u] = acc_fold(acc_add(acc_add2(y[k + 4u], s), rc[k + 4u])); - (*state)[k + 8u] = acc_fold(acc_add(acc_add2(y[k + 8u], s), rc[k + 8u])); + (*state)[k] = acc_fold(acc_add(acc_add2(y[k], s), RC_EXT[rc_index][k])); + (*state)[k + 4u] = acc_fold(acc_add(acc_add2(y[k + 4u], s), RC_EXT[rc_index][k + 4u])); + (*state)[k + 8u] = acc_fold(acc_add(acc_add2(y[k + 8u], s), RC_EXT[rc_index][k + 8u])); } } @@ -224,7 +235,7 @@ fn permute64(state: ptr>) { sbox_lanes(state, select(1u, 12u, is_ext)); } if (is_ext) { - ext_layer64(state, RC_EXT[select(k, k - 22u, k > 26u)]); + ext_layer64(state, select(k, k - 22u, k > 26u)); } else { int_layer64(state, RC_INTERNAL[k - 4u]); if (k == 26u) { @@ -298,60 +309,54 @@ fn mining_main(@builtin(global_invocation_id) global_id: vec3) { for (var i = 0u; i < 8u; i++) { st[i] = gf64_add(st[i], u64(bswap32(current_nonce[7u - i]))); } - // Squeeze-and-compare phases share one inlined permutation: phase 0 - // pads after absorbing, phase 1 yields the most significant 256 bits of - // the hash, which decide hash-vs-target on their own unless they exactly - // equal the target's high half, and only candidates run phase 2 for the - // low half. Byte-swapped hash words are produced on demand. - var hash_le: array; - var cmp = 0u; - var below = false; - for (var phase = 0u; phase < 3u; phase++) { + // The overwhelmingly common rejection needs only the first hash word. + // Delay materializing the full result until that comparison passes. + for (var phase = 0u; phase < 2u; phase++) { permute64(&st); if (phase == 0u) { st[0] = gf64_add(st[0], 1lu); st[1] = gf64_add(st[1], 1lu); - continue; } - var words: array; - for (var i = 0u; i < 4u; i++) { - let c = gf64_canon(st[i]); - words[2u * i] = bswap32(u32(c & EPS64)); - words[2u * i + 1u] = bswap32(u32(c >> 32u)); + } + let first = bswap32(u32(gf64_canon(st[0]) & EPS64)); + if (first > tgt[15]) { + continue; + } + var hash_le: array; + var cmp = 0u; + for (var i = 0u; i < 4u; i++) { + let c = gf64_canon(st[i]); + hash_le[15u - 2u * i] = bswap32(u32(c & EPS64)); + hash_le[14u - 2u * i] = bswap32(u32(c >> 32u)); + } + for (var i = 0u; i < 8u; i++) { + let h = hash_le[15u - i]; + let t = tgt[15u - i]; + if (h != t) { + cmp = select(2u, 1u, h > t); + break; } - let base = select(15u, 7u, phase == 2u); + } + if (cmp == 1u) { + continue; + } + permute64(&st); + for (var i = 0u; i < 4u; i++) { + let c = gf64_canon(st[i]); + hash_le[7u - 2u * i] = bswap32(u32(c & EPS64)); + hash_le[6u - 2u * i] = bswap32(u32(c >> 32u)); + } + var below = cmp == 2u; + if (!below) { for (var i = 0u; i < 8u; i++) { - hash_le[base - i] = words[i]; - } - if (phase == 1u) { - for (var i = 0u; i < 8u; i++) { - let h = words[i]; - let t = tgt[15u - i]; - if (h != t) { - cmp = select(2u, 1u, h > t); - break; - } - } - if (cmp == 1u) { + let h = hash_le[7u - i]; + let t = tgt[7u - i]; + if (h != t) { + below = h < t; break; } - } else { - below = cmp == 2u; - if (!below) { - for (var i = 0u; i < 8u; i++) { - let h = words[i]; - let t = tgt[7u - i]; - if (h != t) { - below = h < t; - break; - } - } - } } } - if (cmp == 1u) { - continue; - } if (below) { if (atomicExchange(&results[0], 1u) == 0u) { @@ -579,7 +584,7 @@ fn state_unpack(v: ptr>, state: ptr>) { var st: array; state_pack(state, &st); - ext_layer64(&st, RC_ZERO); + ext_layer64(&st, 8u); state_unpack(&st, state); } From 2b1235391bea5358ad1138b54367e888c9f6e60a Mon Sep 17 00:00:00 2001 From: userInner <68625791+userInner@users.noreply.github.com> Date: Wed, 9 Sep 2026 21:36:10 +0800 Subject: [PATCH 2/5] Fold Apple field reduction directly from 32-bit limbs Replace carry/borrow comparisons with signed radix folding and explicit matrix doubling shifts. Bounds guarantee one final EPS correction. Extend arithmetic checks with independent four-limb edge cases and optional candidate source input. Validated full workspace build, 58 tests, all-target Clippy, rustfmt, Taplo, 4416 arithmetic cases, CPU parity and range coverage. Offline M4 Pro comparisons gain 2.3-3.2% over the previous commit and 9.8-10.8% over v4.1.0; no mainnet or cross-device performance claim. Code and validation: crates/engine-gpu/src/kernels/mining_u64_apple.wgsl; crates/engine-gpu/examples/arithmetic_edges.rs. PR: https://github.com/Quantus-Network/quantus-miner/pull/96 --- .../engine-gpu/examples/arithmetic_edges.rs | 19 +++++++++++++-- .../src/kernels/mining_u64_apple.wgsl | 24 +++++++++---------- 2 files changed, 29 insertions(+), 14 deletions(-) diff --git a/crates/engine-gpu/examples/arithmetic_edges.rs b/crates/engine-gpu/examples/arithmetic_edges.rs index 77c4c30..6268458 100644 --- a/crates/engine-gpu/examples/arithmetic_edges.rs +++ b/crates/engine-gpu/examples/arithmetic_edges.rs @@ -1,4 +1,5 @@ //! Compare full-width GPU arithmetic against Rust u128, including lazy residues. +//! Optionally pass a kernel source path to validate an experimental Apple kernel. use rand::{RngCore, SeedableRng}; use wgpu::util::DeviceExt; @@ -27,11 +28,26 @@ async fn run() { inputs.extend([a, b]); } } + // Independently exercise all four limbs of the 128-bit reduction input, + // including negative low sums and high sums crossing the radix boundary. + for w0 in [0u64, 1, 0xffff_fffe, 0xffff_ffff] { + for w1 in [0u64, 1, 0xffff_fffe, 0xffff_ffff] { + for w2 in [0u64, 1, 0xffff_fffe, 0xffff_ffff] { + for w3 in [0u64, 1, 0xffff_fffe, 0xffff_ffff] { + inputs.extend([w0 | (w1 << 32), w2 | (w3 << 32)]); + } + } + } + } let mut rng = rand::rngs::StdRng::seed_from_u64(0x20260909); for _ in 0..4096 { inputs.extend([rng.next_u64(), rng.next_u64()]); } let count = inputs.len() / 2; + let kernel_source = std::env::args() + .nth(1) + .map(|path| std::fs::read_to_string(path).expect("read kernel source")) + .unwrap_or_else(|| engine_gpu::Kernel::Apple.source().to_owned()); let source = format!( r#"{} @group(0) @binding(5) var pairs: array; @@ -49,8 +65,7 @@ fn arithmetic_edges(@builtin(global_invocation_id) id: vec3) {{ answer[id.x * 5u + 4u] = gf64_canon(gf64_reduce(U128(a,b))); }} "#, - engine_gpu::Kernel::Apple.source(), - count + kernel_source, count ); let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor { label: None, diff --git a/crates/engine-gpu/src/kernels/mining_u64_apple.wgsl b/crates/engine-gpu/src/kernels/mining_u64_apple.wgsl index 8505e65..f1f0fd9 100644 --- a/crates/engine-gpu/src/kernels/mining_u64_apple.wgsl +++ b/crates/engine-gpu/src/kernels/mining_u64_apple.wgsl @@ -73,16 +73,16 @@ struct U128 { // Reduce a 128-bit value (lo + hi*2^64) mod P using // 2^64 ≡ EPS64 and 2^96 ≡ -1 (mod P). fn gf64_reduce(v: U128) -> u64 { - let hi_hi = v.hi >> 32u; - let folded = (v.hi & EPS64) * EPS64; - let sum = v.lo + folded; - let result = sum - hi_hi; - // Combine the addition carry and subtraction borrow before folding. - // Their signed difference is -1, 0 or 1, so only one EPS64 correction - // is needed; the high-limb bounds prevent a further wrap correction. - let correction = i64(select(0i, 1i, sum < v.lo) - select(0i, 1i, sum < hi_hi)); - let bits = bitcast(correction); - return result + ((bits << 32u) - bits); + // In radix B=2^32, B^2 = B-1 and B^3 = -1 modulo P. Fold + // the four limbs directly, avoiding 64-bit carry/borrow comparisons. + let low = i64(u32(v.lo)) - i64(u32(v.hi)) - i64(u32(v.hi >> 32u)); + let high = i64(u32(v.lo >> 32u)) + i64(u32(v.hi)) + (low >> 32u); + let result = (u64(u32(high)) << 32u) | u64(u32(low)); + // -1 <= high <= 2B-2: the correction is -1, 0, or 1. + // A positive correction has result <= B^2-B-1; a negative one + // has its high limb equal to B-1. Neither requires another fold. + let correction = bitcast(high >> 32u); + return result + ((correction << 32u) - correction); } fn mul_wide(a: u64, b: u64) -> U128 { @@ -148,9 +148,9 @@ fn mds4(x0: u64, x1: u64, x2: u64, x3: u64) -> array { let t01233 = acc_add(t0123, x3); return array( acc_add2(t01123, t01), - acc_add2(t01123, acc_add(Acc(x2, 0u), x2)), + acc_add2(t01123, Acc(x2 << 1u, u32(x2 >> 63u))), acc_add2(t01233, t23), - acc_add2(t01233, acc_add(Acc(x0, 0u), x0)) + acc_add2(t01233, Acc(x0 << 1u, u32(x0 >> 63u))) ); } From eeb273ee8a59b9d79e7df5b99197a0cf3206844d Mon Sep 17 00:00:00 2001 From: userInner <68625791+userInner@users.noreply.github.com> Date: Wed, 9 Sep 2026 21:45:51 +0800 Subject: [PATCH 3/5] Reuse CPU mining midstates and reject high hashes before final squeeze Add a job-local MiningHasher in pow-core and use it in FastCpuEngine. Refresh cached state on high-half nonce changes, return full hashes for candidates, and preserve strict target comparison, search order, counts and cancellation cadence. Keep hash_from_nonce independent for reference verification. Add cache/boundary/golden-vector/cancellation regressions and an offline CPU reference comparison example. M4 Pro reference comparisons improve 2.57-2.59x; full CLI one-worker ABBA mean improves 2.55x. This affects CPU workers only, not GPU throughput or proven mainnet rewards. Validation: full workspace locked build and 62 tests; all-target Clippy, rustfmt, Taplo; 100 GPU/CPU parity tasks passed. Code: crates/pow-core/src/lib.rs and crates/engine-cpu. PR: https://github.com/Quantus-Network/quantus-miner/pull/96 --- .../engine-cpu/examples/compare_reference.rs | 49 ++++++++ crates/engine-cpu/src/lib.rs | 73 +++++++++++- crates/pow-core/src/lib.rs | 111 ++++++++++++++++++ 3 files changed, 230 insertions(+), 3 deletions(-) create mode 100644 crates/engine-cpu/examples/compare_reference.rs diff --git a/crates/engine-cpu/examples/compare_reference.rs b/crates/engine-cpu/examples/compare_reference.rs new file mode 100644 index 0000000..b7e0473 --- /dev/null +++ b/crates/engine-cpu/examples/compare_reference.rs @@ -0,0 +1,49 @@ +//! Offline end-to-end CPU range comparison against the independent full hash. +use engine_cpu::{AtomicBoolCancelCheck, EngineStatus, FastCpuEngine, MinerEngine, Range}; +use pow_core::{hash_from_nonce, JobContext}; +use primitive_types::U512; +use std::{hint::black_box, sync::atomic::AtomicBool, time::Instant}; + +fn main() { + let engine = FastCpuEngine::new(10_000); + let flag = AtomicBool::new(false); + let cancel = AtomicBoolCancelCheck(&flag); + let count = 20_000u64; + let mut totals = [0.0; 2]; + // Warm both paths, then ABBA with the same count and header per pair. + for round in 0..13u64 { + let ctx = JobContext::new([round as u8; 32], U512::MAX); + let start = U512::one() << 200; + for index in [0, 1, 1, 0] { + let timer = Instant::now(); + if index == 0 { + let mut hashes = 0u64; + for offset in 0..count { + let hash = hash_from_nonce(black_box(&ctx), start + U512::from(offset)); + assert!(black_box(hash) >= ctx.target); + hashes += 1; + } + assert_eq!(hashes, count); + } else { + let result = engine.search_range( + black_box(&ctx), + Range { + start, + end: start + U512::from(count - 1), + }, + &cancel, + ); + assert!( + matches!(result, EngineStatus::Exhausted { hash_count } if hash_count == count) + ); + } + if round > 0 { + totals[index] += timer.elapsed().as_secs_f64(); + } + } + } + let hashes = (count * 24) as f64; + println!("reference: {:.4} MH/s", hashes / totals[0] / 1e6); + println!("CPU engine: {:.4} MH/s", hashes / totals[1] / 1e6); + println!("speedup: {:.3}%", (totals[0] / totals[1] - 1.0) * 100.0); +} diff --git a/crates/engine-cpu/src/lib.rs b/crates/engine-cpu/src/lib.rs index 1c8fc91..c5840e7 100644 --- a/crates/engine-cpu/src/lib.rs +++ b/crates/engine-cpu/src/lib.rs @@ -144,7 +144,7 @@ impl MinerEngine for FastCpuEngine { range: Range, cancel: &dyn CancelCheck, ) -> EngineStatus { - use pow_core::{hash_from_nonce, is_valid_hash, step_nonce}; + use pow_core::{step_nonce, MiningHasher}; if range.start > range.end { return EngineStatus::Exhausted { hash_count: 0 }; @@ -155,6 +155,7 @@ impl MinerEngine for FastCpuEngine { // Use decrementing counter to avoid modulo division in hot loop // Initialize to 0 so we check cancellation immediately on first iteration let mut until_check: u64 = 0; + let mut hasher = MiningHasher::new(ctx); loop { // Check for cancellation every batch_size hashes @@ -166,10 +167,10 @@ impl MinerEngine for FastCpuEngine { } until_check -= 1; - let hash = hash_from_nonce(ctx, current); + let hash = hasher.hash_if_valid(current); hash_count = hash_count.saturating_add(1); - if is_valid_hash(ctx, hash) { + if let Some(hash) = hash { let work = current.to_big_endian(); return EngineStatus::Found { candidate: Candidate { @@ -202,6 +203,72 @@ mod tests { use primitive_types::U512; use std::sync::atomic::AtomicBool; + #[test] + fn cached_search_matches_first_reference_solution_across_nonce_carry() { + let engine = FastCpuEngine::new(7); + let flag = AtomicBool::new(false); + let cancel = AtomicBoolCancelCheck(&flag); + for start in [ + (U512::one() << 256) - U512::from(8u64), + U512::MAX - U512::from(31u64), + ] { + let end = start + U512::from(31u64); + let mut ctx = JobContext::new([37; 32], U512::one()); + let hashes: Vec<_> = (0..32u64) + .map(|i| pow_core::hash_from_nonce(&ctx, start + U512::from(i))) + .collect(); + for target in [ + U512::zero(), + *hashes.iter().min().unwrap() + U512::one(), + U512::MAX, + ] { + ctx.target = target; + let expected = hashes.iter().position(|hash| *hash < target); + match ( + engine.search_range(&ctx, Range { start, end }, &cancel), + expected, + ) { + ( + EngineStatus::Found { + candidate, + hash_count, + .. + }, + Some(index), + ) => { + assert_eq!(candidate.nonce, start + U512::from(index)); + assert_eq!(candidate.hash, hashes[index]); + assert_eq!(candidate.work, candidate.nonce.to_big_endian()); + assert_eq!(hash_count, index as u64 + 1); + } + (EngineStatus::Exhausted { hash_count: 32 }, None) => {} + other => panic!("reference mismatch: {other:?}"), + } + } + } + } + + #[test] + fn cached_search_preserves_cancellation_interval() { + struct CancelAfterFirstBatch(std::sync::atomic::AtomicUsize); + impl CancelCheck for CancelAfterFirstBatch { + fn is_cancelled(&self) -> bool { + self.0.fetch_add(1, Ordering::Relaxed) > 0 + } + } + let cancel = CancelAfterFirstBatch(std::sync::atomic::AtomicUsize::new(0)); + let ctx = JobContext::new([8; 32], U512::MAX); + let result = FastCpuEngine::new(13).search_range( + &ctx, + Range { + start: U512::zero(), + end: U512::from(100u64), + }, + &cancel, + ); + assert!(matches!(result, EngineStatus::Cancelled { hash_count: 13 })); + } + #[test] fn engine_returns_exhausted_when_no_solution_in_range() { let header = [2u8; 32]; diff --git a/crates/pow-core/src/lib.rs b/crates/pow-core/src/lib.rs index 41aaf0f..7954205 100644 --- a/crates/pow-core/src/lib.rs +++ b/crates/pow-core/src/lib.rs @@ -230,6 +230,67 @@ pub fn hash_from_nonce(ctx: &JobContext, nonce: U512) -> U512 { qpow_math::get_nonce_hash(ctx.header, nonce_bytes) } +/// Job-local CPU hasher that reuses the header/high-nonce sponge state. +/// +/// The reference `hash_from_nonce` remains independent. This mining path skips +/// the final squeeze when the first half already exceeds the target, but +/// returns the full reference hash for every qualifying nonce. +pub struct MiningHasher { + header: [u8; 32], + target: U512, + target_bytes: [u8; 64], + poseidon: Poseidon2, + cached: Option<([u8; 32], [Goldilocks; SPONGE_WIDTH])>, +} + +impl MiningHasher { + pub fn new(ctx: &JobContext) -> Self { + Self { + header: ctx.header, + target: ctx.target, + target_bytes: ctx.target.to_big_endian(), + poseidon: Poseidon2::new(), + cached: None, + } + } + + /// Return the full hash iff it is strictly below this job's target. + /// Arbitrary nonce order and carries into the high half are supported. + pub fn hash_if_valid(&mut self, nonce: U512) -> Option { + let bytes = nonce.to_big_endian(); + let high: [u8; 32] = bytes[..32].try_into().unwrap(); + let mut state = match self.cached { + Some((cached_high, state)) if cached_high == high => state, + _ => { + let state = mining_midstate(self.header, high).map(Goldilocks::from_u64); + self.cached = Some((high, state)); + state + } + }; + for (felt, chunk) in state.iter_mut().zip(bytes[32..].chunks_exact(4)) { + *felt += Goldilocks::from_u64(u32::from_le_bytes(chunk.try_into().unwrap()) as u64); + } + self.poseidon.permute_mut(&mut state); + state[0] += Goldilocks::ONE; + state[1] += Goldilocks::ONE; + self.poseidon.permute_mut(&mut state); + + let mut hash_bytes = [0u8; 64]; + for (felt, chunk) in state.iter().zip(hash_bytes[..32].chunks_exact_mut(8)) { + chunk.copy_from_slice(&felt.as_canonical_u64().to_le_bytes()); + } + if hash_bytes[..32] > self.target_bytes[..32] { + return None; + } + self.poseidon.permute_mut(&mut state); + for (felt, chunk) in state.iter().zip(hash_bytes[32..].chunks_exact_mut(8)) { + chunk.copy_from_slice(&felt.as_canonical_u64().to_le_bytes()); + } + let hash = U512::from_big_endian(&hash_bytes); + (hash < self.target).then_some(hash) + } +} + /// Check if hash meets difficulty target pub fn is_valid_hash(ctx: &JobContext, hash: U512) -> bool { hash < ctx.target @@ -253,6 +314,56 @@ pub fn mine_nonce_range(ctx: &JobContext, start_nonce: U512, steps: u64) -> Opti mod tests { use super::*; + #[test] + fn mining_hasher_matches_reference_across_cache_changes() { + let boundary = U512::one() << 256; + let mut nonces = vec![U512::zero(), U512::one(), U512::MAX]; + for offset in 0..8u64 { + nonces.push(boundary - U512::from(4u64) + U512::from(offset)); + } + // Jump backwards and between high halves as well as incrementing. + nonces.extend([U512::one(), U512::MAX, boundary, U512::zero()]); + let mut seed = 0x517a9e37u64; + for _ in 0..64 { + let mut bytes = [0u8; 64]; + for chunk in bytes.chunks_exact_mut(8) { + seed ^= seed << 13; + seed ^= seed >> 7; + seed ^= seed << 17; + chunk.copy_from_slice(&seed.to_le_bytes()); + } + nonces.push(U512::from_big_endian(&bytes)); + } + for header in [0u8, 17, 255] { + for target in [U512::zero(), U512::one(), U512::MAX >> 1, U512::MAX] { + let mut ctx = JobContext::new([header; 32], U512::one()); + ctx.target = target; + let mut hasher = MiningHasher::new(&ctx); + for &nonce in &nonces { + let hash = hash_from_nonce(&ctx, nonce); + assert_eq!(hasher.hash_if_valid(nonce), (hash < target).then_some(hash)); + } + } + } + } + + #[test] + fn mining_hasher_preserves_strict_full_hash_comparison() { + for vector in NONCE_HASH_KVS { + let nonce = U512::from_big_endian(&decode64(vector.nonce)); + let mut ctx = JobContext::new(decode32(vector.header), U512::one()); + let hash = U512::from_big_endian(&decode64(vector.hash)); + for target in [hash - U512::one(), hash, hash + U512::one()] { + ctx.target = target; + let mut hasher = MiningHasher::new(&ctx); + // Both a cold cache and a reused midstate must behave identically. + for _ in 0..2 { + assert_eq!(hasher.hash_if_valid(nonce), (hash < target).then_some(hash)); + } + } + } + } + #[test] fn test_job_context_creation() { let header = [1u8; 32]; From 560f65b9b97873812f7e77de2c61e321d0dc8b74 Mon Sep 17 00:00:00 2001 From: userInner <68625791+userInner@users.noreply.github.com> Date: Wed, 9 Sep 2026 22:53:08 +0800 Subject: [PATCH 4/5] Use uniform inputs for Apple GPU mining Move shared Apple kernel inputs into the Metal constant address space through WGSL uniform bindings. Pack words into vec4 arrays without changing byte order and pad dispatch config to 16 bytes. Allocate matching input usages in the host and adapt raw-kernel validation tools; other kernels retain storage inputs. M4 Pro same-parameter ABBA comparisons gain 0.94-1.52% over the previous kernel. Final comparisons against v4.1.0 gain 12.22% and 12.00% cumulatively. These are offline GPU kernel results, not CPU gains or mainnet revenue claims. Validated all three GPU kernel component/end-to-end suites, 100 CPU parity tasks, strict target boundaries, range coverage, full workspace build and 62 tests, Clippy, rustfmt and Taplo. Code: crates/engine-gpu. PR: https://github.com/Quantus-Network/quantus-miner/pull/96 --- crates/engine-gpu/examples/kernel_compare.rs | 8 +++---- .../engine-gpu/examples/trusted_hashrate.rs | 8 +++---- crates/engine-gpu/src/end_to_end_tests.rs | 12 ++++++----- .../src/kernels/mining_u64_apple.wgsl | 15 ++++++------- crates/engine-gpu/src/kernels/mod.rs | 4 +++- crates/engine-gpu/src/lib.rs | 21 ++++++++++++------- 6 files changed, 40 insertions(+), 28 deletions(-) diff --git a/crates/engine-gpu/examples/kernel_compare.rs b/crates/engine-gpu/examples/kernel_compare.rs index dc3df90..2e406a9 100644 --- a/crates/engine-gpu/examples/kernel_compare.rs +++ b/crates/engine-gpu/examples/kernel_compare.rs @@ -86,10 +86,10 @@ impl Runner { Runner { pipeline, results: mk(132, U::STORAGE | U::COPY_SRC | U::COPY_DST), - midstate: mk(96, U::STORAGE | U::COPY_DST), - start_nonce: mk(64, U::STORAGE | U::COPY_DST), - target: mk(64, U::STORAGE | U::COPY_DST), - cfg: mk(12, U::STORAGE | U::COPY_DST), + midstate: mk(96, U::STORAGE | U::UNIFORM | U::COPY_DST), + start_nonce: mk(64, U::STORAGE | U::UNIFORM | U::COPY_DST), + target: mk(64, U::STORAGE | U::UNIFORM | U::COPY_DST), + cfg: mk(16, U::STORAGE | U::UNIFORM | U::COPY_DST), staging: mk(132, U::MAP_READ | U::COPY_DST), device, queue, diff --git a/crates/engine-gpu/examples/trusted_hashrate.rs b/crates/engine-gpu/examples/trusted_hashrate.rs index 65f280e..5b1eaa4 100644 --- a/crates/engine-gpu/examples/trusted_hashrate.rs +++ b/crates/engine-gpu/examples/trusted_hashrate.rs @@ -78,10 +78,10 @@ impl Runner { Runner { pipeline, results: mk(132, U::STORAGE | U::COPY_SRC | U::COPY_DST), - midstate: mk(96, U::STORAGE | U::COPY_DST), - start_nonce: mk(64, U::STORAGE | U::COPY_DST), - target: mk(64, U::STORAGE | U::COPY_DST), - cfg: mk(12, U::STORAGE | U::COPY_DST), + midstate: mk(96, U::STORAGE | U::UNIFORM | U::COPY_DST), + start_nonce: mk(64, U::STORAGE | U::UNIFORM | U::COPY_DST), + target: mk(64, U::STORAGE | U::UNIFORM | U::COPY_DST), + cfg: mk(16, U::STORAGE | U::UNIFORM | U::COPY_DST), staging: mk(132, U::MAP_READ | U::COPY_DST), device, queue, diff --git a/crates/engine-gpu/src/end_to_end_tests.rs b/crates/engine-gpu/src/end_to_end_tests.rs index d125101..b4483e4 100644 --- a/crates/engine-gpu/src/end_to_end_tests.rs +++ b/crates/engine-gpu/src/end_to_end_tests.rs @@ -40,7 +40,7 @@ pub async fn test_end_to_end_mining( let midstate_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor { label: Some("Midstate Buffer"), contents: bytemuck::cast_slice(&midstate_u32s), - usage: wgpu::BufferUsages::STORAGE, + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::UNIFORM, }); // Target Buffer @@ -53,7 +53,7 @@ pub async fn test_end_to_end_mining( let target_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor { label: Some("Target Buffer"), contents: bytemuck::cast_slice(&target_u32s), - usage: wgpu::BufferUsages::STORAGE, + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::UNIFORM, }); // Start Nonce Buffer @@ -67,7 +67,7 @@ pub async fn test_end_to_end_mining( let start_nonce_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor { label: Some("Start Nonce Buffer"), contents: bytemuck::cast_slice(&start_nonce_u32s), - usage: wgpu::BufferUsages::STORAGE, + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::UNIFORM, }); // Results Buffer @@ -89,8 +89,10 @@ pub async fn test_end_to_end_mining( let dispatch_config_data: [u32; 3] = [256, 1, 256]; let dispatch_config_buffer = device.create_buffer(&wgpu::BufferDescriptor { label: Some("Dispatch Config Buffer"), - size: 12, - usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST, + size: 16, + usage: wgpu::BufferUsages::STORAGE + | wgpu::BufferUsages::UNIFORM + | wgpu::BufferUsages::COPY_DST, mapped_at_creation: false, }); queue.write_buffer( diff --git a/crates/engine-gpu/src/kernels/mining_u64_apple.wgsl b/crates/engine-gpu/src/kernels/mining_u64_apple.wgsl index f1f0fd9..7f73ef6 100644 --- a/crates/engine-gpu/src/kernels/mining_u64_apple.wgsl +++ b/crates/engine-gpu/src/kernels/mining_u64_apple.wgsl @@ -4,10 +4,11 @@ @group(0) @binding(0) var results: array>; // Sponge state after absorbing header + high nonce half (12 felts as LE u32 pairs), // precomputed on the host per batch. See pow_core::mining_midstate. -@group(0) @binding(1) var midstate: array; -@group(0) @binding(2) var start_nonce: array; -@group(0) @binding(3) var difficulty_target: array; -@group(0) @binding(4) var dispatch_config: array; +// Pack u32 words into vec4s to preserve the byte layout with uniform alignment. +@group(0) @binding(1) var midstate: array, 6>; +@group(0) @binding(2) var start_nonce: array, 4>; +@group(0) @binding(3) var difficulty_target: array, 4>; +@group(0) @binding(4) var dispatch_config: vec4; const P64: u64 = 0xFFFFFFFF00000001lu; // EPS64 = 2^32 - 1 = 2^64 mod P @@ -266,15 +267,15 @@ fn mining_main(@builtin(global_invocation_id) global_id: vec3) { // Hoist uniform storage reads out of the nonce loop var mid: array; for (var i = 0u; i < 12u; i++) { - mid[i] = (u64(midstate[2u * i + 1u]) << 32u) | u64(midstate[2u * i]); + mid[i] = (u64(midstate[(2u * i + 1u) / 4u][(2u * i + 1u) % 4u]) << 32u) | u64(midstate[(2u * i) / 4u][(2u * i) % 4u]); } var tgt: array; for (var i = 0u; i < 16u; i++) { - tgt[i] = difficulty_target[i]; + tgt[i] = difficulty_target[(i) / 4u][(i) % 4u]; } var nonce_base: array; for (var i = 0u; i < 16u; i++) { - nonce_base[i] = start_nonce[i]; + nonce_base[i] = start_nonce[(i) / 4u][(i) % 4u]; } for (var j = 0u; j < nonces_per_thread; j = j + 1u) { diff --git a/crates/engine-gpu/src/kernels/mod.rs b/crates/engine-gpu/src/kernels/mod.rs index b3fc558..ce89924 100644 --- a/crates/engine-gpu/src/kernels/mod.rs +++ b/crates/engine-gpu/src/kernels/mod.rs @@ -1,6 +1,8 @@ //! Poseidon2 mining kernels. //! -//! Same `mining_main` bindings; must stay bit-exact with `pow_core`. +//! Same `mining_main` binding numbers; must stay bit-exact with `pow_core`. +//! Apple uses uniform bindings for inputs 1..4; other kernels use storage. +//! Input word layouts are identical, with dispatch config padded to 16 bytes. //! //! - Apple Metal + `SHADER_INT64` → Apple Metal u64 (`mining_u64_apple.wgsl`) //! - other GPUs + `SHADER_INT64` → native u64 (`mining_u64.wgsl`) diff --git a/crates/engine-gpu/src/lib.rs b/crates/engine-gpu/src/lib.rs index 8eb8784..016fa37 100644 --- a/crates/engine-gpu/src/lib.rs +++ b/crates/engine-gpu/src/lib.rs @@ -24,6 +24,8 @@ struct GpuContext { queue: wgpu::Queue, pipeline: wgpu::ComputePipeline, + input_usage: wgpu::BufferUsages, + // Cached vendor configuration optimal_workgroups: u32, } @@ -68,7 +70,7 @@ impl GpuContext { let midstate_buffer = self.device.create_buffer(&wgpu::BufferDescriptor { label: Some("Midstate Buffer"), size: 96, - usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST, + usage: self.input_usage | wgpu::BufferUsages::COPY_DST, mapped_at_creation: false, }); @@ -76,7 +78,7 @@ impl GpuContext { let target_buffer = self.device.create_buffer(&wgpu::BufferDescriptor { label: Some("Target Buffer"), size: 64, - usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST, + usage: self.input_usage | wgpu::BufferUsages::COPY_DST, mapped_at_creation: false, }); @@ -84,7 +86,7 @@ impl GpuContext { let start_nonce_buffer = self.device.create_buffer(&wgpu::BufferDescriptor { label: Some("Start Nonce Buffer"), size: 64, - usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST, + usage: self.input_usage | wgpu::BufferUsages::COPY_DST, mapped_at_creation: false, }); @@ -99,11 +101,11 @@ impl GpuContext { mapped_at_creation: false, }); - // Dispatch config: [total_threads, nonces_per_thread, total_nonces] = 3 u32s + // Three dispatch words plus padding for the Apple uniform vec4 binding. let dispatch_config_buffer = self.device.create_buffer(&wgpu::BufferDescriptor { label: Some("Dispatch Config Buffer"), - size: 12, - usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST, + size: 16, + usage: self.input_usage | wgpu::BufferUsages::COPY_DST, mapped_at_creation: false, }); @@ -447,6 +449,11 @@ impl GpuEngine { queue, pipeline, optimal_workgroups, + input_usage: if kernel == Kernel::Apple { + wgpu::BufferUsages::UNIFORM + } else { + wgpu::BufferUsages::STORAGE + }, }), device_type: info.device_type, name: info.name.clone(), @@ -807,7 +814,7 @@ fn run_single_batch( let nonces_per_thread = ((batch_size as u64).div_ceil(total_threads)).max(1) as u32; // Dispatch config: [total_threads, nonces_per_thread, total_nonces] - let dispatch_config = [total_threads as u32, nonces_per_thread, batch_size]; + let dispatch_config = [total_threads as u32, nonces_per_thread, batch_size, 0]; // Write dispatch config gpu_ctx.queue.write_buffer( From 29c10ee6124cac2d712983c776d0848bf168ed63 Mon Sep 17 00:00:00 2001 From: userInner <68625791+userInner@users.noreply.github.com> Date: Wed, 9 Sep 2026 23:11:16 +0800 Subject: [PATCH 5/5] Reduce Apple GPU arithmetic with u32 carries and borrows MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Replace signed 64-bit intermediate reduction sums with explicit u32 limb carries/borrows, and fold matrix accumulators in radix 2^32. Preserve lazy residues, strict target comparison, rounds and dispatch. Document correction bounds in the shader. Expand the independent u128 arithmetic example to 65,920 cases, including accumulator carry edges. M4 Pro offline ABBA: final kernel +18.386% over the previous commit; +30.919–32.357% versus v4.1.0. GPU-only CLI averages 35.9475 to 42.9250 MH/s (+19.41%). Results are hardware-specific, not earnings. Validation: 62 workspace tests, locked builds, all-target Clippy, fmt, Taplo; 65,920 arithmetic cases, all three GPU kernel suites, 100 parity tasks and 30 dispatch coverage cases passed. Mainnet mining stays off. PR: https://github.com/Quantus-Network/quantus-miner/pull/96 Code: crates/engine-gpu/src/kernels/mining_u64_apple.wgsl --- .../engine-gpu/examples/arithmetic_edges.rs | 27 ++++++++----- .../src/kernels/mining_u64_apple.wgsl | 39 ++++++++++++------- 2 files changed, 43 insertions(+), 23 deletions(-) diff --git a/crates/engine-gpu/examples/arithmetic_edges.rs b/crates/engine-gpu/examples/arithmetic_edges.rs index 6268458..1625775 100644 --- a/crates/engine-gpu/examples/arithmetic_edges.rs +++ b/crates/engine-gpu/examples/arithmetic_edges.rs @@ -39,8 +39,13 @@ async fn run() { } } } + for a in edges { + for carries in [0u64, 1, 2, 11, 12, 26, 27, u64::from(u32::MAX)] { + inputs.extend([a, carries]); + } + } let mut rng = rand::rngs::StdRng::seed_from_u64(0x20260909); - for _ in 0..4096 { + for _ in 0..65536 { inputs.extend([rng.next_u64(), rng.next_u64()]); } let count = inputs.len() / 2; @@ -58,11 +63,12 @@ fn arithmetic_edges(@builtin(global_invocation_id) id: vec3) {{ let a = pairs[id.x * 2u]; let b = pairs[id.x * 2u + 1u]; let v = mul_wide(a, b); - answer[id.x * 5u] = v.lo; - answer[id.x * 5u + 1u] = v.hi; - answer[id.x * 5u + 2u] = gf64_canon(gf64_mul(a, b)); - answer[id.x * 5u + 3u] = gf64_canon(gf64_sqr(a)); - answer[id.x * 5u + 4u] = gf64_canon(gf64_reduce(U128(a,b))); + answer[id.x * 6u] = v.lo; + answer[id.x * 6u + 1u] = v.hi; + answer[id.x * 6u + 2u] = gf64_canon(gf64_mul(a, b)); + answer[id.x * 6u + 3u] = gf64_canon(gf64_sqr(a)); + answer[id.x * 6u + 4u] = gf64_canon(gf64_reduce(U128(a,b))); + answer[id.x * 6u + 5u] = gf64_canon(acc_fold(Acc(a,u32(b)))); }} "#, kernel_source, count @@ -84,7 +90,7 @@ fn arithmetic_edges(@builtin(global_invocation_id) id: vec3) {{ contents: bytemuck::cast_slice(&inputs), usage: wgpu::BufferUsages::STORAGE, }); - let size = (count * 5 * 8) as u64; + let size = (count * 6 * 8) as u64; let output = device.create_buffer(&wgpu::BufferDescriptor { label: None, size, @@ -133,18 +139,19 @@ fn arithmetic_edges(@builtin(global_invocation_id) id: vec3) {{ let b = u128::from(inputs[i * 2 + 1]); let product = a * b; assert_eq!( - &answers[i * 5..i * 5 + 5], + &answers[i * 6..i * 6 + 6], &[ product as u64, (product >> 64) as u64, (product % u128::from(p)) as u64, ((a * a) % u128::from(p)) as u64, - (((b << 64) | a) % u128::from(p)) as u64 + (((b << 64) | a) % u128::from(p)) as u64, + ((a + (u128::from(b as u32) << 64)) % u128::from(p)) as u64 ], "pair {i}" ); } drop(mapped); staging.unmap(); - println!("ARITHMETIC OK: {count} operand pairs, wide product / modular multiply / square / arbitrary u128 reduction"); + println!("ARITHMETIC OK: {count} operand pairs, wide product / modular multiply / square / arbitrary u128 reduction / accumulator fold"); } diff --git a/crates/engine-gpu/src/kernels/mining_u64_apple.wgsl b/crates/engine-gpu/src/kernels/mining_u64_apple.wgsl index 7f73ef6..b3ba78e 100644 --- a/crates/engine-gpu/src/kernels/mining_u64_apple.wgsl +++ b/crates/engine-gpu/src/kernels/mining_u64_apple.wgsl @@ -61,9 +61,14 @@ fn acc_add2(a: Acc, b: Acc) -> Acc { } fn acc_fold(a: Acc) -> u64 { - let c = u64(a.carries); - let t = a.lo + ((c << 32u) - c); - return t + select(0lu, EPS64, t < a.lo); + // Add carries * (2^32 - 1) in radix 2^32. The low subtraction + // borrows at most one; delta is nonnegative, even when carries=0. + let low = u32(a.lo) - a.carries; + let delta = a.carries - select(0u, 1u, u32(a.lo) < a.carries); + let high0 = u32(a.lo >> 32u); + let high = high0 + delta; + let result = (u64(high) << 32u) | u64(low); + return result + select(0lu, EPS64, high < high0); } struct U128 { @@ -74,16 +79,24 @@ struct U128 { // Reduce a 128-bit value (lo + hi*2^64) mod P using // 2^64 ≡ EPS64 and 2^96 ≡ -1 (mod P). fn gf64_reduce(v: U128) -> u64 { - // In radix B=2^32, B^2 = B-1 and B^3 = -1 modulo P. Fold - // the four limbs directly, avoiding 64-bit carry/borrow comparisons. - let low = i64(u32(v.lo)) - i64(u32(v.hi)) - i64(u32(v.hi >> 32u)); - let high = i64(u32(v.lo >> 32u)) + i64(u32(v.hi)) + (low >> 32u); - let result = (u64(u32(high)) << 32u) | u64(u32(low)); - // -1 <= high <= 2B-2: the correction is -1, 0, or 1. - // A positive correction has result <= B^2-B-1; a negative one - // has its high limb equal to B-1. Neither requires another fold. - let correction = bitcast(high >> 32u); - return result + ((correction << 32u) - correction); + // With B=2^32, reduce to (w0-w2-w3) + (w1+w2)*B. + // Track the low limb's two possible borrows in u32 instead of + // carrying a signed 64-bit intermediate through the reduction. + let w0 = u32(v.lo); + let w1 = u32(v.lo >> 32u); + let w2 = u32(v.hi); + let w3 = u32(v.hi >> 32u); + let low0 = w0 - w2; + let low = low0 - w3; + let borrow = select(0u, 1u, w0 < w2) + select(0u, 1u, low0 < w3); + let high0 = w1 + w2; + let high = high0 - borrow; + // The mathematical high limb lies in [-1, 2B-2], so carry minus + // borrow is -1, 0 or 1. A positive correction cannot overflow; + // a negative correction has high=B-1 and cannot underflow. + let correction = i64(select(0i, 1i, high0 < w1) - select(0i, 1i, high0 < borrow)); + let bits = bitcast(correction); + return ((u64(high) << 32u) | u64(low)) + ((bits << 32u) - bits); } fn mul_wide(a: u64, b: u64) -> U128 {