diff --git a/crates/engine-cuda/src/kernels/mining.cu b/crates/engine-cuda/src/kernels/mining.cu
index 139e056..bedcdeb 100644
--- a/crates/engine-cuda/src/kernels/mining.cu
+++ b/crates/engine-cuda/src/kernels/mining.cu
@@ -356,6 +356,16 @@ __device__ __forceinline__ u64 gf_canon(u64 a) {
 // 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
@@ -718,7 +728,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]);
@@ -727,15 +742,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);
     }
 }
 
@@ -873,25 +888,26 @@ __device__ __forceinline__ void permute_loop(u64 st[12], bool skip_last_ext, boo
     for (int i = 0; i < 12; i++) st[i] = gf_sbox(st[i]);
     ext_layer(st);
     st[0] = gf_add(st[0], RC_INTERNAL[0]);
+    const u64* d = MDS_DIAG;
 #if LAIR_LOOP_INT_PAIR == 2
 #pragma unroll 1
     for (int r = 0; r + 2 < N_INTERNAL; r += 2) {
         st[0] = gf_sbox(st[0]);
-        int_layer_rc(st, RC_INTERNAL[r + 1], true, nullptr);
+        int_layer_rc_d(st, RC_INTERNAL[r + 1], true, nullptr, d);
         st[0] = gf_sbox(st[0]);
-        int_layer_rc(st, RC_INTERNAL[r + 2], true, nullptr);
+        int_layer_rc_d(st, RC_INTERNAL[r + 2], true, nullptr, d);
     }
     st[0] = gf_sbox(st[0]);
-    int_layer_rc(st, RC_INTERNAL[N_INTERNAL - 1], true, nullptr);
+    int_layer_rc_d(st, RC_INTERNAL[N_INTERNAL - 1], true, nullptr, d);
 #else
 #pragma unroll 1
     for (int r = 0; r + 1 < N_INTERNAL; r++) {
         st[0] = gf_sbox(st[0]);
-        int_layer_rc(st, RC_INTERNAL[r + 1], true, nullptr);
+        int_layer_rc_d(st, RC_INTERNAL[r + 1], true, nullptr, d);
     }
 #endif
     st[0] = gf_sbox(st[0]);
-    int_layer_rc(st, 0ull, false, RC_TERMINAL[0]);
+    int_layer_rc_d(st, 0ull, false, RC_TERMINAL[0], d);
 #pragma unroll 1
     for (int r = 0; r + 1 < N_EXTERNAL_HALF; r++) {
 #pragma unroll