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
@@ -88,17 +93,24 @@ jobs:
           # `kernel` is the kernel id validate expects on the per-device metric:
           # cuda (engine-cuda, #3) on every NVIDIA host once the build carries
           # the fat binary; u64 would mean the miner silently fell back to wgpu.
+          # `cuda_kernel` selects the kernel entry (MINER_CUDA_KERNEL, #29):
+          # unrolled is the default; loop measured +1.7% on the 4090 and
+          # -1.8% / -3.4% on the 3060 / 5090 (2026-09-14), so only benjy runs
+          # it. `kernel` must be the id that choice reports: cuda or cuda-loop.
           - host: benjy.hanzalova.internal
             node: bob.hanzalova.internal
             gpu_devices: "1"        # 1x RTX 4090 (sm_89); reference card for #2
-            kernel: cuda
+            cuda_kernel: loop
+            kernel: cuda-loop
           - host: quadbrat.hanzalova.internal
             node: bob.hanzalova.internal
             gpu_devices: "1"        # 1x RTX 3060 (sm_86)
+            cuda_kernel: unrolled
             kernel: cuda
           - host: beast.hanzalova.internal
             node: bob.hanzalova.internal
             gpu_devices: "2"        # 2x RTX 5090 (sm_120)
+            cuda_kernel: unrolled
             kernel: cuda
     steps:
       - uses: actions/checkout@v4
@@ -146,6 +158,7 @@ jobs:
         if: ${{ github.event.inputs.mode != 'validate' }}
         env:
           GPU_DEVICES: ${{ matrix.gpu_devices }}
+          CUDA_KERNEL: ${{ matrix.cuda_kernel }}
         run: |
           set -euo pipefail
           SSHOPTS="-i $HOME/.ssh/id_gitea_ci -o StrictHostKeyChecking=accept-new"
@@ -210,6 +223,7 @@ jobs:
           t = pathlib.Path("deploy/miner.env.tmpl").read_text()
           t = t.replace("{{QUANTUS_NODE_ADDR}}", os.environ["NODE_ADDR"])
           t = t.replace("{{QUANTUS_GPU_DEVICES}}", os.environ["GPU_DEVICES"])
+          t = t.replace("{{MINER_CUDA_KERNEL}}", os.environ["CUDA_KERNEL"])
           pathlib.Path("miner.env").write_text(t)
           PY
           push --chown=root:quantus-miner --chmod=F0640 \

README.md

@@ -15,7 +15,7 @@ enforced power cap, memory clock locked at 810 MHz, CUDA 13.0 driver 580:
 | card | power cap | this fork | upstream `--cuda-gpu` | quantusminer.com qpow-cuda 1.0.7 | quanpool-miner 6.2.4 (5% fee) |
 |---|---|---|---|---|---|
 | RTX 5090 | 400 W | **1104 MH/s** | 944 | 1171 | 1234 (1172 after fee) |
-| RTX 4090 | 250 W | **711 MH/s** | | | |
+| RTX 4090 | 250 W | **727 MH/s** (loop kernel) | | | |
 | RTX 3060 | 130 W | **141 MH/s** | | | |
 
 The 5090 column was measured on the same card in the same session
