Merge pull request 'deploy: require nvcc at build time' (#32) from ci/deploy-require-cuda into main

The CUDA build script warns and embeds an empty kernel when nvcc is
absent, which is right for the rust runner's test job and wrong for a
deploy build. MINER_CUDA_REQUIRE=1 turns that into a build failure
before any host is touched; the release workflow already sets it.

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

.gitea/workflows/deploy.yaml

@@ -61,6 +61,11 @@ jobs:
           # Embedded by crates/miner-cli/build.rs into --version; validate below
           # asserts the host runs exactly this commit.
           MINER_BUILD_SHA: ${{ github.sha }}
+          # Fail the build here if nvcc is missing rather than shipping a
+          # binary that falls back to wgpu and failing validate on the hosts.
+          # (The same warning in ci.yml's test job is expected: the `rust`
+          # runner has no nvcc and only lints and tests.)
+          MINER_CUDA_REQUIRE: "1"
         run: |
           set -euo pipefail
           cargo build --release --locked -p miner-cli

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

@@ -353,9 +353,24 @@ __device__ __forceinline__ u64 gf_canon(u64 a) {
 #else
 #define LAIR_LOOP_LAUNCH_BOUNDS __launch_bounds__(LAIR_LOOP_TPB)
 #endif
+// Loop kernel: take the diagonal from the launch parameters instead of the
+// __constant__ table (see MiningUniforms::diag).
+#ifndef LAIR_LOOP_PARAM_DIAG
+#define LAIR_LOOP_PARAM_DIAG 1
+#endif
 // Internal rounds per loop iteration in the loop kernel (1 or 2). One
 // measured 2105 MH/s against 1976 for two on beast: the pair doubles the
 // live state across the iteration for no scheduling gain.
+//
+// Where the loop kernel's internal round (239 instructions at 128
+// registers) still exceeds the quanpool kernel's 199 (2026-09-14, static
+// screen on sm_120): 12 reloads of the diagonal constants per round, which
+// ptxas issues even with registers to spare (13 LDC at 192 registers too;
+// hoisting them in C or passing them as 32-bit halves compiles to identical
+// SASS); ~12 in the product carry assembly; ~8 in the EPS multiply of the
+// reductions; ~5 in the sum-plus-constant fold; ~6 moves. Five multiply-add
+// formulations (LAIR_MULADD) and two constant deliveries were screened and
+// none beat the shipped form.
 #ifndef LAIR_LOOP_INT_PAIR
 #define LAIR_LOOP_INT_PAIR 1
 #endif
@@ -480,8 +495,15 @@ __device__ __forceinline__ u64 acc_reduce(Acc a) {
     return r;
 }
 
-// a * b + c mod P. The addend rides in the two `mad.wide` partial products
-// (a0*b0 + c_lo and a1*b0 + c_hi both fit in 64 bits), so it costs nothing.
+// a * b + c mod P. Formulations are screened by the loop kernel's internal
+// round body (static SASS count, target quanpool's 199) before any card
+// time: LAIR_MULADD selects one. 0 is the shipped form.
+#ifndef LAIR_MULADD
+#define LAIR_MULADD 0
+#endif
+#if LAIR_MULADD == 0
+// The addend rides in the two `mad.wide` partial products (a0*b0 + c_lo and
+// a1*b0 + c_hi both fit in 64 bits), so it costs nothing.
 __device__ __forceinline__ u64 gf_mul_add(u64 a, u64 b, u64 c) {
     u64 r;
     u32 e = LAIR_EPS32;
@@ -510,6 +532,120 @@ __device__ __forceinline__ u64 gf_mul_add(u64 a, u64 b, u64 c) {
         : "=l"(r) : "l"(a), "l"(b), "l"(c0), "l"(c1), "r"(e));
     return r;
 }
+#elif LAIR_MULADD == 1
+// Whole 64-bit addend folded into a0*b0 with a carry-out (mad.lo.cc /
+// madc.hi.cc); the middle products accumulated the same way. Measured worse
+// than 0 on 2026-09-14 (loop internal body 268 against 254).
+__device__ __forceinline__ u64 gf_mul_add(u64 a, u64 b, u64 c) {
+    u64 r;
+    u32 e = LAIR_EPS32;
+    asm("{\n\t"
+        ".reg .b32 a0, a1, b0, b1, s0, s1, m0, m1, cw, c2, ll, lh, hl, hh, c;\n\t"
+        ".reg .b64 m;\n\t"
+        "mov.b64 {a0, a1}, %1;\n\t"
+        "mov.b64 {b0, b1}, %2;\n\t"
+        "mov.b64 {s0, s1}, %3;\n\t"
+        "mul.wide.u32 m, a1, b0;\n\t"
+        "mov.b64 {m0, m1}, m;\n\t"
+        "mad.lo.cc.u32 m0, a0, b1, m0;\n\t"
+        "madc.hi.cc.u32 m1, a0, b1, m1;\n\t"
+        "addc.u32 cw, 0, 0;\n\t"
+        "mad.lo.cc.u32 ll, a0, b0, s0;\n\t"
+        "madc.hi.cc.u32 lh, a0, b0, s1;\n\t"
+        "addc.u32 c2, 0, 0;\n\t"
+        "add.cc.u32 lh, lh, m0;\n\t"
+        "addc.cc.u32 hl, m1, c2;\n\t"
+        "addc.u32 hh, cw, 0;\n\t"
+        "mad.lo.cc.u32 hl, a1, b1, hl;\n\t"
+        "madc.hi.u32 hh, a1, b1, hh;\n\t"
+        LAIR_REDUCE_PTX("%4")
+        "mov.b64 %0, {ll, lh};\n\t"
+        "}"
+        : "=l"(r) : "l"(a), "l"(b), "l"(c), "r"(e));
+    return r;
+}
+#elif LAIR_MULADD == 2
+// nvcc's own 128-bit product (a*b, __umul64hi), the addend folded in C, the
+// reduction as the PTX chain over C temporaries.
+__device__ __forceinline__ u64 gf_mul_add(u64 a, u64 b, u64 c) {
+    u64 lo = a * b + c;
+    u64 hi = __umul64hi(a, b) + (lo < c ? 1ull : 0ull);
+    u32 ll = (u32)lo, lh = (u32)(lo >> 32), hl = (u32)hi, hh = (u32)(hi >> 32), cc;
+    u32 e = LAIR_EPS32;
+    asm("mad.lo.cc.u32 %0, %3, %5, %0;\n\tmadc.hi.cc.u32 %1, %3, %5, %1;\n\taddc.u32 %2, %4, 0;\n\taddc.u32 %1, %1, 0;\n\tsub.cc.u32 %0, %0, %2;\n\tsubc.u32 %1, %1, 0;"
+        : "+r"(ll), "+r"(lh), "=r"(cc) : "r"(hl), "r"(hh), "r"(e));
+    return ((u64)lh << 32) | ll;
+}
+#elif LAIR_MULADD == 3
+// The shipped product chain (gf_mul's), then the addend added to the 128-bit
+// product as a 32-bit carry chain before the reduction.
+__device__ __forceinline__ u64 gf_mul_add(u64 a, u64 b, u64 c) {
+    u64 r;
+    u32 e = LAIR_EPS32;
+    asm("{\n\t"
+        ".reg .b32 a0, a1, b0, b1, s0, s1, p0l, p0h, m0, m1, cw, ll, lh, hl, hh, c;\n\t"
+        ".reg .b64 p0, m, m2, p3, t, hi;\n\t"
+        "mov.b64 {a0, a1}, %1;\n\t"
+        "mov.b64 {b0, b1}, %2;\n\t"
+        "mov.b64 {s0, s1}, %3;\n\t"
+        "mul.wide.u32 p0, a0, b0;\n\t"
+        "mul.wide.u32 m, a1, b0;\n\t"
+        "mul.wide.u32 m2, a0, b1;\n\t"
+        "add.cc.u64 m, m, m2;\n\t"
+        "addc.u32 cw, 0, 0;\n\t"
+        "mov.b64 {p0l, p0h}, p0;\n\t"
+        "mov.b64 {m0, m1}, m;\n\t"
+        "mov.b64 t, {m1, cw};\n\t"
+        "mul.wide.u32 p3, a1, b1;\n\t"
+        "add.cc.u32 lh, p0h, m0;\n\t"
+        "addc.u64 hi, p3, t;\n\t"
+        "mov.b64 {hl, hh}, hi;\n\t"
+        "add.cc.u32 ll, p0l, s0;\n\t"
+        "addc.cc.u32 lh, lh, s1;\n\t"
+        "addc.cc.u32 hl, hl, 0;\n\t"
+        "addc.u32 hh, hh, 0;\n\t"
+        LAIR_REDUCE_PTX("%4")
+        "mov.b64 %0, {ll, lh};\n\t"
+        "}"
+        : "=l"(r) : "l"(a), "l"(b), "l"(c), "r"(e));
+    return r;
+}
+#elif LAIR_MULADD == 4
+// Halves of the addend as mad.wide addends (as 0) but with the product's
+// carry assembly in 32-bit chains instead of add.cc.u64.
+__device__ __forceinline__ u64 gf_mul_add(u64 a, u64 b, u64 c) {
+    u64 r;
+    u32 e = LAIR_EPS32;
+    u64 c0 = c & EPS, c1 = c >> 32;
+    asm("{\n\t"
+        ".reg .b32 a0, a1, b0, b1, p0l, p0h, p3l, p3h, m0, m1, m2l, m2h, cw, ll, lh, hl, hh, c;\n\t"
+        ".reg .b64 p0, m, m2, p3;\n\t"
+        "mov.b64 {a0, a1}, %1;\n\t"
+        "mov.b64 {b0, b1}, %2;\n\t"
+        "mad.wide.u32 p0, a0, b0, %3;\n\t"
+        "mad.wide.u32 m, a1, b0, %4;\n\t"
+        "mul.wide.u32 m2, a0, b1;\n\t"
+        "mul.wide.u32 p3, a1, b1;\n\t"
+        "mov.b64 {p0l, p0h}, p0;\n\t"
+        "mov.b64 {p3l, p3h}, p3;\n\t"
+        "mov.b64 {m0, m1}, m;\n\t"
+        "mov.b64 {m2l, m2h}, m2;\n\t"
+        "add.cc.u32 m0, m0, m2l;\n\t"
+        "addc.cc.u32 m1, m1, m2h;\n\t"
+        "addc.u32 cw, p3h, 0;\n\t"
+        "add.cc.u32 lh, p0h, m0;\n\t"
+        "addc.cc.u32 hl, p3l, m1;\n\t"
+        "addc.u32 hh, cw, 0;\n\t"
+        "mov.b32 ll, p0l;\n\t"
+        LAIR_REDUCE_PTX("%5")
+        "mov.b64 %0, {ll, lh};\n\t"
+        "}"
+        : "=l"(r) : "l"(a), "l"(b), "l"(c0), "l"(c1), "r"(e));
+    return r;
+}
+#else
+#error "unknown LAIR_MULADD"
+#endif
 
 // a * b + (c.lo + c.hi * 2^64) mod P: reduce the addend first (4 ops), then
 // fold it into the product for free.
@@ -597,7 +733,12 @@ __device__ __forceinline__ u64 ext_layer_out0(const u64 st[12]) {
 // 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_d(u64 st[12], u64 rc0, bool has_rc0, const u64* rc_row, const u64* d);
 __device__ __forceinline__ void int_layer_rc(u64 st[12], u64 rc0, bool has_rc0, const u64* rc_row) {
+    int_layer_rc_d(st, rc0, has_rc0, rc_row, MDS_DIAG);
+}
+#define LAIR_DMA(x, i, add) gf_mul_add((x), d[i], (add))
+__device__ __forceinline__ void int_layer_rc_d(u64 st[12], u64 rc0, bool has_rc0, const u64* rc_row, const u64* d) {
     Acc s = acc_of(st[0]);
 #pragma unroll
     for (int i = 1; i < 12; i++) s = acc_add(s, st[i]);
@@ -606,15 +747,15 @@ __device__ __forceinline__ void int_layer_rc(u64 st[12], u64 rc0, bool has_rc0,
         // The constant joins the unreduced sum: one more limb add, then the
         // same single reduction.
 #pragma unroll
-        for (int i = 0; i < 12; i++) st[i] = gf_mul_add_acc(st[i], MDS_DIAG[i], acc_add(s, rc_row[i]));
+        for (int i = 0; i < 12; i++) st[i] = LAIR_DMA(st[i], i, acc_reduce(acc_add(s, rc_row[i])));
     } else {
         if (has_rc0) {
-            st[0] = gf_mul_add_acc(st[0], MDS_DIAG[0], acc_add(s, rc0));
+            st[0] = LAIR_DMA(st[0], 0, acc_reduce(acc_add(s, rc0)));
         } else {
-            st[0] = gf_mul_add(st[0], MDS_DIAG[0], sum);
+            st[0] = LAIR_DMA(st[0], 0, sum);
         }
 #pragma unroll
-        for (int i = 1; i < 12; i++) st[i] = gf_mul_add(st[i], MDS_DIAG[i], sum);
+        for (int i = 1; i < 12; i++) st[i] = LAIR_DMA(st[i], i, sum);
     }
 }
 
@@ -740,7 +881,7 @@ __device__ __forceinline__ void permute(u64 st[12], bool skip_last_ext = false,
 // each round stay unrolled so the state lives in registers.
 #if LAIR_LAZY_ADD && LAIR_FOLD_RC
 static_assert(N_INTERNAL % 2 == 0, "the paired internal loop assumes an even round count");
-__device__ __forceinline__ void permute_loop(u64 st[12], bool skip_last_ext, bool skip_first_ext) {
+__device__ __forceinline__ void permute_loop(u64 st[12], bool skip_last_ext, bool skip_first_ext, const u64* diag) {
     if (!skip_first_ext) ext_layer_rc(st, RC_INITIAL[0]);
 #pragma unroll 1
     for (int r = 0; r + 1 < N_EXTERNAL_HALF; r++) {
@@ -752,25 +893,31 @@ __device__ __forceinline__ void permute_loop(u64 st[12], bool skip_last_ext, boo
     for (int i = 0; i < 12; i++) st[i] = gf_sbox(st[i]);
     ext_layer(st);
     st[0] = gf_add(st[0], RC_INTERNAL[0]);
+#if LAIR_LOOP_PARAM_DIAG
+    const u64* d = diag;
+#else
+    const u64* d = MDS_DIAG;
+    (void)diag;
+#endif
 #if LAIR_LOOP_INT_PAIR == 2
 #pragma unroll 1
     for (int r = 0; r + 2 < N_INTERNAL; r += 2) {
         st[0] = gf_sbox(st[0]);
-        int_layer_rc(st, RC_INTERNAL[r + 1], true, nullptr);
+        int_layer_rc_d(st, RC_INTERNAL[r + 1], true, nullptr, d);
         st[0] = gf_sbox(st[0]);
-        int_layer_rc(st, RC_INTERNAL[r + 2], true, nullptr);
+        int_layer_rc_d(st, RC_INTERNAL[r + 2], true, nullptr, d);
     }
     st[0] = gf_sbox(st[0]);
-    int_layer_rc(st, RC_INTERNAL[N_INTERNAL - 1], true, nullptr);
+    int_layer_rc_d(st, RC_INTERNAL[N_INTERNAL - 1], true, nullptr, d);
 #else
 #pragma unroll 1
     for (int r = 0; r + 1 < N_INTERNAL; r++) {
         st[0] = gf_sbox(st[0]);
-        int_layer_rc(st, RC_INTERNAL[r + 1], true, nullptr);
+        int_layer_rc_d(st, RC_INTERNAL[r + 1], true, nullptr, d);
     }
 #endif
     st[0] = gf_sbox(st[0]);
-    int_layer_rc(st, 0ull, false, RC_TERMINAL[0]);
+    int_layer_rc_d(st, 0ull, false, RC_TERMINAL[0], d);
 #pragma unroll 1
     for (int r = 0; r + 1 < N_EXTERNAL_HALF; r++) {
 #pragma unroll
@@ -782,14 +929,15 @@ __device__ __forceinline__ void permute_loop(u64 st[12], bool skip_last_ext, boo
     if (!skip_last_ext) ext_layer(st);
 }
 #else
-__device__ __forceinline__ void permute_loop(u64 st[12], bool skip_last_ext, bool skip_first_ext) {
+__device__ __forceinline__ void permute_loop(u64 st[12], bool skip_last_ext, bool skip_first_ext, const u64* diag) {
+    (void)diag;
     permute(st, skip_last_ext, skip_first_ext);
 }
 #endif
 
 template <bool LOOP>
-__device__ __forceinline__ void run_permute(u64 st[12], bool skip_last_ext = false, bool skip_first_ext = false) {
-    if (LOOP) permute_loop(st, skip_last_ext, skip_first_ext);
+__device__ __forceinline__ void run_permute(u64 st[12], const u64* diag, bool skip_last_ext = false, bool skip_first_ext = false) {
+    if (LOOP) permute_loop(st, skip_last_ext, skip_first_ext, diag);
     else permute(st, skip_last_ext, skip_first_ext);
 }
 
@@ -817,8 +965,12 @@ struct MiningUniforms {
     // State after the first external layer + its round constant for the
     // batch's first nonce (host: first_layer_after_absorb).
     u64 layer0_base[12];
+    // The internal layer's diagonal, same values as MDS_DIAG, for the loop
+    // kernel: a parameter-bank operand costs no load and no register, where
+    // ptxas reloaded the __constant__ table every round.
+    u64 diag[12];
 };
-static_assert(sizeof(MiningUniforms) == 320, "MiningUniforms must match the host layout");
+static_assert(sizeof(MiningUniforms) == 416, "MiningUniforms must match the host layout");
 
 #if LAIR_NONCE_DIR
 // Column 7 of the external matrix circ(2*M4, M4, M4): what the layer adds per
@@ -883,13 +1035,13 @@ __device__ __forceinline__ void mine_nonces(u32* __restrict__ results,
             for (int i = 0; i < 8; i++) st[i] = gf_add(st[i], (u64)bswap32(current_nonce[7 - i]));
             ext_layer_rc(st, RC_INITIAL[0]);
         }
-        run_permute<LOOP>(st, false, true);
+        run_permute<LOOP>(st, uni.diag, false, true);
 #else
 #pragma unroll
         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]));
-        run_permute<LOOP>(st);
+        run_permute<LOOP>(st, uni.diag);
 #endif
         st[0] = gf_add(st[0], 1ull);
         st[1] = gf_add(st[1], 1ull);
@@ -897,7 +1049,7 @@ __device__ __forceinline__ void mine_nonces(u32* __restrict__ results,
         // Element 0 of the final state is the hash's top 64 bits. Compute it
         // alone, and only nonces that do not already exceed the target there
         // pay for the other eleven outputs of the last linear layer.
-        run_permute<LOOP>(st, true);
+        run_permute<LOOP>(st, uni.diag, true);
         {
             u64 c0 = gf_canon(ext_layer_out0(st));
             u32 h0 = bswap32((u32)(c0 & EPS));
@@ -907,7 +1059,7 @@ __device__ __forceinline__ void mine_nonces(u32* __restrict__ results,
         }
         ext_layer(st);
 #else
-        run_permute<LOOP>(st);
+        run_permute<LOOP>(st, uni.diag);
 #endif
 
         // First squeeze: most significant 256 bits of the hash.
@@ -930,7 +1082,7 @@ __device__ __forceinline__ void mine_nonces(u32* __restrict__ results,
         u32 hash_le[16];
 #pragma unroll
         for (int i = 0; i < 8; i++) hash_le[15 - i] = bswap32(first[i]);
-        run_permute<LOOP>(st);
+        run_permute<LOOP>(st, uni.diag);
 #pragma unroll
         for (int i = 0; i < 4; i++) {
             u64 c = gf_canon(st[i]);

crates/engine-cuda/src/lib.rs

@@ -173,9 +173,13 @@ struct MiningUniforms {
     /// derives each nonce's state from this by one scalar times a fixed
     /// small-integer column instead of recomputing the layer (quantus/miner#3).
     layer0_base: [u64; 12],
+    /// lair: the internal layer's diagonal, for the loop kernel (#27). In the
+    /// parameter bank ptxas reads it straight from the instruction operand;
+    /// from `__constant__` it reloaded all twelve values every round.
+    diag: [u64; 12],
 }
 
-// SAFETY: `#[repr(C)]` plain-old-data with no padding (12 x u64, 32 x u32, 12 x u64),
+// SAFETY: `#[repr(C)]` plain-old-data with no padding (12 x u64, 32 x u32, 24 x u64),
 // 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 {}
@@ -183,7 +187,7 @@ 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>() == 320);
+const _: () = assert!(std::mem::size_of::<MiningUniforms>() == 416);
 
 pub struct CudaEngine {
     engine_id: usize,
@@ -539,6 +543,7 @@ impl CudaEngine {
             start_nonce: start_u32s,
             target: res.target_u32s,
             layer0_base,
+            diag: qp_poseidon_constants::POSEIDON2_MATRIX_DIAG_12_RAW,
         };
 
         let stream = res.stream.clone();