Merge pull request 'engine-cuda: deferred carries, carry-flag arithmetic, one nonce per thread' (#17) from cuda/tuning into main

.gitea/workflows/bench.yaml

@@ -38,6 +38,7 @@ on:
   pull_request:
     paths:
       - crates/engine-gpu/**
+      - crates/engine-cuda/**
       - crates/engine-cpu/**
       - crates/pow-core/**
       - crates/miner-service/**

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,41 @@ 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
+// 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.
+#ifndef LAIR_FOLD_RC
+#define LAIR_FOLD_RC 1
+#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 +84,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 +106,159 @@ __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
+// 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.
+#ifndef LAIR_FOLD_RC
+#define LAIR_FOLD_RC 1
+#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);
+}
+
+// a * b + (c.lo + c.hi * 2^64) mod P, for a deferred-carry addend.
+__device__ __forceinline__ u64 gf_mul_add_acc(u64 a, u64 b, Acc c) {
+    u64 lo = a * b;
+    u64 hi = __umul64hi(a, b);
+    u64 s = lo + c.lo;
+    hi += ((s < lo) ? 1ull : 0ull) + (u64)c.hi;
+    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.
+// `rc` is the next round's constant row to fold into the outputs, or nullptr.
+__device__ __forceinline__ void ext_layer_rc(u64 st[12], const u64* rc) {
+    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++) {
+        Acc o = acc_add2(acc[i], sums[i & 3]);
+        if (rc != nullptr) o = acc_add(o, rc[i]);
+        st[i] = acc_reduce(o);
+    }
+}
+
+__device__ __forceinline__ void ext_layer(u64 st[12]) { ext_layer_rc(st, nullptr); }
+
+// Internal linear layer: one deferred sum, folded into each diagonal multiply.
+// `rc0` is the next internal round's constant for element 0, folded into that
+// element's multiply-add; `rc_row` a full row (the first terminal round's).
+__device__ __forceinline__ void int_layer_rc(u64 st[12], u64 rc0, bool has_rc0, const u64* rc_row) {
+    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);
+    if (rc_row != nullptr) {
+#pragma unroll
+        for (int i = 0; i < 12; i++) st[i] = gf_mul_add_acc(st[i], MDS_DIAG[i], acc_add(acc_of(sum), rc_row[i]));
+    } else {
+        if (has_rc0) {
+            st[0] = gf_mul_add_acc(st[0], MDS_DIAG[0], acc_add(acc_of(sum), rc0));
+        } else {
+            st[0] = gf_mul_add(st[0], MDS_DIAG[0], sum);
+        }
+#pragma unroll
+        for (int i = 1; i < 12; i++) st[i] = gf_mul_add(st[i], MDS_DIAG[i], sum);
+    }
+}
+
+__device__ __forceinline__ void int_layer(u64 st[12]) { int_layer_rc(st, 0ull, false, nullptr); }
+#else
 // External linear layer: 4x4 MDS on each chunk, then circulant sums.
 __device__ __forceinline__ void ext_layer(u64 st[12]) {
 #pragma unroll
@@ -105,8 +294,49 @@ __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
+#if LAIR_LAZY_ADD && LAIR_FOLD_RC
+    // Each linear layer folds the constant of the round that follows it.
+    ext_layer_rc(st, RC_INITIAL[0]);
+#pragma unroll
+    for (int r = 0; r < N_EXTERNAL_HALF; r++) {
+#pragma unroll
+        for (int i = 0; i < 12; i++) st[i] = gf_sbox(st[i]);
+        if (r + 1 < N_EXTERNAL_HALF) {
+            ext_layer_rc(st, RC_INITIAL[r + 1]);
+        } else {
+            ext_layer(st);
+            st[0] = gf_add(st[0], RC_INTERNAL[0]);
+        }
+    }
+    LAIR_UNROLL(LAIR_INT_UNROLL)
+    for (int r = 0; r < N_INTERNAL; r++) {
+        st[0] = gf_sbox(st[0]);
+        if (r + 1 < N_INTERNAL) {
+            int_layer_rc(st, RC_INTERNAL[r + 1], true, nullptr);
+        } else {
+            int_layer_rc(st, 0ull, false, RC_TERMINAL[0]);
+        }
+    }
+#pragma unroll
+    for (int r = 0; r < N_EXTERNAL_HALF; r++) {
+#pragma unroll
+        for (int i = 0; i < 12; i++) st[i] = gf_sbox(st[i]);
+        if (r + 1 < N_EXTERNAL_HALF) {
+            ext_layer_rc(st, RC_TERMINAL[r + 1]);
+        } else {
+            ext_layer(st);
+        }
+    }
+#else
     ext_layer(st);
 #pragma unroll
     for (int r = 0; r < N_EXTERNAL_HALF; r++) {
@@ -114,7 +344,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);
@@ -125,6 +355,7 @@ __device__ __forceinline__ void permute(u64 st[12]) {
         for (int i = 0; i < 12; i++) st[i] = gf_sbox(gf_add(st[i], RC_TERMINAL[r][i]));
         ext_layer(st);
     }
+#endif
 }
 
 __device__ __forceinline__ u32 bswap32(u32 v) {

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")