@@ -23,6 +23,12 @@ The 5090 column was measured on the same card in the same session
 the pool kernels get their remaining edge, and what has been tried, is
 written up in [issue #27](https://git.lair.cafe/quantus/miner/issues/27).
 
+Since v4.0.2-lair.2 the binary carries two kernels. The default (`unrolled`)
+is fastest on the RTX 5090 and 3060; RTX 40-series cards measure 1.7% better
+on the rolled-loop one, selected with `MINER_CUDA_KERNEL=loop` in the
+environment (it reports as kernel `cuda-loop` on the metrics). Both are
+bit-exact with the CPU reference; try each with `quantus-bench` on your card.
+
 **Binaries** are on the [releases page](https://git.lair.cafe/quantus/miner/releases):
 one Linux x86_64 `quantus-miner` carrying cubins for sm_86 (RTX 30), sm_89
 (RTX 40) and sm_120 (RTX 50) plus PTX for anything newer, built on Fedora 43

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

@@ -333,6 +333,47 @@ __device__ __forceinline__ u64 gf_canon(u64 a) {
 #else
 #define LAIR_LAUNCH_BOUNDS __launch_bounds__(LAIR_TPB)
 #endif
+// lair: the loop-structured kernel (`mining_loop`, quantus/miner#27). Same
+// arithmetic and same host contract as `mining_main`, but the round loops
+// are rolled with the constants fetched from constant memory by index, the
+// way the fastest closed kernel is built, so the register footprint can drop
+// and more warps fit per SM. Selected at run time (MINER_CUDA_KERNEL=loop);
+// the unrolled kernel stays the default until this one measures better.
+// 512 measured best for this kernel on beast (2131 MH/s against 2095 at
+// 256 and 2099 at 128 x 4 blocks); the host picks the same block size when
+// the loop kernel is selected.
+#ifndef LAIR_LOOP_TPB
+#define LAIR_LOOP_TPB 512
+#endif
+#ifndef LAIR_LOOP_MINBLOCKS
+#define LAIR_LOOP_MINBLOCKS 0
+#endif
+#if LAIR_LOOP_MINBLOCKS
+#define LAIR_LOOP_LAUNCH_BOUNDS __launch_bounds__(LAIR_LOOP_TPB, LAIR_LOOP_MINBLOCKS)
+#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
 // Decide most nonces on element 0 of the final state before computing the
 // rest of the last linear layer: at mainnet difficulty the top ~47 bits of
 // the hash must be zero, and they live in that element.
@@ -454,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;
@@ -484,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.
@@ -571,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]);
@@ -580,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);
     }
 }
 
@@ -708,6 +875,72 @@ __device__ __forceinline__ void permute(u64 st[12], bool skip_last_ext = false,
 #endif
 }
 
+// lair: rolled-loop twin of permute for the loop kernel. Round-constant
+// rows are addressed by the loop variable, so they are constant-memory loads
+// with a uniform offset instead of immediates; the per-element loops inside
+// 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, 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++) {
+#pragma unroll
+        for (int i = 0; i < 12; i++) st[i] = gf_sbox(st[i]);
+        ext_layer_rc(st, RC_INITIAL[r + 1]);
+    }
+#pragma unroll
+    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_d(st, RC_INTERNAL[r + 1], true, nullptr, d);
+        st[0] = gf_sbox(st[0]);
+        int_layer_rc_d(st, RC_INTERNAL[r + 2], true, nullptr, d);
+    }
+    st[0] = gf_sbox(st[0]);
+    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_d(st, RC_INTERNAL[r + 1], true, nullptr, d);
+    }
+#endif
+    st[0] = gf_sbox(st[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
+        for (int i = 0; i < 12; i++) st[i] = gf_sbox(st[i]);
+        ext_layer_rc(st, RC_TERMINAL[r + 1]);
+    }
+#pragma unroll
+    for (int i = 0; i < 12; i++) st[i] = gf_sbox(st[i]);
+    if (!skip_last_ext) ext_layer(st);
+}
+#else
+__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], 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);
+}
+
 __device__ __forceinline__ u32 bswap32(u32 v) {
     return __byte_perm(v, 0, 0x0123);
 }
@@ -732,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
@@ -742,12 +979,14 @@ static_assert(sizeof(MiningUniforms) == 320, "MiningUniforms must match the host
 __device__ __constant__ u64 LAIR_DIR7[12] = {1, 1, 3, 2, 2, 2, 6, 4, 1, 1, 3, 2};
 #endif
 
-extern "C" __global__ void LAIR_LAUNCH_BOUNDS
-mining_main(u32* __restrict__ results,
-            const MiningUniforms uni,
-            u32 total_threads,
-            u32 nonces_per_thread,
-            u32 total_nonces)
+// The per-nonce search, shared by both kernel entries; LOOP picks the
+// permutation implementation at compile time.
+template <bool LOOP>
+__device__ __forceinline__ void mine_nonces(u32* __restrict__ results,
+                                            const MiningUniforms& uni,
+                                            u32 total_threads,
+                                            u32 nonces_per_thread,
+                                            u32 total_nonces)
 {
     if (*((volatile u32*)results) != 0u) return;
     u32 tid = blockIdx.x * blockDim.x + threadIdx.x;
@@ -796,13 +1035,13 @@ mining_main(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]);
         }
