Optimize CPU mining and Apple GPU hashing

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);
+}

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];

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<storage, read> pairs: array<u64>;
+@group(0) @binding(6) var<storage, read_write> answer: array<u64>;
+@compute @workgroup_size(64)
+fn arithmetic_edges(@builtin(global_invocation_id) id: vec3<u32>) {{
+    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");
+}

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");
+}