engine-cuda: LAIR_MINBLOCKS launch-bounds knob, with the occupancy results

No default change. Records on the knob what forcing the register budget
measured on beast against the quanpool kernel (quantus/miner#27): 64
registers by any launch shape spills and loses 5%, 80 is neutral.

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

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,42 @@ __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
+// 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 +363,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 +433,64 @@ __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. 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;
+}
+
+// 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 +515,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,6 +552,22 @@ __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).
@@ -302,11 +577,13 @@ __device__ __forceinline__ void int_layer_rc(u64 st[12], u64 rc0, bool has_rc0,
     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] = gf_mul_add_acc(st[i], MDS_DIAG[i], 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] = gf_mul_add_acc(st[0], MDS_DIAG[0], acc_add(s, rc0));
         } else {
             st[0] = gf_mul_add(st[0], MDS_DIAG[0], sum);
         }
@@ -344,6 +621,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 +643,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 +682,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,7 +703,7 @@ __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
 }
@@ -437,10 +729,20 @@ 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];
 };
-static_assert(sizeof(MiningUniforms) == 224, "MiningUniforms must match the host layout");
+static_assert(sizeof(MiningUniforms) == 320, "MiningUniforms must match the host layout");
 
-extern "C" __global__ void __launch_bounds__(256)
+#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
+
+extern "C" __global__ void LAIR_LAUNCH_BOUNDS
 mining_main(u32* __restrict__ results,
             const MiningUniforms uni,
             u32 total_threads,
@@ -477,14 +779,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]);
+        }
+        permute(st, 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);
+#endif
         st[0] = gf_add(st[0], 1ull);
         st[1] = gf_add(st[1], 1ull);
+#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.
+        permute(st, 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
         permute(st);
+#endif
 
         // First squeeze: most significant 256 bits of the hash.
         u32 first[8];

crates/engine-cuda/src/lib.rs

@@ -27,7 +27,44 @@ 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;
+/// 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;
+
+fn threads_per_block() -> 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)
+}
+
+/// 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 +93,7 @@ struct Device {
     func: CudaFunction,
     name: String,
     sm_count: u32,
+    threads_per_block: u32,
     threads_per_sm: u32,
     metrics: metrics::DeviceMetrics,
 }
@@ -86,9 +124,15 @@ 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],
 }
 
-// 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, 12 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 +140,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>() == 320);
 
 pub struct CudaEngine {
     engine_id: usize,
@@ -184,6 +228,7 @@ impl Device {
             name,
             sm_count,
             threads_per_sm: threads_per_sm(),
+            threads_per_block: threads_per_block(),
             metrics: metrics::DeviceMetrics::new(ordinal, KERNEL_ID),
         })
     }
@@ -340,8 +385,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 +472,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 +487,12 @@ 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,
         };
 
         let stream = res.stream.clone();
@@ -445,7 +500,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 +562,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}");
+        }
+    }
+}