-        permute(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]));
-        permute(st);
+        run_permute<LOOP>(st, uni.diag);
 #endif
         st[0] = gf_add(st[0], 1ull);
         st[1] = gf_add(st[1], 1ull);
@@ -810,7 +1049,7 @@ mining_main(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.
-        permute(st, true);
+        run_permute<LOOP>(st, uni.diag, true);
         {
             u64 c0 = gf_canon(ext_layer_out0(st));
             u32 h0 = bswap32((u32)(c0 & EPS));
@@ -820,7 +1059,7 @@ mining_main(u32* __restrict__ results,
         }
         ext_layer(st);
 #else
-        permute(st);
+        run_permute<LOOP>(st, uni.diag);
 #endif
 
         // First squeeze: most significant 256 bits of the hash.
@@ -843,7 +1082,7 @@ mining_main(u32* __restrict__ results,
         u32 hash_le[16];
 #pragma unroll
         for (int i = 0; i < 8; i++) hash_le[15 - i] = bswap32(first[i]);
-        permute(st);
+        run_permute<LOOP>(st, uni.diag);
 #pragma unroll
         for (int i = 0; i < 4; i++) {
             u64 c = gf_canon(st[i]);
@@ -874,3 +1113,23 @@ mining_main(u32* __restrict__ results,
         }
     }
 }
+
+extern "C" __global__ void LAIR_LAUNCH_BOUNDS
+mining_main(u32* __restrict__ results,
+            const MiningUniforms uni,
+            u32 total_threads,
+            u32 nonces_per_thread,
+            u32 total_nonces)
+{
+    mine_nonces<false>(results, uni, total_threads, nonces_per_thread, total_nonces);
+}
+
+extern "C" __global__ void LAIR_LOOP_LAUNCH_BOUNDS
+mining_loop(u32* __restrict__ results,
+            const MiningUniforms uni,
+            u32 total_threads,
+            u32 nonces_per_thread,
+            u32 total_nonces)
+{
+    mine_nonces<true>(results, uni, total_threads, nonces_per_thread, total_nonces);
+}

crates/engine-cuda/src/lib.rs

@@ -26,18 +26,60 @@ use std::time::Instant;
 const FATBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/mining.fatbin"));
 /// SM list the fat binary carries, for `--version` and the config metric.
 pub const CUDA_ARCHS: &str = env!("MINER_CUDA_ARCHS");
-const KERNEL_ID: &str = "cuda";
+/// lair: which kernel entry the module runs (quantus/miner#27). `unrolled`
+/// is `mining_main`, the deployed kernel; `loop` is `mining_loop`, the
+/// rolled-round twin under evaluation. Selected with MINER_CUDA_KERNEL so a
+/// host can be switched without a rebuild, and reported as the kernel id on
+/// the per-device metrics so the deploy validate and the dashboards can tell
+/// them apart.
+#[derive(Clone, Copy, PartialEq, Eq, Debug)]
+enum KernelChoice {
+    Unrolled,
+    Loop,
+}
+
+impl KernelChoice {
+    fn from_env() -> Self {
+        match std::env::var("MINER_CUDA_KERNEL").as_deref() {
+            Ok("loop") => KernelChoice::Loop,
+            Ok("unrolled") | Ok("") | Err(_) => KernelChoice::Unrolled,
+            Ok(other) => {
+                log::warn!(target: "cuda_engine", "MINER_CUDA_KERNEL={other:?} is not unrolled|loop; using unrolled");
+                KernelChoice::Unrolled
+            }
+        }
+    }
+    fn function(self) -> &'static str {
+        match self {
+            KernelChoice::Unrolled => "mining_main",
+            KernelChoice::Loop => "mining_loop",
+        }
+    }
+    /// Kernel id on metrics: `cuda` for the deployed kernel (what the deploy
+    /// validate asserts), `cuda-loop` for the candidate.
+    fn id(self) -> &'static str {
+        match self {
+            KernelChoice::Unrolled => "cuda",
+            KernelChoice::Loop => "cuda-loop",
+        }
+    }
+}
 /// Threads per block. Must match the kernel's `__launch_bounds__` (LAIR_TPB,
 /// default 256); override for experiments with MINER_CUDA_THREADS_PER_BLOCK
 /// on a kernel compiled with the same value.
 const THREADS_PER_BLOCK_DEFAULT: u32 = 256;
