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

@@ -20,6 +20,7 @@ on:
       - "**.md"
       - .gitea/workflows/ci.yml
       - .gitea/workflows/bench.yaml
+      - .gitea/workflows/release.yaml
   workflow_dispatch:
     inputs:
       mode:
@@ -60,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
@@ -87,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
@@ -145,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"
@@ -209,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 \

.gitea/workflows/release.yaml

@@ -0,0 +1,78 @@
+---
+# Release binaries for independent miners (quantus/miner#27).
+#
+# Pushing a tag `v<upstream version>-lair.<n>` builds the miner and the bench
+# harness on the cuda-13.0 runner (Fedora 43 like the mining hosts, and the
+# only runner with nvcc, so the binary carries the sm_86/sm_89/sm_120 cubins
+# and PTX), records what was built, and publishes a Gitea release with the
+# binaries, their checksums and the build report attached. The tag's commit
+# is embedded in `--version` so a downloaded binary identifies itself.
+#
+# Deploying to the fleet is deploy.yaml's job and follows main; a release is
+# only a public snapshot and changes nothing on the hosts.
+name: release
+
+on:
+  push:
+    tags: ["v*-lair.*"]
+
+env:
+  CARGO_TERM_COLOR: always
+
+jobs:
+  build:
+    runs-on: cuda-13.0
+    steps:
+      - uses: actions/checkout@v4
+      - name: build quantus-miner and quantus-bench
+        env:
+          MINER_BUILD_SHA: ${{ github.sha }}
+          MINER_CUDA_REQUIRE: "1"
+        run: |
+          set -euo pipefail
+          cargo build --release --locked -p miner-cli -p bench-harness
+          ./target/release/quantus-miner --version
+          ./target/release/quantus-bench --version || true
+      - name: package
+        run: |
+          set -euo pipefail
+          tag="${GITHUB_REF_NAME}"
+          dir="quantus-miner-${tag}-linux-x86_64"
+          mkdir -p "dist/${dir}"
+          cp target/release/quantus-miner target/release/quantus-bench "dist/${dir}/"
+          cp LICENSE "dist/${dir}/"
+          {
+            echo "quantus-miner ${tag}"
+            echo "commit: ${GITHUB_SHA}"
+            echo "built: $(date -u +%Y-%m-%dT%H:%M:%SZ) on $(source /etc/os-release && echo "$PRETTY_NAME"), $(ldd --version | head -1)"
+            echo "cuda: $(nvcc --version | grep -oE 'release [0-9.]+' | head -1)"
+            echo "cubins: sm_86 sm_89 sm_120 + compute_120 PTX"
+            echo "version string: $(./target/release/quantus-miner --version)"
+          } > "dist/${dir}/BUILD.txt"
+          (cd dist && tar -czf "${dir}.tar.gz" "${dir}" && sha256sum "${dir}.tar.gz" > "${dir}.tar.gz.sha256")
+          cat "dist/${dir}/BUILD.txt" "dist/${dir}.tar.gz.sha256"
+      - name: publish release
+        run: |
+          set -euo pipefail
+          tag="${GITHUB_REF_NAME}"
+          dir="quantus-miner-${tag}-linux-x86_64"
+          api="${GITHUB_API_URL}/repos/${GITHUB_REPOSITORY}"
+          auth="Authorization: token ${{ secrets.GITHUB_TOKEN }}"
+          body=$(python3 - "$tag" "dist/${dir}/BUILD.txt" "dist/${dir}.tar.gz.sha256" <<'PY'
+          import json, sys, pathlib
+          tag, build, sha = sys.argv[1], pathlib.Path(sys.argv[2]).read_text(), pathlib.Path(sys.argv[3]).read_text().strip()
+          text = (
+              f"quantus-miner {tag}: Linux x86_64, NVIDIA RTX 30/40/50 (sm_86, sm_89, sm_120, plus PTX), "
+              "glibc 2.42 or newer, driver with CUDA 13.0 support.\n\n"
+              "No dev fee, no licence server, no telemetry. Measured rates and how to reproduce them are in the README.\n\n"
+              f"```\n{build}```\n\nsha256: `{sha}`\n"
+          )
+          print(json.dumps({"tag_name": tag, "name": f"quantus-miner {tag}", "body": text, "draft": False, "prerelease": False}))
+          PY
+          )
+          id=$(curl -fsS -X POST -H "$auth" -H "Content-Type: application/json" "${api}/releases" -d "$body" | python3 -c 'import json,sys; print(json.load(sys.stdin)["id"])')
+          echo "release id ${id}"
+          for f in "dist/${dir}.tar.gz" "dist/${dir}.tar.gz.sha256" "dist/${dir}/BUILD.txt"; do
+            curl -fsS -X POST -H "$auth" "${api}/releases/${id}/assets?name=$(basename "$f")" -F "attachment=@${f}" > /dev/null
+            echo "attached $(basename "$f")"
+          done

