engine-cuda: deferred carries, carry-flag arithmetic, one nonce per thread

quantus/miner#3 step 5, measured on benjy (RTX 4090, 250 W, miner paused,
interleaved rounds, parity verified on every variant):

  variant                                  MH/s
  straight port (PR #16)                   297-305
  + grid cap 32768 threads/SM              304   (1 nonce per thread up to 4M)
  + deferred-carry linear layers           344   (+13%)
  + carry-flag gf_add / gf_reduce (PTX)    385   (+12%)
  + carry-flag accumulators                414   (+7%)

414 MH/s is 2.86x origin's wgpu build (144.4) on the same card and limit.
Fewer instructions per hash also let the power-capped card clock higher
(1920 -> 1995 MHz), which is part of the gain.

Rejected by measurement: __noinline__ permute (-13%: code size was not the
limit, the call convention was), -maxrregcount 80/64 (flat), internal-round
unroll 1 or 22 (flat, within noise of 2).

The batch sweep found batch size irrelevant once the grid keeps one nonce
per thread (304-306 MH/s from 1M to 16M); batch stays at 1M for the
smallest stale cost. All knobs remain as -D macros and MINER_NVCC_FLAGS
so the harness can keep testing variants.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01CBgs2nSi4H2mdh8kD8vMX5

crates/engine-cuda/build.rs

@@ -18,6 +18,7 @@
 //!   MINER_NVCC_CCBIN                     host compiler for nvcc (-ccbin)
 //!   MINER_CUDA_ALLOW_UNSUPPORTED_COMPILER=1  pass -allow-unsupported-compiler
 //!   MINER_CUDA_REQUIRE=1                 fail the build instead of warning when nvcc is missing
+//!   MINER_NVCC_FLAGS                     extra nvcc arguments (whitespace-separated), e.g. -DLAIR_LAZY_ADD=0
 
 use std::env;
 use std::fs;
@@ -118,6 +119,7 @@ fn main() {
         "MINER_NVCC_CCBIN",
         "MINER_CUDA_ALLOW_UNSUPPORTED_COMPILER",
         "MINER_CUDA_REQUIRE",
+        "MINER_NVCC_FLAGS",
     ] {
         println!("cargo:rerun-if-env-changed={v}");
     }
@@ -163,6 +165,11 @@ fn main() {
     {
         cmd.arg("-allow-unsupported-compiler");
     }
+    if let Ok(extra) = env::var("MINER_NVCC_FLAGS") {
+        for a in extra.split_whitespace() {
+            cmd.arg(a);
+        }
+    }
     for sm in &sms {
         cmd.arg("--generate-code")
             .arg(format!("arch=compute_{sm},code=sm_{sm}"));

crates/engine-cuda/src/kernels/mining.cu

@@ -28,6 +28,36 @@ typedef unsigned int u32;
 // 2^64 mod P = 2^32 - 1
 #define EPS 0xFFFFFFFFull
 
+#ifndef LAIR_PTX_CARRY
+#define LAIR_PTX_CARRY 1
+#endif
+#ifndef LAIR_PTX_ACC
+#define LAIR_PTX_ACC LAIR_PTX_CARRY
+#endif
+
+#if LAIR_PTX_CARRY
+// Carry-flag arithmetic: the wrap is read from the condition code instead of
+// a compare, and folded with a multiply-add (c * EPS is exact for c in {0,1}).
+__device__ __forceinline__ u64 gf_add(u64 a, u64 b) {
+    u64 s0, c1, s1, c2;
+    asm("add.cc.u64 %0, %2, %3;\n\taddc.u64 %1, 0, 0;" : "=l"(s0), "=l"(c1) : "l"(a), "l"(b));
+    asm("add.cc.u64 %0, %2, %3;\n\taddc.u64 %1, 0, 0;" : "=l"(s1), "=l"(c2) : "l"(s0), "l"(c1 * EPS));
+    return s1 + c2 * EPS;
+}
+
+__device__ __forceinline__ u64 gf_reduce(u64 lo, u64 hi) {
+    u64 hi_hi = hi >> 32;
+    u64 hi_lo = hi & EPS;
+    u64 t0, bm;
+    // bm = -borrow: all ones when lo < hi_hi.
+    asm("sub.cc.u64 %0, %2, %3;\n\tsubc.u64 %1, 0, 0;" : "=l"(t0), "=l"(bm) : "l"(lo), "l"(hi_hi));
+    t0 -= (bm & EPS);
+    u64 t1 = hi_lo * EPS;
+    u64 t2, c;
+    asm("add.cc.u64 %0, %2, %3;\n\taddc.u64 %1, 0, 0;" : "=l"(t2), "=l"(c) : "l"(t0), "l"(t1));
+    return t2 + c * EPS;
+}
+#else
 // a + b mod P, lazy. A wrapped carry folds back as 2^64 = EPS (mod P); the
 // second fold can only be needed when the first one wrapped.
 __device__ __forceinline__ u64 gf_add(u64 a, u64 b) {
@@ -49,6 +79,7 @@ __device__ __forceinline__ u64 gf_reduce(u64 lo, u64 hi) {
     if (t2 < t0) t2 += EPS;
     return t2;
 }
+#endif
 
 __device__ __forceinline__ u64 gf_mul(u64 a, u64 b) {
     return gf_reduce(a * b, __umul64hi(a, b));
@@ -70,6 +101,126 @@ __device__ __forceinline__ u64 gf_canon(u64 a) {
     return a - (a >= P64 ? P64 : 0ull);
 }
 
+#ifndef LAIR_LAZY_ADD
+#define LAIR_LAZY_ADD 1
+#endif
+#ifndef LAIR_INT_UNROLL
+#define LAIR_INT_UNROLL 2
+#endif
+// nvcc does not macro-expand `#pragma unroll N`; stringise through _Pragma.
+#define LAIR_PRAGMA(x) _Pragma(#x)
+#define LAIR_UNROLL(n) LAIR_PRAGMA(unroll n)
+#ifndef LAIR_PTX_CARRY
+#define LAIR_PTX_CARRY 1
+#endif
+#ifndef LAIR_PTX_ACC
+#define LAIR_PTX_ACC LAIR_PTX_CARRY
+#endif
+#ifndef LAIR_NOINLINE_PERMUTE
+#define LAIR_NOINLINE_PERMUTE 0
+#endif
+
+// Deferred-carry accumulator: the value is lo + hi * 2^64 with hi small. Adds
+// cost an add-with-carry instead of two compare-and-fold steps; the folds are
+// paid once per output when the accumulator is reduced. Sound as long as
+// hi < 2^32, which a linear layer's sums (at most a few dozen terms) satisfy.
+struct Acc {
+    u64 lo;
+    u32 hi;
+};
+
+__device__ __forceinline__ Acc acc_of(u64 a) {
+    Acc r; r.lo = a; r.hi = 0u; return r;
+}
+
+#if LAIR_PTX_ACC
+__device__ __forceinline__ Acc acc_add(Acc a, u64 b) {
+    u64 s; u32 c;
+    asm("add.cc.u64 %0, %2, %3;\n\taddc.u32 %1, 0, 0;" : "=l"(s), "=r"(c) : "l"(a.lo), "l"(b));
+    a.lo = s;
+    a.hi += c;
+    return a;
+}
+
+__device__ __forceinline__ Acc acc_add2(Acc a, Acc b) {
+    u64 s; u32 c;
+    asm("add.cc.u64 %0, %2, %3;\n\taddc.u32 %1, %4, %5;" : "=l"(s), "=r"(c) : "l"(a.lo), "l"(b.lo), "r"(a.hi), "r"(b.hi));
+    a.lo = s;
+    a.hi = c;
+    return a;
+}
+#else
+__device__ __forceinline__ Acc acc_add(Acc a, u64 b) {
+    u64 s = a.lo + b;
+    a.hi += (s < a.lo) ? 1u : 0u;
+    a.lo = s;
+    return a;
+}
+
+__device__ __forceinline__ Acc acc_add2(Acc a, Acc b) {
+    u64 s = a.lo + b.lo;
+    a.hi += b.hi + ((s < a.lo) ? 1u : 0u);
+    a.lo = s;
+    return a;
+}
+#endif
+
+// lo + hi * 2^64 = lo + hi * EPS (mod P); hi * EPS < 2^64, so this is one
+// lazy gf_add.
+__device__ __forceinline__ u64 acc_reduce(Acc a) {
+    return gf_add(a.lo, (u64)a.hi * EPS);
+}
+
+// a * b + c mod P: the addend is folded into the 128-bit product before the
+// single reduction, saving a gf_add per term.
+__device__ __forceinline__ u64 gf_mul_add(u64 a, u64 b, u64 c) {
+    u64 lo = a * b;
+    u64 hi = __umul64hi(a, b);
+    u64 s = lo + c;
+    hi += (s < lo) ? 1ull : 0ull;
+    return gf_reduce(s, hi);
+}
+
+#if LAIR_LAZY_ADD
+// External linear layer with deferred carries: every output is a sum of a
+// handful of inputs, reduced once.
+__device__ __forceinline__ void ext_layer(u64 st[12]) {
+    Acc acc[12];
+#pragma unroll
+    for (int chunk = 0; chunk < 3; chunk++) {
+        int o = chunk * 4;
+        u64 x0 = st[o], x1 = st[o + 1], x2 = st[o + 2], x3 = st[o + 3];
+        Acc t01 = acc_add(acc_of(x0), x1);
+        Acc t23 = acc_add(acc_of(x2), x3);
+        Acc t0123 = acc_add2(t01, t23);
+        Acc t01123 = acc_add(t0123, x1);
+        Acc t01233 = acc_add(t0123, x3);
+        acc[o + 3] = acc_add(acc_add(t01233, x0), x0);
+        acc[o + 1] = acc_add(acc_add(t01123, x2), x2);
+        acc[o] = acc_add2(t01123, t01);
+        acc[o + 2] = acc_add2(t01233, t23);
+    }
+    Acc sums[4];
+#pragma unroll
+    for (int k = 0; k < 4; k++) {
+        sums[k] = acc_add2(acc_add2(acc[k], acc[k + 4]), acc[k + 8]);
+    }
+#pragma unroll
+    for (int i = 0; i < 12; i++) {
+        st[i] = acc_reduce(acc_add2(acc[i], sums[i & 3]));
+    }
+}
+
+// Internal linear layer: one deferred sum, folded into each diagonal multiply.
+__device__ __forceinline__ void int_layer(u64 st[12]) {
+    Acc s = acc_of(st[0]);
+#pragma unroll
+    for (int i = 1; i < 12; i++) s = acc_add(s, st[i]);
+    u64 sum = acc_reduce(s);
+#pragma unroll
+    for (int i = 0; i < 12; i++) st[i] = gf_mul_add(st[i], MDS_DIAG[i], sum);
+}
+#else
 // External linear layer: 4x4 MDS on each chunk, then circulant sums.
 __device__ __forceinline__ void ext_layer(u64 st[12]) {
 #pragma unroll
@@ -105,8 +256,15 @@ __device__ __forceinline__ void int_layer(u64 st[12]) {
 #pragma unroll
     for (int i = 0; i < 12; i++) st[i] = gf_add(gf_mul(st[i], MDS_DIAG[i]), sum);
 }
+#endif
 
+#if LAIR_NOINLINE_PERMUTE
+// One copy of the (fully unrolled) permutation instead of three inlined ones:
+// 17k instructions instead of 51k, for instruction-cache pressure.
+__device__ __noinline__ void permute(u64 st[12]) {
+#else
 __device__ __forceinline__ void permute(u64 st[12]) {
+#endif
     ext_layer(st);
 #pragma unroll
     for (int r = 0; r < N_EXTERNAL_HALF; r++) {
@@ -114,7 +272,7 @@ __device__ __forceinline__ void permute(u64 st[12]) {
         for (int i = 0; i < 12; i++) st[i] = gf_sbox(gf_add(st[i], RC_INITIAL[r][i]));
         ext_layer(st);
     }
-#pragma unroll 2
+    LAIR_UNROLL(LAIR_INT_UNROLL)
     for (int r = 0; r < N_INTERNAL; r++) {
         st[0] = gf_sbox(gf_add(st[0], RC_INTERNAL[r]));
         int_layer(st);

crates/engine-cuda/src/lib.rs

@@ -29,7 +29,10 @@ const KERNEL_ID: &str = "cuda";
 const THREADS_PER_BLOCK: u32 = 256;
 /// Threads per SM the grid is sized for (override: MINER_CUDA_THREADS_PER_SM).
 /// Sets how many nonces each thread loops over for a given batch; tuned with #2.
-const THREADS_PER_SM_DEFAULT: u32 = 8192;
+// Measured on the 4090 (#3): 8192 caps the grid at ~1M threads, so batches
+// above 1M loop nonces per thread and lose up to 4%; 32768 keeps one nonce per
+// thread up to 4M and is flat across batch sizes at ~305 MH/s.
+const THREADS_PER_SM_DEFAULT: u32 = 32768;
 
 fn threads_per_sm() -> u32 {
     std::env::var("MINER_CUDA_THREADS_PER_SM")