+/// The loop kernel's `LAIR_LOOP_TPB`.
+const LOOP_THREADS_PER_BLOCK_DEFAULT: u32 = 512;
 
-fn threads_per_block() -> u32 {
+fn threads_per_block(kernel: KernelChoice) -> u32 {
     std::env::var("MINER_CUDA_THREADS_PER_BLOCK")
         .ok()
         .and_then(|v| v.parse().ok())
         .filter(|v: &u32| *v > 0 && v.is_multiple_of(32))
-        .unwrap_or(THREADS_PER_BLOCK_DEFAULT)
+        .unwrap_or(match kernel {
+            KernelChoice::Unrolled => THREADS_PER_BLOCK_DEFAULT,
+            KernelChoice::Loop => LOOP_THREADS_PER_BLOCK_DEFAULT,
+        })
 }
 
 /// lair: size every batch to whole grids. The CLI default batch (1M nonces)
@@ -93,6 +135,7 @@ struct Device {
     func: CudaFunction,
     name: String,
     sm_count: u32,
+    kernel: KernelChoice,
     threads_per_block: u32,
     threads_per_sm: u32,
     metrics: metrics::DeviceMetrics,
@@ -130,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 {}
@@ -140,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,
@@ -171,9 +218,10 @@ impl CudaEngine {
                 Ok(d) => {
                     log::info!(
                         target: "cuda_engine",
-                        "CUDA device {ordinal}: {} ({} SMs) using {KERNEL_ID} kernel [sm_{{{}}}]",
+                        "CUDA device {ordinal}: {} ({} SMs) using {} kernel [sm_{{{}}}]",
                         d.name,
                         d.sm_count,
+                        d.kernel.id(),
                         CUDA_ARCHS
                     );
                     devices.push(Arc::new(d));
@@ -221,15 +269,17 @@ impl Device {
             cudarc::driver::sys::CUdevice_attribute::CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT,
         )? as u32;
         let module = ctx.load_module(Ptx::from_binary(FATBIN.to_vec()))?;
-        let func = module.load_function("mining_main")?;
+        let kernel = KernelChoice::from_env();
+        let func = module.load_function(kernel.function())?;
         Ok(Self {
             ctx,
             func,
             name,
             sm_count,
+            kernel,
             threads_per_sm: threads_per_sm(),
-            threads_per_block: threads_per_block(),
-            metrics: metrics::DeviceMetrics::new(ordinal, KERNEL_ID),
+            threads_per_block: threads_per_block(kernel),
+            metrics: metrics::DeviceMetrics::new(ordinal, kernel.id()),
         })
     }
 
@@ -493,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();

deploy/miner.env.tmpl

@@ -4,3 +4,5 @@
 # from the node host as separate 0640 files.
 QUANTUS_NODE_ADDR={{QUANTUS_NODE_ADDR}}
 QUANTUS_GPU_DEVICES={{QUANTUS_GPU_DEVICES}}
+# Kernel entry in the CUDA engine: unrolled (mining_main) or loop (mining_loop), quantus/miner#29.
+MINER_CUDA_KERNEL={{MINER_CUDA_KERNEL}}