From c1cac9141983ff7edcee06579109f0fd0e7cf006 Mon Sep 17 00:00:00 2001 From: rob thijssen Date: Mon, 14 Sep 2026 00:49:11 +0300 Subject: [PATCH] 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 Claude-Session: https://claude.ai/code/session_01Ue5ZZm1Hiv5zPnucykKKuF --- crates/engine-cuda/src/kernels/mining.cu | 18 +++++++++++++++++- 1 file changed, 17 insertions(+), 1 deletion(-) diff --git a/crates/engine-cuda/src/kernels/mining.cu b/crates/engine-cuda/src/kernels/mining.cu index d55a659..a2fc02d 100644 --- a/crates/engine-cuda/src/kernels/mining.cu +++ b/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,