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/src/kernels/mining.cu

@@ -317,6 +317,22 @@ __device__ __forceinline__ u64 gf_canon(u64 a) {
 #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.
@@ -726,7 +742,7 @@ 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 __launch_bounds__(LAIR_TPB)
+extern "C" __global__ void LAIR_LAUNCH_BOUNDS
 mining_main(u32* __restrict__ results,
             const MiningUniforms uni,
             u32 total_threads,