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/engine-gpu/examples/arithmetic_edges.rs b/crates/engine-gpu/examples/arithmetic_edges.rs new file mode 100644 index 0000000..1625775 --- /dev/null +++ b/crates/engine-gpu/examples/arithmetic_edges.rs @@ -0,0 +1,157 @@ +//! 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; + +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]); + } + } + // 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)]); + } + } + } + } + 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..65536 { + 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; +@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 * 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 + ); + 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 * 6 * 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 * 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, + ((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 / accumulator fold"); +} 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..2e406a9 --- /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::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, + } + } + + 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/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 7feb7cf..b3ba78e 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 @@ -60,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 { @@ -73,13 +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 { - 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); + // 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 { @@ -91,8 +108,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 +134,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 { @@ -137,16 +162,16 @@ 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))) ); } // 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 +183,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 +249,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) { @@ -255,15 +280,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) { @@ -298,60 +323,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 +598,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); } 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( 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];