diff --git a/crates/engine-cuda/src/kernels/mining.cu b/crates/engine-cuda/src/kernels/mining.cu index 5242041..2bcb56e 100644 --- a/crates/engine-cuda/src/kernels/mining.cu +++ b/crates/engine-cuda/src/kernels/mining.cu @@ -28,11 +28,43 @@ typedef unsigned int u32; // 2^64 mod P = 2^32 - 1 #define EPS 0xFFFFFFFFull +// lair: whether gf_add and gf_reduce should use inline-PTX carry arithmetic or +// plain C is architecture-dependent, and the difference is large in both +// directions. An `asm` block is opaque to nvcc's optimiser, so carry arithmetic +// written as PTX blocks common-subexpression elimination and strength reduction +// across neighbouring field operations. On sm_120 that costs a quarter of the +// kernel; on sm_86/89 the hand-written sequence still wins. +// +// Whole-kernel instructions / ALU ops (ptxas + cuobjdump, CUDA 13.0) and +// measured hashrate, three interleaved rounds each, parity 40/40 vs CPU: +// +// sm_120 sm_86/89 +// CARRY=1 37,432 / 26,688 1126.6 37,027 benjy 461.6 quadbrat 85.5 +// CARRY=0 27,837 / 19,519 1333.3 41,196 benjy 401.5 quadbrat 76.3 +// +// So +18.3% on the 5090s and -13.0% / -10.8% on the 4090 and 3060. Hence the +// per-arch default. Unlike LAIR_INT_UNROLL -- where the static measures pointed +// one way and the cards' power envelope the other -- this one was measured on +// hardware on all three architectures, and the static count predicted the sign +// correctly on each. +// +// LAIR_PTX_ACC stays 1 everywhere: in the Acc accumulators the add-with-carry +// pair is the whole operation, so there is nothing around it to optimise. On +// sm_120, carry off with Acc off measured 1285.8 against 1333.3 with Acc on. +// +// PR #17 measured the PTX carry path as a win on every card, and that was true +// of the kernel it was written for -- which spilled 104 bytes per thread. Once +// the spills went (see MiningUniforms) the trade reversed on Blackwell only. +// Re-measure per architecture before touching this. #ifndef LAIR_PTX_CARRY -#define LAIR_PTX_CARRY 1 +# if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 1200 +# define LAIR_PTX_CARRY 0 +# else +# define LAIR_PTX_CARRY 1 +# endif #endif #ifndef LAIR_PTX_ACC -#define LAIR_PTX_ACC LAIR_PTX_CARRY +#define LAIR_PTX_ACC 1 #endif // Fold the next round's constant into the linear layer's final reduction // instead of a separate gf_add per element before the S-box. @@ -90,6 +122,20 @@ __device__ __forceinline__ u64 gf_mul(u64 a, u64 b) { return gf_reduce(a * b, __umul64hi(a, b)); } +// Squaring deliberately uses the general product. nvcc does not specialise +// `a * a`: it emits the same six IMAD.WIDE in 48 instructions as gf_mul +// (measured, sm_120). The textbook saving is real but does not pay here -- +// a = a1*2^32 + a0 gives a^2 = a0^2 + a0*a1*2^33 + a1^2*2^64, three 32x32 +// multiplies instead of four, and it was implemented, verified bit-exact +// against __int128 over 40M values, and measured 5.7% SLOWER on beast's 5090s +// (1072.9 vs 1137.7 MH/s, three interleaved rounds). It saves 6.7% of the +// kernel's widening multiplies and costs 11% more instructions: 5 multiplies +// in 56 instructions against 6 in 48. Four formulations (condition-code +// carry, compare carry, and two mad.wide.u32 variants) all compiled to the +// same 5/56, so 56 is the floor. The scarce resource in this kernel is +// instruction issue, not multiply throughput -- which is also why raising +// occupancy lost (see LAIR_INT_UNROLL). Anything that wins here has to remove +// work, not re-express it. __device__ __forceinline__ u64 gf_sqr(u64 a) { return gf_reduce(a * a, __umul64hi(a, a)); } @@ -109,6 +155,14 @@ __device__ __forceinline__ u64 gf_canon(u64 a) { #ifndef LAIR_LAZY_ADD #define LAIR_LAZY_ADD 1 #endif +// lair: 2 stays the default on every architecture. Unroll 1 looks better on +// sm_120 by every static measure -- 78 registers instead of 106, no spill +// either way, a third resident block per SM instead of two -- and measured +// 0.92% SLOWER on beast's 5090s (1128.4 vs 1138.8 MH/s, three interleaved +// rounds, spread <=0.1%). These cards are power-bound, not latency-bound: at +// the 400 W cap the extra resident warps buy power draw, and the card clocks +// down to pay for it (2325/2385 MHz against 2355/2415). Occupancy is not the +// lever here; work per hash is. Do not "fix" this without a measurement. #ifndef LAIR_INT_UNROLL #define LAIR_INT_UNROLL 2 #endif @@ -116,10 +170,14 @@ __device__ __forceinline__ u64 gf_canon(u64 a) { #define LAIR_PRAGMA(x) _Pragma(#x) #define LAIR_UNROLL(n) LAIR_PRAGMA(unroll n) #ifndef LAIR_PTX_CARRY -#define LAIR_PTX_CARRY 1 +# if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 1200 +# define LAIR_PTX_CARRY 0 +# else +# define LAIR_PTX_CARRY 1 +# endif #endif #ifndef LAIR_PTX_ACC -#define LAIR_PTX_ACC LAIR_PTX_CARRY +#define LAIR_PTX_ACC 1 #endif // Fold the next round's constant into the linear layer's final reduction // instead of a separate gf_add per element before the S-box. @@ -367,11 +425,24 @@ __device__ __forceinline__ u32 bswap32(u32 v) { // midstate 12 felts as u64 // start_nonce 16 u32 LE limbs // target 16 u32 LE limbs +// lair: the launch-uniform inputs ride in the parameter bank instead of being +// copied into per-thread registers (quantus/miner#3). Keeping midstate (24 +// registers), target (16) and the nonce base (16) live across the whole nonce +// loop cost 56 registers of residency that the permutation's own working set +// needs more: ptxas spilled 104 bytes per thread on sm_120 and 44 on sm_86/89. +// Parameters are constant memory -- broadcast and cached, no register cost -- +// and they arrive with the launch, so the host no longer copies midstate, +// start_nonce or target to the device once per batch. +struct MiningUniforms { + u64 midstate[12]; + u32 start_nonce[16]; + u32 target[16]; +}; +static_assert(sizeof(MiningUniforms) == 224, "MiningUniforms must match the host layout"); + extern "C" __global__ void __launch_bounds__(256) mining_main(u32* __restrict__ results, - const u64* __restrict__ midstate, - const u32* __restrict__ start_nonce, - const u32* __restrict__ target, + const MiningUniforms uni, u32 total_threads, u32 nonces_per_thread, u32 total_nonces) @@ -381,43 +452,33 @@ mining_main(u32* __restrict__ results, if (tid >= total_threads) return; u32 base_index = tid * nonces_per_thread; - u64 mid[12]; -#pragma unroll - for (int i = 0; i < 12; i++) mid[i] = midstate[i]; - u32 tgt[16]; -#pragma unroll - for (int i = 0; i < 16; i++) tgt[i] = target[i]; - u32 nonce_base[16]; -#pragma unroll - for (int i = 0; i < 16; i++) nonce_base[i] = start_nonce[i]; - for (u32 j = 0; j < nonces_per_thread; j++) { u32 logical_index = base_index + j; if (logical_index >= total_nonces) break; if (j > 0u && *((volatile u32*)results) != 0u) return; - // Low 256 bits only; the host guarantees no carry into limbs 8..15. - u32 current_nonce[16]; - u32 val0 = nonce_base[0]; + // Low 256 bits only; the host guarantees no carry into limbs 8..15, so + // the high half is never materialised here -- on a hit it is written + // straight from the parameter bank. + u32 current_nonce[8]; + u32 val0 = uni.start_nonce[0]; u32 sum0 = val0 + logical_index; current_nonce[0] = sum0; u32 carry = sum0 < val0 ? 1u : 0u; #pragma unroll for (int i = 1; i < 8; i++) { - u32 val = nonce_base[i]; + u32 val = uni.start_nonce[i]; u32 sum = val + carry; current_nonce[i] = sum; carry = sum < val ? 1u : 0u; } -#pragma unroll - for (int i = 8; i < 16; i++) current_nonce[i] = nonce_base[i]; // Resume the sponge from the midstate: absorb the low nonce half, // pad, squeeze twice (3 permutations instead of 5; the second squeeze // only for candidates). u64 st[12]; #pragma unroll - for (int i = 0; i < 12; i++) st[i] = mid[i]; + for (int i = 0; i < 12; i++) st[i] = uni.midstate[i]; #pragma unroll for (int i = 0; i < 8; i++) st[i] = gf_add(st[i], (u64)bswap32(current_nonce[7 - i])); permute(st); @@ -437,7 +498,7 @@ mining_main(u32* __restrict__ results, #pragma unroll for (int i = 0; i < 8; i++) { u32 h = bswap32(first[i]); - u32 t = tgt[15 - i]; + u32 t = uni.target[15 - i]; if (h != t) { cmp = h > t ? 1u : 2u; break; } } if (cmp == 1u) continue; @@ -457,7 +518,7 @@ mining_main(u32* __restrict__ results, #pragma unroll for (int i = 0; i < 8; i++) { u32 h = hash_le[7 - i]; - u32 t = tgt[7 - i]; + u32 t = uni.target[7 - i]; if (h != t) { below = h < t; break; } } } @@ -465,10 +526,11 @@ mining_main(u32* __restrict__ 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]; - } + for (int i = 0; i < 8; i++) results[1 + i] = current_nonce[i]; +#pragma unroll + for (int i = 8; i < 16; i++) results[1 + i] = uni.start_nonce[i]; +#pragma unroll + for (int i = 0; i < 16; i++) results[17 + i] = hash_le[i]; __threadfence(); } return; diff --git a/crates/engine-cuda/src/lib.rs b/crates/engine-cuda/src/lib.rs index 5e5ffd7..c1c2253 100644 --- a/crates/engine-cuda/src/lib.rs +++ b/crates/engine-cuda/src/lib.rs @@ -10,7 +10,8 @@ //! driver is present, and falls back to wgpu otherwise. use cudarc::driver::{ - CudaContext, CudaFunction, CudaSlice, CudaStream, DriverError, LaunchConfig, PushKernelArg, + CudaContext, CudaFunction, CudaSlice, CudaStream, DeviceRepr, DriverError, LaunchConfig, + PushKernelArg, }; use cudarc::nvrtc::Ptx; use engine_cpu::{CancelCheck, Candidate, EngineStatus, FoundOrigin, MinerEngine, Range}; @@ -62,16 +63,41 @@ struct Device { struct WorkerResources { stream: Arc, results: CudaSlice, - midstate: CudaSlice, - start_nonce: CudaSlice, - target: CudaSlice, host_results: Vec, - /// Target written for this job context; rewritten when it changes. + /// Target limbs for this job context, recomputed when the target changes. + target_u32s: [u32; 16], + /// Target the limbs above were derived from. target_written: Option, } const RESULTS_LEN: usize = 1 + 16 + 16; +/// lair: the kernel's launch-uniform inputs, passed by value so they land in +/// the parameter bank (constant memory) instead of per-thread registers. +/// +/// Layout must match `struct MiningUniforms` in `kernels/mining.cu`. Passing +/// them as parameters rather than device buffers took the kernel from 128 +/// registers with 104 bytes of per-thread spill on sm_120 (44 on sm_86/89) to +/// 106 / 80 registers with none, and removed three host-to-device copies per +/// batch. +#[repr(C)] +#[derive(Clone, Copy)] +struct MiningUniforms { + midstate: [u64; 12], + start_nonce: [u32; 16], + target: [u32; 16], +} + +// SAFETY: `#[repr(C)]` plain-old-data with no padding (12 x u64 then 32 x u32), +// matching the kernel's parameter layout; there is nothing to validate beyond +// the layout, which is what this marker asserts. +unsafe impl DeviceRepr for MiningUniforms {} + +// A mismatch here would silently feed the kernel the wrong midstate or target, +// so it is a build error rather than a mining bug. The kernel carries the same +// assertion. +const _: () = assert!(std::mem::size_of::() == 224); + pub struct CudaEngine { engine_id: usize, devices: Vec>, @@ -166,10 +192,8 @@ impl Device { let stream = self.ctx.new_stream()?; Ok(WorkerResources { results: stream.alloc_zeros::(RESULTS_LEN)?, - midstate: stream.alloc_zeros::(12)?, - start_nonce: stream.alloc_zeros::(16)?, - target: stream.alloc_zeros::(16)?, host_results: vec![0u32; RESULTS_LEN], + target_u32s: [0u32; 16], target_written: None, stream, }) @@ -263,21 +287,14 @@ impl CudaEngine { range: Range, cancel: &dyn CancelCheck, ) -> EngineStatus { - // Target is per job; write it once per context. + // Target is per job; convert it to limbs once per context. It reaches + // the kernel in the parameter bank, so there is nothing to upload. if res.target_written != Some(ctx.target) { let target_bytes = ctx.target.to_little_endian(); - let mut target_u32s = [0u32; 16]; for i in 0..16 { - target_u32s[i] = + res.target_u32s[i] = u32::from_le_bytes(target_bytes[i * 4..(i + 1) * 4].try_into().unwrap()); } - if res - .stream - .memcpy_htod(&target_u32s, &mut res.target) - .is_err() - { - return self.device_lost(dev, 0); - } res.target_written = Some(ctx.target); } @@ -417,10 +434,14 @@ impl CudaEngine { let nonce_be = batch_start.to_big_endian(); let midstate = pow_core::mining_midstate(ctx.header, nonce_be[..32].try_into().unwrap()); + let uni = MiningUniforms { + midstate, + start_nonce: start_u32s, + target: res.target_u32s, + }; + let stream = res.stream.clone(); let r: Result<(), DriverError> = (|| { - stream.memcpy_htod(&start_u32s, &mut res.start_nonce)?; - stream.memcpy_htod(&midstate, &mut res.midstate)?; stream.memset_zeros(&mut res.results)?; let cfg = LaunchConfig { grid_dim: (num_blocks, 1, 1), @@ -430,9 +451,7 @@ impl CudaEngine { let mut launch = stream.launch_builder(&dev.func); launch .arg(&mut res.results) - .arg(&res.midstate) - .arg(&res.start_nonce) - .arg(&res.target) + .arg(&uni) .arg(&total_threads_u32) .arg(&nonces_per_thread) .arg(&batch_size);