Merge pull request 'perf: parameter-bank uniforms and per-arch carry path (+20.5% sm_120)' (#22) from perf/sm120-spills into main

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;

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<CudaStream>,
     results: CudaSlice<u32>,
-    midstate: CudaSlice<u64>,
-    start_nonce: CudaSlice<u32>,
-    target: CudaSlice<u32>,
     host_results: Vec<u32>,
-    /// 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<U512>,
 }
 
 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::<MiningUniforms>() == 224);
+
 pub struct CudaEngine {
     engine_id: usize,
     devices: Vec<Arc<Device>>,
@@ -166,10 +192,8 @@ impl Device {
         let stream = self.ctx.new_stream()?;
         Ok(WorkerResources {
             results: stream.alloc_zeros::<u32>(RESULTS_LEN)?,
-            midstate: stream.alloc_zeros::<u64>(12)?,
-            start_nonce: stream.alloc_zeros::<u32>(16)?,
-            target: stream.alloc_zeros::<u32>(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);