From 7804d9d442bfc0fb5375b6f520536d5e399760b1 Mon Sep 17 00:00:00 2001 From: rob thijssen Date: Sun, 13 Sep 2026 18:32:20 +0300 Subject: [PATCH 1/3] engine-cuda: launch uniforms in the parameter bank, no spills The kernel copied midstate, target and the nonce base into per-thread register arrays and held all 56 registers live across the whole nonce loop, which is exactly what the permutation's own working set needed. ptxas spilled 104 bytes per thread on sm_120 and 44 on sm_86/89; the SASS carried 26 LDL + 26 STL inside the hot loop. They are launch-uniform, so they now travel as a by-value MiningUniforms parameter. Parameters live in constant memory -- broadcast and cached, no register residency -- and arrive with the launch, so three memcpy_htod calls per batch are gone with them. Only the low eight nonce limbs are materialised now; on a hit the high half is written from the parameter bank. 128 -> 80/80/106 registers on sm_86/89/120, zero spill everywhere. Both sides assert the 224-byte layout, so a mismatch is a build error rather than a silent mining bug. Measured A/B on the fleet, same session, same silicon, parity verified against the CPU (25/25 on quadbrat and benjy, 40/40 on beast): quadbrat 3060 83.45 -> 84.60 MH/s +1.38% benjy 4090 457.87 -> 461.77 MH/s +0.85% beast 2x5090 1133.97 -> 1138.79 MH/s +0.43% fleet 1675.29 -> 1685.16 MH/s +0.59% Not taken: a per-arch LAIR_INT_UNROLL of 1 on sm_120. It is better by every static measure -- 78 registers instead of 106, three resident blocks per SM instead of two, no spill either way -- and measured 0.92% slower on beast over three interleaved rounds. These cards are power-bound, not latency-bound: the extra warps buy power draw and the card clocks down to pay for it. Benjy shows the same thing from the other side, delivering +0.85% at a lower clock (2085 vs 2160 MHz), +4.4% per MHz. Occupancy is not the lever here; work per hash is. Co-Authored-By: Claude Opus 5 (1M context) Claude-Session: https://claude.ai/code/session_017ytCkttecH9SNTUEw8VT9D --- crates/engine-cuda/src/kernels/mining.cu | 64 +++++++++++++---------- crates/engine-cuda/src/lib.rs | 65 +++++++++++++++--------- 2 files changed, 80 insertions(+), 49 deletions(-) diff --git a/crates/engine-cuda/src/kernels/mining.cu b/crates/engine-cuda/src/kernels/mining.cu index 5242041..c1188d0 100644 --- a/crates/engine-cuda/src/kernels/mining.cu +++ b/crates/engine-cuda/src/kernels/mining.cu @@ -109,6 +109,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 @@ -367,11 +375,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 +402,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 +448,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 +468,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 +476,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); From 206aa0d8764440d299465dd4647a33867c895fcf Mon Sep 17 00:00:00 2001 From: rob thijssen Date: Sun, 13 Sep 2026 18:42:48 +0300 Subject: [PATCH 2/3] engine-cuda: record why gf_sqr uses the general product The three-multiply squaring is a real algebraic saving and a measured loss. Implemented, verified bit-exact against __int128 over 40M random values plus limb and power-of-two boundaries, then benchmarked on beast: 1072.9 vs 1137.7 MH/s over three interleaved rounds, 5.7% slower. It trades 6.7% of the kernel's widening multiplies for 11% more instructions -- 5 IMAD.WIDE in 56 instructions against the general product's 6 in 48 -- and four formulations (condition-code carry, compare carry, two mad.wide.u32 variants) all compiled to the same 5/56, so that is the floor rather than a tuning failure. Taken with the LAIR_INT_UNROLL result, the picture is consistent: this kernel is bound by instruction issue and by the cards' power envelope, not by multiply throughput or occupancy. Re-expressing the same work cannot win; only removing work can. Left as a comment so the next person does not spend the afternoon rediscovering it. Co-Authored-By: Claude Opus 5 (1M context) Claude-Session: https://claude.ai/code/session_017ytCkttecH9SNTUEw8VT9D --- crates/engine-cuda/src/kernels/mining.cu | 14 ++++++++++++++ 1 file changed, 14 insertions(+) diff --git a/crates/engine-cuda/src/kernels/mining.cu b/crates/engine-cuda/src/kernels/mining.cu index c1188d0..b55cbbf 100644 --- a/crates/engine-cuda/src/kernels/mining.cu +++ b/crates/engine-cuda/src/kernels/mining.cu @@ -90,6 +90,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)); } From 575c4d0abd7438392feb07e684a18c94d59a2707 Mon Sep 17 00:00:00 2001 From: rob thijssen Date: Sun, 13 Sep 2026 19:35:32 +0300 Subject: [PATCH 3/3] engine-cuda: per-arch carry path, +20.5% on sm_120 An `asm` block is opaque to nvcc's optimiser, so carry arithmetic written as inline PTX blocks common-subexpression elimination and strength reduction across neighbouring field operations. On Blackwell that costs a quarter of the kernel. On Ada and Ampere the hand-written sequence still wins, by a similar margin in the other direction, so the default is now chosen per device pass. Whole-kernel instructions (ptxas + cuobjdump, CUDA 13.0) and measured hashrate, three interleaved rounds each: sm_120 sm_86/89 PTX_CARRY=1 37,432 1127.1 MH/s 37,027 benjy 461.6 quadbrat 85.5 PTX_CARRY=0 27,837 1358.0 MH/s 41,196 benjy 401.5 quadbrat 76.3 sm_86/89 keep the exact configuration they had, so those cards are unchanged by construction -- identical macro values, identical 37,027 instructions. Parity 60/60 against the CPU on beast. LAIR_PTX_ACC is now 1 everywhere rather than following LAIR_PTX_CARRY. In the Acc accumulators the add-with-carry pair is the whole operation, so there is nothing around it for the optimiser to fold; on sm_120 with the carry path off, Acc PTX on measured 1333.3 against 1285.8 with it off. PR #17 measured the PTX carry path as a win on every card. That was true of the kernel it was written for, which spilled 104 bytes per thread; once the spills went the trade reversed on Blackwell only. Found by profiling rather than reading: ncu puts us at 43,700 executed instructions per nonce against the rival binary's 22,229 on the same card, a 1.97x ratio that matches the 2.05x throughput gap, with 82% of the excess on the ALU pipe rather than the multiply pipe. This lands 35,785. The remaining gap is still ALU: 26,586 against their 14,413, while the multiply pipe is now close to parity at 8,405 against 7,030. Co-Authored-By: Claude Opus 5 (1M context) Claude-Session: https://claude.ai/code/session_017ytCkttecH9SNTUEw8VT9D --- crates/engine-cuda/src/kernels/mining.cu | 44 +++++++++++++++++++++--- 1 file changed, 40 insertions(+), 4 deletions(-) diff --git a/crates/engine-cuda/src/kernels/mining.cu b/crates/engine-cuda/src/kernels/mining.cu index b55cbbf..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. @@ -138,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.