diff --git a/README.md b/README.md index 5b0800e..7b27b0a 100644 --- a/README.md +++ b/README.md @@ -58,7 +58,7 @@ node's chain config dir (`/chains//`): | `--cpu-workers ` | `MINER_CPU_WORKERS` | Number of CPU worker threads | Auto-detect | | `--gpu-devices ` | `MINER_GPU_DEVICES` | Number of GPU devices | Auto-detect | | `--cuda-gpu` | `MINER_CUDA_GPU` | Use native CUDA instead of wgpu/Vulkan (NVIDIA) | off | -| `--gpu-batch-size ` | `MINER_GPU_BATCH_SIZE` | GPU batch size in nonces | 1000000 | +| `--gpu-batch-size ` | `MINER_GPU_BATCH_SIZE` | GPU batch size in nonces | 1000000 (32000000 with `--cuda-gpu`) | | `--cpu-batch-size ` | `MINER_CPU_BATCH_SIZE` | CPU batch size in hashes | 10000 | | `--gpu-throttle-ms ` | `MINER_GPU_THROTTLE_MS` | Sleep duration (ms) between GPU batches | 0 | | `--metrics-port ` | `MINER_METRICS_PORT` | Prometheus metrics port | 9900 | diff --git a/crates/engine-cuda/src/kernels/mining.cu b/crates/engine-cuda/src/kernels/mining.cu index 58c21a3..e53f3cc 100644 --- a/crates/engine-cuda/src/kernels/mining.cu +++ b/crates/engine-cuda/src/kernels/mining.cu @@ -2,7 +2,7 @@ typedef unsigned int u32; typedef unsigned long long u64; struct MiningParams { - u32 midstate[24]; + u32 prestate[24]; u32 start_nonce[16]; u32 difficulty_target[16]; u32 dispatch_config[3]; @@ -271,13 +271,8 @@ __device__ __forceinline__ void int_layer64(u64 *state, const u64 *rc12, } } -__device__ __forceinline__ void permute64(u64 *state) { +__device__ __forceinline__ void permute64_after_initial(u64 *state) { u64 rc[12]; - #pragma unroll - for (int i = 0; i < 12; i++) { - rc[i] = RC_INITIAL[0][i]; - } - ext_layer64(state, rc, 0ULL); #pragma unroll 1 for (int r = 0; r < 4; r++) { #pragma unroll @@ -313,6 +308,35 @@ __device__ __forceinline__ void permute64(u64 *state) { } } +__device__ __forceinline__ void permute64(u64 *state) { + u64 rc[12]; + #pragma unroll + for (int i = 0; i < 12; i++) { + rc[i] = RC_INITIAL[0][i]; + } + ext_layer64(state, rc, 0ULL); + permute64_after_initial(state); +} + +__device__ __forceinline__ void permute64_twice_after_initial(u64 *state) { + #pragma unroll 1 + for (int pass = 0; pass < 2; pass++) { + if (pass != 0) { + u64 rc[12]; + #pragma unroll + for (int i = 0; i < 12; i++) { + rc[i] = RC_INITIAL[0][i]; + } + ext_layer64(state, rc, 0ULL); + } + permute64_after_initial(state); + if (pass == 0) { + state[0] = gf64_add(state[0], 1ULL); + state[1] = gf64_add(state[1], 1ULL); + } + } +} + __device__ __forceinline__ u32 bswap32(u32 v) { return ((v & 0xFFu) << 24) | ((v & 0xFF00u) << 8) | ((v >> 8) & 0xFF00u) | (v >> 24); @@ -402,9 +426,6 @@ extern "C" __global__ void __launch_bounds__(256, 4) hash_nonces(u32 *hashes, co extern "C" __global__ void __launch_bounds__(256, 4) mining_main(u32 *results, const MiningParams params) { - if (*((volatile u32 *)results) != 0u) { - return; - } u32 thread_id = blockIdx.x * blockDim.x + threadIdx.x; u32 total_threads = params.dispatch_config[0]; u32 nonces_per_thread = params.dispatch_config[1]; @@ -417,18 +438,15 @@ extern "C" __global__ void __launch_bounds__(256, 4) mining_main(u32 *results, u64 mid[12]; #pragma unroll for (int i = 0; i < 12; i++) { - mid[i] = ((u64)params.midstate[2 * i + 1] << 32) | (u64)params.midstate[2 * i]; + mid[i] = ((u64)params.prestate[2 * i + 1] << 32) | (u64)params.prestate[2 * i]; } u32 tgt[16]; #pragma unroll for (int i = 0; i < 16; i++) { tgt[i] = params.difficulty_target[i]; } - u32 nonce_base[16]; - #pragma unroll - for (int i = 0; i < 16; i++) { - nonce_base[i] = params.start_nonce[i]; - } + u64 nonce_base_low = ((u64)params.start_nonce[1] << 32) | + (u64)params.start_nonce[0]; #pragma unroll @@ -437,26 +455,39 @@ extern "C" __global__ void __launch_bounds__(256, 4) mining_main(u32 *results, if (logical_index >= total_nonces) { break; } - if (j > 0u && *((volatile u32 *)results) != 0u) { - return; - } - - u32 current_nonce[16]; - nonce_from_index(nonce_base, logical_index, current_nonce); - u64 st[12]; #pragma unroll for (int i = 0; i < 12; i++) { st[i] = mid[i]; } - #pragma unroll - for (int i = 0; i < 8; i++) { - st[i] = gf64_add(st[i], (u64)bswap32(current_nonce[7 - i])); - } - permute64(st); - st[0] = gf64_add(st[0], 1ULL); - st[1] = gf64_add(st[1], 1ULL); - permute64(st); + u64 nonce_low = nonce_base_low + (u64)logical_index; + u64 x6 = (u64)bswap32((u32)(nonce_low >> 32)); + u64 x7 = (u64)bswap32((u32)nonce_low); + u64 x6_2 = x6 + x6; + u64 x6_3 = x6_2 + x6; + u64 x6_4 = x6_2 + x6_2; + u64 x6_6 = x6_3 + x6_3; + u64 x7_2 = x7 + x7; + u64 x7_3 = x7_2 + x7; + u64 x7_4 = x7_2 + x7_2; + u64 x7_6 = x7_3 + x7_3; + u64 c0 = x6 + x7; + u64 c1 = x6_3 + x7; + u64 c2 = x6_2 + x7_3; + u64 c3 = x6 + x7_2; + st[0] = gf64_add(st[0], c0); + st[1] = gf64_add(st[1], c1); + st[2] = gf64_add(st[2], c2); + st[3] = gf64_add(st[3], c3); + st[4] = gf64_add(st[4], x6_2 + x7_2); + st[5] = gf64_add(st[5], x6_6 + x7_2); + st[6] = gf64_add(st[6], x6_4 + x7_6); + st[7] = gf64_add(st[7], x6_2 + x7_4); + st[8] = gf64_add(st[8], c0); + st[9] = gf64_add(st[9], c1); + st[10] = gf64_add(st[10], c2); + st[11] = gf64_add(st[11], c3); + permute64_twice_after_initial(st); u32 first[8]; #pragma unroll @@ -506,11 +537,7 @@ extern "C" __global__ void __launch_bounds__(256, 4) mining_main(u32 *results, if (below) { if (atomicExch(&results[0], 1u) == 0u) { - #pragma unroll - for (int i = 0; i < 16; i++) { - results[1 + i] = current_nonce[i]; - results[17 + i] = hash_le[i]; - } + results[1] = logical_index; } return; } diff --git a/crates/engine-cuda/src/lib.rs b/crates/engine-cuda/src/lib.rs index 94fa132..6cb7397 100644 --- a/crates/engine-cuda/src/lib.rs +++ b/crates/engine-cuda/src/lib.rs @@ -10,16 +10,17 @@ use primitive_types::U512; use std::cell::RefCell; use std::sync::atomic::{AtomicUsize, Ordering}; use std::sync::Arc; +use std::time::{Duration, Instant}; const KERNEL_SRC: &str = include_str!("kernels/mining.cu"); const THREADS_PER_BLOCK: u32 = 256; const MAX_BLOCKS: u32 = 4096; -const RESULTS_U32S: usize = 1 + 16 + 16; +const RESULTS_U32S: usize = 2; #[repr(C)] #[derive(Clone, Copy)] struct MiningParams { - midstate: [u32; 24], + prestate: [u32; 24], start_nonce: [u32; 16], difficulty_target: [u32; 16], dispatch_config: [u32; 3], @@ -43,6 +44,8 @@ struct WorkerBuffers { hashes: CudaSlice, mine: CudaFunction, hash: CudaFunction, + busy: Duration, + busy_since: Instant, } pub struct CudaEngine { @@ -102,6 +105,9 @@ impl CudaEngine { for ordinal in 0..count as usize { let ctx = CudaContext::new(ordinal) .map_err(|e| format!("Failed to create CUDA context for device {ordinal}: {e}"))?; + ctx.set_blocking_synchronize().map_err(|e| { + format!("Failed to enable blocking CUDA synchronization on device {ordinal}: {e}") + })?; let name = ctx .name() .unwrap_or_else(|_| format!("cuda-device-{ordinal}")); @@ -218,9 +224,20 @@ fn create_buffers( mine, hash, stream, + busy: Duration::ZERO, + busy_since: Instant::now(), }) } +fn take_gpu_duty_cycle(buffers: &mut WorkerBuffers) -> Option { + let wall = buffers.busy_since.elapsed(); + let duty = (wall >= Duration::from_secs(1)) + .then(|| 100.0 * buffers.busy.as_secs_f64() / wall.as_secs_f64()); + buffers.busy = Duration::ZERO; + buffers.busy_since = Instant::now(); + duty +} + fn worker_buffers( engine: &CudaEngine, device_index: usize, @@ -254,6 +271,10 @@ enum BatchResult { DeviceLost, } +fn nonces_until_low64_carry(nonce: U512) -> U512 { + (U512::one() << 64).saturating_sub(U512::from(nonce.low_u64())) +} + fn run_single_batch( buffers: &mut WorkerBuffers, ctx: &JobContext, @@ -264,17 +285,18 @@ fn run_single_batch( let total_threads = num_blocks * THREADS_PER_BLOCK; let nonces_per_thread = batch_size.div_ceil(total_threads).max(1); let dispatch = [total_threads, nonces_per_thread, batch_size]; - let start_limbs = pow_core::u512_to_le_u32s(batch_start); let nonce_be = batch_start.to_big_endian(); - let mid = pow_core::mining_midstate_u32s(ctx.header, nonce_be[..32].try_into().unwrap()); + let start_limbs = pow_core::u512_to_le_u32s(batch_start); + let prestate = pow_core::mining_prestate_low64_u32s(ctx.header, nonce_be); let target = pow_core::u512_to_le_u32s(ctx.target); let params = MiningParams { - midstate: mid, + prestate, start_nonce: start_limbs, difficulty_target: target, dispatch_config: dispatch, }; + let launch_start = Instant::now(); if let Err(e) = (|| { buffers.stream.memset_zeros(&mut buffers.results)?; let cfg = LaunchConfig { @@ -294,6 +316,7 @@ fn run_single_batch( log::error!(target: "cuda_engine", "CUDA batch failed: {e}"); return BatchResult::DeviceLost; } + buffers.busy += launch_start.elapsed(); let result_u32s = match buffers.stream.clone_dtoh(&buffers.results) { Ok(v) => v, @@ -305,26 +328,33 @@ fn run_single_batch( let dispatched = (total_threads as u64 * nonces_per_thread as u64).min(batch_size as u64); if result_u32s[0] != 0 { - let mut nonce_limbs = [0u32; 16]; - let mut hash_limbs = [0u32; 16]; - nonce_limbs.copy_from_slice(&result_u32s[1..17]); - hash_limbs.copy_from_slice(&result_u32s[17..33]); - let nonce = pow_core::u512_from_le_u32s(nonce_limbs); - let hash = pow_core::u512_from_le_u32s(hash_limbs); - let hashes_computed = if nonce >= batch_start { - let logical_index = (nonce - batch_start).as_u64(); - let winning_iteration = logical_index % (nonces_per_thread as u64); - (total_threads as u64 * (winning_iteration + 1)).min(dispatched) - } else { - dispatched - }; + let logical_index = result_u32s[1] as u64; + if logical_index >= dispatched { + log::error!( + target: "cuda_engine", + "CUDA returned out-of-range candidate index {logical_index} for {dispatched} dispatched nonces" + ); + return BatchResult::DeviceLost; + } + let nonce = batch_start + U512::from(logical_index); + let hash = pow_core::hash_from_nonce(ctx, nonce); + if hash >= ctx.target { + log::error!( + target: "cuda_engine", + "CUDA candidate verification failed for nonce {}: hash {} is not below target {}", + format_u512(nonce), + format_u512(hash), + format_u512(ctx.target) + ); + return BatchResult::DeviceLost; + } return BatchResult::Found { candidate: Candidate { nonce, work: nonce.to_big_endian(), hash, }, - hash_count: hashes_computed, + hash_count: dispatched, }; } @@ -393,18 +423,26 @@ impl MinerEngine for CudaEngine { return EngineStatus::DeviceLost { hash_count: 0 }; } - let search_start = std::time::Instant::now(); + let search_start = Instant::now(); let mut total_hashes: u64 = 0; let mut current_start = range.start; let mut batch_num = 0u64; + let duty = WORKER_BUFFERS.with(|cell| { + take_gpu_duty_cycle( + cell.borrow_mut() + .as_mut() + .expect("CUDA buffers initialized"), + ) + }); log::info!( target: "cuda_engine", - "CUDA {} search started: range {}..{}, batch size: {} nonces", + "CUDA {} search started: range {}..{}, batch size: {} nonces{}", device_index, format_u512(range.start), format_u512(range.end), - self.batch_size + self.batch_size, + duty.map_or(String::new(), |d| format!(", GPU busy {d:.1}% since previous search")) ); while current_start <= range.end { @@ -418,8 +456,7 @@ impl MinerEngine for CudaEngine { .end .saturating_sub(current_start) .saturating_add(U512::one()); - let headroom = - (U512::one() << 256) - (current_start & ((U512::one() << 256) - U512::one())); + let headroom = nonces_until_low64_carry(current_start); let cap = remaining.min(headroom); let batch_size_u512 = U512::from(self.batch_size); let this_batch_size: u32 = if cap > batch_size_u512 { @@ -513,7 +550,11 @@ mod tests { } fn engine_or_skip() -> Option { - match CudaEngine::try_new(1024, 0) { + engine_or_skip_with(1024) + } + + fn engine_or_skip_with(batch_size: u32) -> Option { + match CudaEngine::try_new(batch_size, 0) { Ok(e) => Some(e), Err(e) => { let msg = e.to_string(); @@ -527,6 +568,19 @@ mod tests { } } + #[test] + fn low64_batches_stop_before_carry() { + assert_eq!(nonces_until_low64_carry(U512::from(u64::MAX)), U512::one()); + assert_eq!( + nonces_until_low64_carry(U512::from(u64::MAX - 7)), + U512::from(8u64) + ); + assert_eq!( + nonces_until_low64_carry((U512::one() << 192) + U512::from(1u64)), + (U512::one() << 64) - U512::one() + ); + } + #[test] fn cuda_matches_nonce_hash_golden_vectors() { let Some(engine) = engine_or_skip() else { @@ -672,6 +726,41 @@ mod tests { CudaEngine::clear_worker_resources(); } + #[test] + fn cuda_found_batch_counts_every_dispatched_nonce() { + let batch_size = 4_000_000u32; + let Some(engine) = engine_or_skip_with(batch_size) else { + return; + }; + let total_threads = + batch_size.div_ceil(THREADS_PER_BLOCK).min(MAX_BLOCKS) * THREADS_PER_BLOCK; + assert!( + batch_size > total_threads, + "batch must span several nonces per thread" + ); + let header = decode32(pow_core::NONCE_HASH_KVS[1].header); + let ctx = engine.prepare_context(header, U512::one()); + let start = U512::from(0xfeed_face_0000_0000u64); + let end = start + U512::from(batch_size - 1); + let cancel = AtomicBool::new(false); + match engine.search_range(&ctx, Range { start, end }, &AtomicBoolCancelCheck(&cancel)) { + EngineStatus::Found { + candidate, + hash_count, + .. + } => { + assert!(candidate.nonce >= start && candidate.nonce <= end); + assert_eq!( + pow_core::hash_from_nonce(&ctx, candidate.nonce), + candidate.hash + ); + assert_eq!(hash_count, batch_size as u64); + } + other => panic!("expected Found, got {other:?}"), + } + CudaEngine::clear_worker_resources(); + } + #[test] fn cuda_search_finds_cpu_verified_solution() { let Some(engine) = engine_or_skip() else { diff --git a/crates/miner-cli/src/main.rs b/crates/miner-cli/src/main.rs index e237385..c456c20 100644 --- a/crates/miner-cli/src/main.rs +++ b/crates/miner-cli/src/main.rs @@ -11,6 +11,7 @@ use std::time::{Duration, Instant}; // CLI defaults const DEFAULT_GPU_BATCH_SIZE: u32 = 1_000_000; +const DEFAULT_CUDA_BATCH_SIZE: u32 = 32_000_000; const DEFAULT_CPU_BATCH_SIZE: u64 = 10_000; #[derive(Subcommand, Debug)] @@ -58,8 +59,9 @@ enum Command { gpu_devices: Option, /// GPU batch size in nonces - controls how often GPU checks for cancellation - #[arg(long = "gpu-batch-size", env = "MINER_GPU_BATCH_SIZE", default_value_t = DEFAULT_GPU_BATCH_SIZE, value_parser = clap::value_parser!(u32).range(1..))] - gpu_batch_size: u32, + /// (default: 1000000, or 32000000 with --cuda-gpu) + #[arg(long = "gpu-batch-size", env = "MINER_GPU_BATCH_SIZE", value_parser = clap::value_parser!(u32).range(1..))] + gpu_batch_size: Option, /// CPU batch size in hashes - controls how often CPU checks for cancellation #[arg(long = "cpu-batch-size", env = "MINER_CPU_BATCH_SIZE", default_value_t = DEFAULT_CPU_BATCH_SIZE, value_parser = clap::value_parser!(u64).range(1..))] @@ -107,8 +109,9 @@ enum Command { gpu_devices: Option, /// GPU batch size in nonces - controls how often GPU checks for cancellation - #[arg(long = "gpu-batch-size", env = "MINER_GPU_BATCH_SIZE", default_value_t = DEFAULT_GPU_BATCH_SIZE, value_parser = clap::value_parser!(u32).range(1..))] - gpu_batch_size: u32, + /// (default: 1000000, or 32000000 with --cuda-gpu) + #[arg(long = "gpu-batch-size", env = "MINER_GPU_BATCH_SIZE", value_parser = clap::value_parser!(u32).range(1..))] + gpu_batch_size: Option, /// CPU batch size in hashes - controls how often CPU checks for cancellation #[arg(long = "cpu-batch-size", env = "MINER_CPU_BATCH_SIZE", default_value_t = DEFAULT_CPU_BATCH_SIZE, value_parser = clap::value_parser!(u64).range(1..))] @@ -212,7 +215,7 @@ async fn main() { tls_cert_sha256, cpu_workers, gpu_devices, - gpu_batch_size, + gpu_batch_size: resolve_gpu_batch_size(gpu_batch_size, cuda_gpu), cpu_batch_size, gpu_throttle_ms, allow_integrated, @@ -239,7 +242,7 @@ async fn main() { run_benchmark( cpu_workers, gpu_devices, - gpu_batch_size, + resolve_gpu_batch_size(gpu_batch_size, cuda_gpu), cpu_batch_size, duration, allow_integrated, @@ -250,6 +253,14 @@ async fn main() { } } +fn resolve_gpu_batch_size(explicit: Option, cuda_gpu: bool) -> u32 { + explicit.unwrap_or(if cuda_gpu { + DEFAULT_CUDA_BATCH_SIZE + } else { + DEFAULT_GPU_BATCH_SIZE + }) +} + fn resolve_auth_token( auth_token: Option, auth_token_file: Option, diff --git a/crates/pow-core/src/lib.rs b/crates/pow-core/src/lib.rs index 41aaf0f..a29e8f9 100644 --- a/crates/pow-core/src/lib.rs +++ b/crates/pow-core/src/lib.rs @@ -1,4 +1,5 @@ use primitive_types::U512; +use qp_poseidon_core::poseidon2::INITIAL_EXTERNAL_CONSTANTS; use qp_poseidon_core::{Goldilocks, Poseidon2}; pub use qp_poseidon_core::SPONGE_WIDTH; @@ -34,6 +35,61 @@ pub fn mining_midstate_u32s(header: [u8; 32], nonce_high_be: [u8; 32]) -> [u32; out } +/// Poseidon2 state after the nonce-invariant part of the next permutation's +/// initial linear layer and first round constants. +/// +/// The low 64 bits of `nonce_be` are deliberately excluded. A GPU kernel can +/// inject their sparse linear contribution per nonce, then resume immediately +/// before the first external-round S-box. Callers must split batches before the +/// low 64 bits carry. +pub fn mining_prestate_low64_u32s(header: [u8; 32], nonce_be: [u8; 64]) -> [u32; 24] { + let prestate = mining_prestate_low64(header, nonce_be); + let mut out = [0u32; 24]; + for (i, felt) in prestate.iter().enumerate() { + out[2 * i] = *felt as u32; + out[2 * i + 1] = (*felt >> 32) as u32; + } + out +} + +/// Canonical-felt form of [`mining_prestate_low64_u32s`]. +pub fn mining_prestate_low64(header: [u8; 32], nonce_be: [u8; 64]) -> [u64; SPONGE_WIDTH] { + let mut state = + mining_midstate(header, nonce_be[..32].try_into().unwrap()).map(Goldilocks::from_u64); + + // Absorb the fixed upper 192 bits of the nonce's low 256-bit half. The + // final two words are the low 64 bits specialized by the GPU kernel. + for (i, chunk) in nonce_be[32..56].chunks_exact(4).enumerate() { + state[i] += Goldilocks::from_u64(u32::from_le_bytes(chunk.try_into().unwrap()) as u64); + } + external_linear_layer(&mut state); + for (felt, constant) in state.iter_mut().zip(INITIAL_EXTERNAL_CONSTANTS[0]) { + *felt += Goldilocks::from_u64(constant); + } + state.map(|felt| felt.as_canonical_u64()) +} + +fn external_linear_layer(state: &mut [Goldilocks; SPONGE_WIDTH]) { + for chunk in state.chunks_exact_mut(4) { + let chunk: &mut [Goldilocks; 4] = chunk.try_into().unwrap(); + let t01 = chunk[0] + chunk[1]; + let t23 = chunk[2] + chunk[3]; + let t0123 = t01 + t23; + let t01123 = t0123 + chunk[1]; + let t01233 = t0123 + chunk[3]; + chunk[3] = t01233 + chunk[0] + chunk[0]; + chunk[1] = t01123 + chunk[2] + chunk[2]; + chunk[0] = t01123 + t01; + chunk[2] = t01233 + t23; + } + + let sums: [Goldilocks; 4] = + std::array::from_fn(|offset| (offset..SPONGE_WIDTH).step_by(4).map(|i| state[i]).sum()); + for (i, felt) in state.iter_mut().enumerate() { + *felt += sums[i % 4]; + } +} + /// Sponge state (canonical u64 felts) after absorbing the 32-byte header and the /// high 32 bytes of the big-endian nonce — the first two of the five Poseidon2 /// permutations of `get_nonce_hash`. This state is identical for every nonce in a