README.md

@@ -2,6 +2,58 @@
 
 High-performance external mining service for Quantus Network with support for CPU, GPU, and hybrid CPU+GPU mining.
 
+## This fork
+
+Like the official miner, but fast: a native CUDA engine for NVIDIA cards, kept
+bit-exact with the reference hash and measured against every other miner we
+can get hold of. No dev fee, no licence server, no telemetry, solo or pool.
+Apache-2.0, like upstream.
+
+**Measured hashrate**, three interleaved 30-second rounds per card at the
+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 | **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
+(2026-09-13); the closed pool binaries were run in benchmark mode only. How
+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
+(glibc 2.42 or newer), and `quantus-bench`, the measurement harness. Nothing
+is compiled at run time; the driver needs CUDA 13.0 support (580.x).
+
+**Reproduce the numbers.** `quantus-bench` pins the power limit it expects,
+warms up, reports the median of timed windows with their spread, and checks
+GPU hashes against the CPU reference for random jobs:
+
+```bash
+# hashrate: median of 5 x 30 s windows, one worker thread per card
+./quantus-bench --duration-secs 30 --runs 5 --workers 1 --expect-power-limit 400
+
+# parity: 2000 random jobs, every found hash recomputed on the CPU
+./quantus-bench --duration-secs 3 --runs 1 --parity-jobs 2000
+```
+
+Compare builds with interleaved rounds (A, B, A, B, ...) rather than one
+block of runs each: on a power-capped card, clock drift between blocks looks
+like a 0.5% code change.
+
+**Solo mining** against your own node is the `serve` command below with the
+node's auth token and certificate fingerprint; nothing else is needed.
+
 ## Building
 
 ```bash

crates/engine-cuda/Cargo.toml

@@ -12,6 +12,7 @@ metrics = { path = "../metrics" }
 primitive-types = { workspace = true }
 log = { workspace = true }
 cudarc = { version = "0.19", default-features = false, features = ["driver", "nvrtc", "dynamic-loading", "cuda-13000"] }
+qp-poseidon-constants = "1.1"
 
 [build-dependencies]
 qp-poseidon-constants = "1.1"

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

@@ -118,6 +118,129 @@ __device__ __forceinline__ u64 gf_reduce(u64 lo, u64 hi) {
 }
 #endif
 
+// lair: field arithmetic as self-contained PTX blocks (quantus/miner#3).
+//
+// Measured against the closed-source qpow-cuda 1.0.7 kernel on the same
+// silicon: its multiply-plus-reduce is 11 SASS instructions on sm_120 and a
+// squaring 10, against ~14 per multiply here, and it folds the internal
+// layer's row sum into the diagonal multiply's addend for free. The whole
+// difference is in how the 128-bit product is reduced: `mad.lo.cc.u32` +
+// `madc.hi.cc.u32` become one IMAD.WIDE.U32 with a carry-out predicate, and
+// the 2^96 = -1 term is folded with three more ALU ops. Each block below is
+// one complete field operation, so there is nothing across its boundary for
+// nvcc to optimise -- unlike LAIR_PTX_CARRY, whose fragments starved the
+// optimiser on sm_120.
+//
+// LAIR_EXACT_REDUCE: the closed-source kernel drops the final borrow of the
+// reduction. That is wrong only when the 128-bit product has bits 64..95 all
+// zero and its low 64 bits below its top 32 bits: about 2^-64 of multiplies
+// with field-element inputs (verified by emulating the chain against a
+// reference multiply: 0 mismatches in 3M random products, and the constructed
+// product 2^48 * 2^48 = 2^96 does fail). The other apparent hazard, the
+// 32-bit wrap of hh + carry, cannot occur: hh = 2^32 - 1 needs both operands
+// within 2^32 of 2^64, and then the low half is too small to carry. So with 0
+// the kernel computes a wrong hash for about 1472 * 2^-64 ~ 1e-16 of nonces,
+// which can neither hide a real solution nor pass a false one in practice
+// (a wrong hash still has to beat the target, 2^-47 at mainnet difficulty,
+// to be submitted, and the node would reject it); with 1 every hash is
+// bit-exact, for three more instructions per multiply. Measured on
+// beast's 5090s, three interleaved rounds: 0 = 1996 MH/s, 1 = 1695 MH/s,
+// against 1356 for the C reduction -- exactness costs 15%.
+#ifndef LAIR_FUSED_MUL
+#define LAIR_FUSED_MUL 1
+#endif
+#ifndef LAIR_EXACT_REDUCE
+#define LAIR_EXACT_REDUCE 0
+#endif
+
+#if LAIR_FUSED_MUL
+// EPS as an operand ptxas cannot see the value of: with an immediate it would
+// rewrite `x * EPS` into shift-and-subtract, which loses on this hardware.
+__device__ __constant__ u32 LAIR_EPS32 = 0xFFFFFFFFu;
+
+// Reduce {ll, lh} + 2^64 * {hl, hh} in place, lazily (result < 2^64, not
+// canonical), using 2^64 = EPS and 2^96 = -1 (mod P). E is the EPS operand.
+//   {ll,lh} += hl * EPS, carry cy            (mad.lo.cc / madc.hi.cc)
+//   c = hh + cy; lh += cy                    (+cy * 2^32 ...)
+//   {ll,lh} -= c                             (... - cy - hh = +cy*EPS - hh)
+#if LAIR_EXACT_REDUCE
+#define LAIR_REDUCE_PTX(E) \
+    "mad.lo.cc.u32 ll, hl, " E ", ll;\n\t" \
+    "madc.hi.cc.u32 lh, hl, " E ", lh;\n\t" \
+    "addc.u32 c, hh, 0;\n\t" \
+    "addc.u32 lh, lh, 0;\n\t" \
+    "sub.cc.u32 ll, ll, c;\n\t" \
+    "subc.cc.u32 lh, lh, 0;\n\t" \
+    "subc.u32 c, 0, 0;\n\t" \
+    "sub.cc.u32 ll, ll, c;\n\t" \
+    "subc.u32 lh, lh, 0;\n\t"
+#else
+#define LAIR_REDUCE_PTX(E) \
+    "mad.lo.cc.u32 ll, hl, " E ", ll;\n\t" \
+    "madc.hi.cc.u32 lh, hl, " E ", lh;\n\t" \
+    "addc.u32 c, hh, 0;\n\t" \
+    "addc.u32 lh, lh, 0;\n\t" \
+    "sub.cc.u32 ll, ll, c;\n\t" \
+    "subc.u32 lh, lh, 0;\n\t"
+#endif
+
+__device__ __forceinline__ u64 gf_mul(u64 a, u64 b) {
+    u64 r;
+    u32 e = LAIR_EPS32;
+    asm("{\n\t"
+        ".reg .b32 a0, a1, b0, b1, 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"
+        "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"
+        "mov.b32 ll, p0l;\n\t"
+        LAIR_REDUCE_PTX("%3")
+        "mov.b64 %0, {ll, lh};\n\t"
+        "}"
+        : "=l"(r) : "l"(a), "l"(b), "r"(e));
+    return r;
+}
+
+// a^2 with three 32x32 products: a0^2 + 2*a0*a1*2^32 + a1^2*2^64.
+__device__ __forceinline__ u64 gf_sqr(u64 a) {
+    u64 r;
+    u32 e = LAIR_EPS32;
+    asm("{\n\t"
+        ".reg .b32 a0, a1, p0l, p0h, m2l, m2h, c32, ll, lh, hl, hh, c;\n\t"
+        ".reg .b64 p0, m, m2, p3, t, hi;\n\t"
+        "mov.b64 {a0, a1}, %1;\n\t"
+        "mul.wide.u32 p0, a0, a0;\n\t"
+        "mul.wide.u32 m, a0, a1;\n\t"
+        "shl.b64 m2, m, 1;\n\t"
+        "shr.u64 t, m, 63;\n\t"
+        "mov.b64 {m2l, m2h}, m2;\n\t"
+        "mov.b64 {p0l, p0h}, p0;\n\t"
+        "add.cc.u32 lh, p0h, m2l;\n\t"
+        "cvt.u32.u64 c32, t;\n\t"
+        "mov.b64 t, {m2h, c32};\n\t"
+        "mul.wide.u32 p3, a1, a1;\n\t"
+        "addc.u64 hi, p3, t;\n\t"
+        "mov.b64 {hl, hh}, hi;\n\t"
+        "mov.b32 ll, p0l;\n\t"
+        LAIR_REDUCE_PTX("%2")
+        "mov.b64 %0, {ll, lh};\n\t"
+        "}"
+        : "=l"(r) : "l"(a), "r"(e));
+    return r;
+}
+
+#else
 __device__ __forceinline__ u64 gf_mul(u64 a, u64 b) {
     return gf_reduce(a * b, __umul64hi(a, b));
 }
@@ -139,8 +262,11 @@ __device__ __forceinline__ u64 gf_mul(u64 a, u64 b) {
 __device__ __forceinline__ u64 gf_sqr(u64 a) {
     return gf_reduce(a * a, __umul64hi(a, a));
 }
+#endif // LAIR_FUSED_MUL
 
-// x^7
+// x^7. Writing the four steps as one PTX block in 32-bit halves was tried
+// (2026-09-13) and compiles to byte-identical SASS: the 64-bit pair moves
+// between blocks are free, so there is nothing to save here.
 __device__ __forceinline__ u64 gf_sbox(u64 x) {
     u64 x2 = gf_sqr(x);
     u64 x4 = gf_sqr(x2);
@@ -187,6 +313,83 @@ __device__ __forceinline__ u64 gf_canon(u64 a) {
 #ifndef LAIR_NOINLINE_PERMUTE
 #define LAIR_NOINLINE_PERMUTE 0
 #endif
+// Threads per block; the host's MINER_CUDA_THREADS_PER_BLOCK must agree.
+#ifndef LAIR_TPB
+#define LAIR_TPB 256
+#endif
+// Minimum resident blocks per SM the compiler must fit (caps registers at
+// 65536 / (LAIR_TPB * LAIR_MINBLOCKS)); 0 leaves the register budget to ptxas.
+// Measured 2026-09-13 on beast against the 110-register default (2205 MH/s):
+// 64 registers by any route (1024 threads, or 256 x 4, or 512 x 2) spills
+// 216-264 bytes and loses 5% (2091-2107); 80 registers (256 x 3, 52 bytes of
+// spill) is neutral (2207). The quanpool kernel runs 1024 threads per SM at
+// 64 registers without spilling because its rounds are loops with the
+// constants in constant memory; this unrolled kernel cannot get there.
+#ifndef LAIR_MINBLOCKS
+#define LAIR_MINBLOCKS 0
+#endif
+#if LAIR_MINBLOCKS
+#define LAIR_LAUNCH_BOUNDS __launch_bounds__(LAIR_TPB, LAIR_MINBLOCKS)
+#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.
+#ifndef LAIR_EARLY_REJECT
+#define LAIR_EARLY_REJECT 1
+#endif
+// Derive each nonce's state after the first linear layer from the host's
+// precomputed base state plus one scalar times column 7 of the external
+// matrix, instead of absorbing and recomputing the layer per nonce. Falls
+// back to the full path for the nonces past a carry out of the low limb.
+#ifndef LAIR_NONCE_DIR
+#define LAIR_NONCE_DIR 1
+#endif
+#if LAIR_NONCE_DIR && !(LAIR_LAZY_ADD && LAIR_FOLD_RC && LAIR_FUSED_MUL)
+#error "LAIR_NONCE_DIR needs LAIR_LAZY_ADD, LAIR_FOLD_RC and LAIR_FUSED_MUL"
+#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
@@ -201,7 +404,45 @@ __device__ __forceinline__ Acc acc_of(u64 a) {
     Acc r; r.lo = a; r.hi = 0u; return r;
 }
 
-#if LAIR_PTX_ACC
+#if LAIR_PTX_ACC && LAIR_FUSED_MUL
+// 32-bit limb chains. On sm_120 ptxas turns `add.cc.u64` + `addc.u32` into
+// two IADD.64 (one for the carry, one for the sum) and a select, five
+// instructions per add; the 32-bit chain is IADD, IADD.X and one IADD3.X
+// that ptxas merges with the neighbouring carry, under three.
+__device__ __forceinline__ Acc acc_add(Acc a, u64 b) {
+    u64 s; u32 c;
+    asm("{\n\t"
+        ".reg .b32 al, ah, bl, bh;\n\t"
+        "mov.b64 {al, ah}, %2;\n\t"
+        "mov.b64 {bl, bh}, %3;\n\t"
+        "add.cc.u32 al, al, bl;\n\t"
+        "addc.cc.u32 ah, ah, bh;\n\t"
+        "addc.u32 %1, %4, 0;\n\t"
+        "mov.b64 %0, {al, ah};\n\t"
+        "}"
+        : "=l"(s), "=r"(c) : "l"(a.lo), "l"(b), "r"(a.hi));
+    a.lo = s;
+    a.hi = c;
+    return a;
+}
+
+__device__ __forceinline__ Acc acc_add2(Acc a, Acc b) {
+    u64 s; u32 c;
+    asm("{\n\t"
+        ".reg .b32 al, ah, bl, bh;\n\t"
+        "mov.b64 {al, ah}, %2;\n\t"
+        "mov.b64 {bl, bh}, %3;\n\t"
+        "add.cc.u32 al, al, bl;\n\t"
+        "addc.cc.u32 ah, ah, bh;\n\t"
+        "addc.u32 %1, %4, %5;\n\t"
+        "mov.b64 %0, {al, ah};\n\t"
+        "}"
+        : "=l"(s), "=r"(c) : "l"(a.lo), "l"(b.lo), "r"(a.hi), "r"(b.hi));
+    a.lo = s;
+    a.hi = c;
+    return a;
+}
+#elif 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));
@@ -233,6 +474,185 @@ __device__ __forceinline__ Acc acc_add2(Acc a, Acc b) {
 }
 #endif
 
+#if LAIR_FUSED_MUL
+// lo + hi * 2^64 = lo + hi * EPS (mod P), one fused reduction. Exact: with
+// no 2^96 term there is no borrow to drop.
+__device__ __forceinline__ u64 acc_reduce(Acc a) {
+    u64 r;
+    u32 e = LAIR_EPS32;
+    asm("{\n\t"
+        ".reg .b32 ll, lh, c;\n\t"
+        "mov.b64 {ll, lh}, %1;\n\t"
+        "mad.lo.cc.u32 ll, %2, %3, ll;\n\t"
+        "madc.hi.cc.u32 lh, %2, %3, lh;\n\t"
+        "addc.u32 c, 0, 0;\n\t"
+        "addc.u32 lh, lh, 0;\n\t"
+        "sub.cc.u32 ll, ll, c;\n\t"
+        "subc.u32 lh, lh, 0;\n\t"
+        "mov.b64 %0, {ll, lh};\n\t"
+        "}"
+        : "=l"(r) : "l"(a.lo), "r"(a.hi), "r"(e));
+    return r;
+}
+
+// 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;
+    u64 c0 = c & EPS, c1 = c >> 32;
+    asm("{\n\t"
+        ".reg .b32 a0, a1, b0, b1, p0l, p0h, p3l, p3h, m0, m1, 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"
+        "add.cc.u64 m, m, m2;\n\t"
+        "addc.u32 cw, p3h, 0;\n\t"
+        "mov.b64 {m0, m1}, m;\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;
+}
+#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.
+__device__ __forceinline__ u64 gf_mul_add_acc(u64 a, u64 b, Acc c) {
+    return gf_mul_add(a, b, acc_reduce(c));
+}
+#else
 // 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) {
@@ -257,6 +677,7 @@ __device__ __forceinline__ u64 gf_mul_add_acc(u64 a, u64 b, Acc c) {
     hi += ((s < lo) ? 1ull : 0ull) + (u64)c.hi;
     return gf_reduce(s, hi);
 }
+#endif // LAIR_FUSED_MUL
 
 #if LAIR_LAZY_ADD
 // External linear layer with deferred carries: every output is a sum of a
@@ -293,25 +714,48 @@ __device__ __forceinline__ void ext_layer_rc(u64 st[12], const u64* rc) {
 
 __device__ __forceinline__ void ext_layer(u64 st[12]) { ext_layer_rc(st, nullptr); }
 
+// Output 0 of the external linear layer alone, inputs untouched: the chunk
+// sums it needs (about a tenth of the full layer).
+__device__ __forceinline__ u64 ext_layer_out0(const u64 st[12]) {
+    Acc a[3];
+#pragma unroll
+    for (int chunk = 0; chunk < 3; chunk++) {
+        int o = chunk * 4;
+        Acc t01 = acc_add(acc_of(st[o]), st[o + 1]);
+        Acc t23 = acc_add(acc_of(st[o + 2]), st[o + 3]);
+        Acc t01123 = acc_add(acc_add2(t01, t23), st[o + 1]);
+        a[chunk] = acc_add2(t01123, t01);
+    }
+    Acc sum0 = acc_add2(acc_add2(a[0], a[1]), a[2]);
+    return acc_reduce(acc_add2(a[0], sum0));
+}
+
 // 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]);
     u64 sum = acc_reduce(s);
     if (rc_row != nullptr) {
+        // 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(acc_of(sum), 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(acc_of(sum), 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);
     }
 }
 
@@ -344,6 +788,18 @@ __device__ __forceinline__ void ext_layer(u64 st[12]) {
     }
 }
 
+__device__ __forceinline__ u64 ext_layer_out0(const u64 st[12]) {
+    u64 a[3];
+#pragma unroll
+    for (int chunk = 0; chunk < 3; chunk++) {
+        int o = chunk * 4;
+        u64 t01 = gf_add(st[o], st[o + 1]);
+        u64 t23 = gf_add(st[o + 2], st[o + 3]);
+        a[chunk] = gf_add(gf_add(gf_add(t01, t23), st[o + 1]), t01);
+    }
+    return gf_add(a[0], gf_add(gf_add(a[0], a[1]), a[2]));
+}
+
 // Internal linear layer: diagonal matrix plus full sum.
 __device__ __forceinline__ void int_layer(u64 st[12]) {
     u64 sum = st[0];
@@ -354,16 +810,19 @@ __device__ __forceinline__ void int_layer(u64 st[12]) {
 }
 #endif
 
+// `skip_last_ext`: leave the state before the final external linear layer,
+// for the caller to finish with ext_layer_out0 / ext_layer.
 #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]) {
+__device__ __noinline__ void permute(u64 st[12], bool skip_last_ext = false, bool skip_first_ext = false) {
 #else
-__device__ __forceinline__ void permute(u64 st[12]) {
+__device__ __forceinline__ void permute(u64 st[12], bool skip_last_ext = false, bool skip_first_ext = false) {
 #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]);
+    // `skip_first_ext`: the caller supplies the state after this layer.
+    if (!skip_first_ext) ext_layer_rc(st, RC_INITIAL[0]);
 #pragma unroll
     for (int r = 0; r < N_EXTERNAL_HALF; r++) {
 #pragma unroll
@@ -390,7 +849,7 @@ __device__ __forceinline__ void permute(u64 st[12]) {
         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 {
+        } else if (!skip_last_ext) {
             ext_layer(st);
         }
     }
@@ -411,11 +870,77 @@ __device__ __forceinline__ void permute(u64 st[12]) {
     for (int r = 0; r < N_EXTERNAL_HALF; r++) {
 #pragma unroll
         for (int i = 0; i < 12; i++) st[i] = gf_sbox(gf_add(st[i], RC_TERMINAL[r][i]));
-        ext_layer(st);
+        if (r + 1 < N_EXTERNAL_HALF || !skip_last_ext) ext_layer(st);
     }
 #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);
 }
@@ -437,15 +962,31 @@ struct MiningUniforms {
     u64 midstate[12];
     u32 start_nonce[16];
     u32 target[16];
+    // 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) == 224, "MiningUniforms must match the host layout");
+static_assert(sizeof(MiningUniforms) == 416, "MiningUniforms must match the host layout");
 
-extern "C" __global__ void __launch_bounds__(256)
-mining_main(u32* __restrict__ results,
-            const MiningUniforms uni,
-            u32 total_threads,
-            u32 nonces_per_thread,
-            u32 total_nonces)
+#if LAIR_NONCE_DIR
+// Column 7 of the external matrix circ(2*M4, M4, M4): what the layer adds per
+// unit of element 7, which is where the low nonce limb is absorbed. Checked by
+// engine-cuda's `column_seven_of_external_matrix` test.
+__device__ __constant__ u64 LAIR_DIR7[12] = {1, 1, 3, 2, 2, 2, 6, 4, 1, 1, 3, 2};
+#endif
+
+// 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;
@@ -477,14 +1018,49 @@ mining_main(u32* __restrict__ results,
         // pad, squeeze twice (3 permutations instead of 5; the second squeeze
         // only for candidates).
         u64 st[12];
+#if LAIR_NONCE_DIR
+        if (current_nonce[1] == uni.start_nonce[1]) {
+            // Only the low limb differs from the batch's first nonce, so only
+            // element 7 of the absorbed state does, by d; the layer is linear.
+            u64 a = (u64)bswap32(current_nonce[0]);
+            u64 b = (u64)bswap32(uni.start_nonce[0]);
+            u64 d = a >= b ? a - b : a + (P64 - b);
+#pragma unroll
+            for (int i = 0; i < 12; i++) st[i] = gf_mul_add(d, LAIR_DIR7[i], uni.layer0_base[i]);
+        } else {
+            // Past a carry out of the low limb: absorb and apply the layer.
+#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]));
+            ext_layer_rc(st, RC_INITIAL[0]);
+        }
+        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);
-        permute(st);
+#if LAIR_EARLY_REJECT
+        // 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, uni.diag, true);
+        {
+            u64 c0 = gf_canon(ext_layer_out0(st));
+            u32 h0 = bswap32((u32)(c0 & EPS));
+            u32 t0 = uni.target[15];
+            if (h0 > t0) continue;
+            if (h0 == t0 && bswap32((u32)(c0 >> 32)) > uni.target[14]) continue;
+        }
+        ext_layer(st);
+#else
+        run_permute<LOOP>(st, uni.diag);
+#endif
 
         // First squeeze: most significant 256 bits of the hash.
         u32 first[8];
@@ -506,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]);
@@ -537,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,8 +26,87 @@ 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";
-const THREADS_PER_BLOCK: u32 = 256;
+/// 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(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(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)
+/// fills a 5090 about eleven waves deep with a half-empty tail wave and pays
+/// the per-launch gap every 1M hashes. Rounding the batch up to a multiple of
+/// the grid (SMs x threads per SM, one nonce per thread) measured +1.6% on
+/// beast, and two grids per batch (each thread loops over two nonces) a
+/// further +0.4% there and +0.5% on the 4090, neutral on the 3060
+/// (quantus/miner#3, 2026-09-13). A batch is then ~10 ms on every card, so a
+/// job switch discards about 0.1% of a block interval. Overrides:
+/// MINER_CUDA_BATCH_ALIGN=0 for the CLI batch as given, MINER_CUDA_BATCH_WAVES
+/// for the minimum number of grids per batch.
+const BATCH_WAVES_DEFAULT: u32 = 2;
+
+fn batch_align() -> bool {
+    std::env::var("MINER_CUDA_BATCH_ALIGN")
+        .map(|v| v != "0")
+        .unwrap_or(true)
+}
+
+fn batch_waves() -> u32 {
+    std::env::var("MINER_CUDA_BATCH_WAVES")
+        .ok()
+        .and_then(|v| v.parse().ok())
+        .filter(|v: &u32| *v > 0)
+        .unwrap_or(BATCH_WAVES_DEFAULT)
+}
 /// 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.
 // Measured on the 4090 (#3): 8192 caps the grid at ~1M threads, so batches
@@ -56,6 +135,8 @@ struct Device {
     func: CudaFunction,
     name: String,
     sm_count: u32,
+    kernel: KernelChoice,
+    threads_per_block: u32,
     threads_per_sm: u32,
     metrics: metrics::DeviceMetrics,
 }
@@ -86,9 +167,19 @@ struct MiningUniforms {
     midstate: [u64; 12],
     start_nonce: [u32; 16],
     target: [u32; 16],
+    /// lair: the state after the first external linear layer (and its round
+    /// constant) for the batch's first nonce. That layer is linear, and only
+    /// the least significant nonce limb changes within a batch, so the kernel
+    /// 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 then 32 x u32),
+// 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 {}
@@ -96,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>() == 224);
+const _: () = assert!(std::mem::size_of::<MiningUniforms>() == 416);
 
 pub struct CudaEngine {
     engine_id: usize,
@@ -127,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));
@@ -177,14 +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(),
-            metrics: metrics::DeviceMetrics::new(ordinal, KERNEL_ID),
+            threads_per_block: threads_per_block(kernel),
+            metrics: metrics::DeviceMetrics::new(ordinal, kernel.id()),
         })
     }
 
@@ -340,8 +435,15 @@ impl CudaEngine {
             let headroom =
                 (U512::one() << 256) - (current_start & ((U512::one() << 256) - U512::one()));
             let cap = remaining.min(headroom);
-            let this_batch: u32 = if cap > U512::from(self.batch_size) {
+            let want: u32 = if batch_align() {
+                let grid = (dev.sm_count * dev.threads_per_sm).max(1);
+                let waves = self.batch_size.div_ceil(grid).max(batch_waves());
+                waves.saturating_mul(grid)
+            } else {
                 self.batch_size
+            };
+            let this_batch: u32 = if cap > U512::from(want) {
+                want
             } else {
                 cap.low_u32()
             };
@@ -420,8 +522,9 @@ impl CudaEngine {
         // the remainder per thread.
         let max_threads = (dev.sm_count * dev.threads_per_sm) as u64;
         let logical_threads = (batch_size as u64).min(max_threads).max(1);
-        let num_blocks = ((logical_threads as u32).div_ceil(THREADS_PER_BLOCK)).max(1);
-        let total_threads = (num_blocks * THREADS_PER_BLOCK) as u64;
+        let tpb = dev.threads_per_block;
+        let num_blocks = ((logical_threads as u32).div_ceil(tpb)).max(1);
+        let total_threads = (num_blocks * tpb) as u64;
         let nonces_per_thread = ((batch_size as u64).div_ceil(total_threads)).max(1) as u32;
         let total_threads_u32 = total_threads as u32;
 
@@ -434,10 +537,13 @@ 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 layer0_base = first_layer_after_absorb(&midstate, &start_u32s[..8]);
         let uni = MiningUniforms {
             midstate,
             start_nonce: start_u32s,
             target: res.target_u32s,
+            layer0_base,
+            diag: qp_poseidon_constants::POSEIDON2_MATRIX_DIAG_12_RAW,
         };
 
         let stream = res.stream.clone();
@@ -445,7 +551,7 @@ impl CudaEngine {
             stream.memset_zeros(&mut res.results)?;
             let cfg = LaunchConfig {
                 grid_dim: (num_blocks, 1, 1),
-                block_dim: (THREADS_PER_BLOCK, 1, 1),
+                block_dim: (tpb, 1, 1),
                 shared_mem_bytes: 0,
             };
             let mut launch = stream.launch_builder(&dev.func);
@@ -507,3 +613,69 @@ impl CudaEngine {
 fn bytemuck_cast(words: &[u32]) -> Vec<u8> {
     words.iter().flat_map(|w| w.to_le_bytes()).collect()
 }
+
+// lair: host-side twin of the kernel's absorb + first external linear layer,
+// in canonical Goldilocks arithmetic. Bit-exact with `ext_layer_rc` in
+// kernels/mining.cu modulo p (the kernel keeps lazy representatives).
+const GOLDILOCKS_P: u128 = 0xFFFF_FFFF_0000_0001;
+
+fn gf_add(a: u64, b: u64) -> u64 {
+    ((a as u128 + b as u128) % GOLDILOCKS_P) as u64
+}
+
+/// State after absorbing the low nonce limbs into the midstate and applying
+/// the first external layer plus the first round constant, for the batch's
+/// first nonce. `nonce_le` are the low 8 little-endian u32 limbs.
+fn first_layer_after_absorb(midstate: &[u64; 12], nonce_le: &[u32]) -> [u64; 12] {
+    let mut st = *midstate;
+    for i in 0..8 {
+        st[i] = gf_add(st[i], nonce_le[7 - i].swap_bytes() as u64);
+    }
+    let mut out = [0u64; 12];
+    for chunk in 0..3 {
+        let o = chunk * 4;
+        let (x0, x1, x2, x3) = (st[o], st[o + 1], st[o + 2], st[o + 3]);
+        let t01 = gf_add(x0, x1);
+        let t23 = gf_add(x2, x3);
+        let t0123 = gf_add(t01, t23);
+        let t01123 = gf_add(t0123, x1);
+        let t01233 = gf_add(t0123, x3);
+        out[o + 3] = gf_add(t01233, gf_add(x0, x0));
+        out[o + 1] = gf_add(t01123, gf_add(x2, x2));
+        out[o] = gf_add(t01123, t01);
+        out[o + 2] = gf_add(t01233, t23);
+    }
+    let mut sums = [0u64; 4];
+    for k in 0..4 {
+        sums[k] = gf_add(gf_add(out[k], out[k + 4]), out[k + 8]);
+    }
+    let rc = &qp_poseidon_constants::POSEIDON2_INITIAL_EXTERNAL_CONSTANTS_RAW[0];
+    for i in 0..12 {
+        out[i] = gf_add(gf_add(out[i], sums[i & 3]), rc[i]);
+    }
+    out
+}
+
+#[cfg(test)]
+mod layer0_tests {
+    use super::*;
+
+    /// The kernel's fast path assumes column 7 of the external matrix is this
+    /// vector: applying the layer to the unit vector e7 must reproduce it.
+    #[test]
+    fn column_seven_of_external_matrix() {
+        let zero = [0u64; 12];
+        // absorb puts bswap(limb 0) into element 7: choose limb 0 so that it becomes 1
+        let mut nonce = [0u32; 8];
+        nonce[0] = 1u32.swap_bytes();
+        let with = first_layer_after_absorb(&zero, &nonce);
+        let without = first_layer_after_absorb(&zero, &[0u32; 8]);
+        let rc = &qp_poseidon_constants::POSEIDON2_INITIAL_EXTERNAL_CONSTANTS_RAW[0];
+        let expect: [u64; 12] = [1, 1, 3, 2, 2, 2, 6, 4, 1, 1, 3, 2];
+        for i in 0..12 {
+            assert_eq!(without[i], rc[i]);
+            let diff = (with[i] as u128 + GOLDILOCKS_P - without[i] as u128) % GOLDILOCKS_P;
+            assert_eq!(diff as u64, expect[i], "column entry {i}");
+        }
+    }
+}

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}}