From 338b4a7d3215279986c8d2dede102afe488fe8a5 Mon Sep 17 00:00:00 2001 From: rob thijssen Date: Mon, 14 Sep 2026 19:40:59 +0300 Subject: [PATCH 1/3] engine-cuda: LAIR_MULADD screening variants (0 stays; 1-4 measured worse) Static screen on sm_120, loop kernel internal round body: 0 = 239, 1 = 260, 2 = 283, 3 = 275, 4 = 265 instructions. Co-Authored-By: Claude Fable 5.1 Claude-Session: https://claude.ai/code/session_01Ue5ZZm1Hiv5zPnucykKKuF --- crates/engine-cuda/src/kernels/mining.cu | 125 ++++++++++++++++++++++- 1 file changed, 123 insertions(+), 2 deletions(-) diff --git a/crates/engine-cuda/src/kernels/mining.cu b/crates/engine-cuda/src/kernels/mining.cu index f6e4ffd..139e056 100644 --- a/crates/engine-cuda/src/kernels/mining.cu +++ b/crates/engine-cuda/src/kernels/mining.cu @@ -480,8 +480,15 @@ __device__ __forceinline__ u64 acc_reduce(Acc a) { 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. +// 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; @@ -510,6 +517,120 @@ __device__ __forceinline__ u64 gf_mul_add(u64 a, u64 b, u64 c) { : "=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. -- 2.52.0 From 4572a5e56743b3120140438fbc96046ee45562c8 Mon Sep 17 00:00:00 2001 From: rob thijssen Date: Mon, 14 Sep 2026 19:43:59 +0300 Subject: [PATCH 2/3] engine-cuda: attribute the loop kernel's internal-round excess; drop two no-op knobs Constant hoisting in C and 32-bit-halves delivery of the diagonal both compile to byte-identical SASS and are removed. The remaining 40 instructions per round over the quanpool kernel are itemised in the kernel comment. Co-Authored-By: Claude Fable 5.1 Claude-Session: https://claude.ai/code/session_01Ue5ZZm1Hiv5zPnucykKKuF --- crates/engine-cuda/src/kernels/mining.cu | 34 +++++++++++++++++------- 1 file changed, 25 insertions(+), 9 deletions(-) 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 -- 2.52.0 From 1fea51cfc3786036f0cfefbb53368dd23a18f02d Mon Sep 17 00:00:00 2001 From: rob thijssen Date: Mon, 14 Sep 2026 19:51:22 +0300 Subject: [PATCH 3/3] engine-cuda: loop kernel takes the diagonal from the launch parameters MiningUniforms carries the twelve diagonal constants (416 bytes now) and the loop kernel reads them there (LAIR_LOOP_PARAM_DIAG=1) instead of from the __constant__ table. Static: still 13 loads per internal round (ptxas reloads either way) but spill 88 -> 32 bytes. Measured, three interleaved rounds, loop kernel only: beast 2135 -> 2162 MH/s (+1.3%), benjy 725 -> 730 (+0.7%); parity 300/300 and 200/200. The unrolled kernel is byte- identical (21,120 instructions, 110 registers). Co-Authored-By: Claude Fable 5.1 Claude-Session: https://claude.ai/code/session_01Ue5ZZm1Hiv5zPnucykKKuF --- crates/engine-cuda/src/kernels/mining.cu | 35 +++++++++++++++++------- crates/engine-cuda/src/lib.rs | 9 ++++-- 2 files changed, 32 insertions(+), 12 deletions(-) diff --git a/crates/engine-cuda/src/kernels/mining.cu b/crates/engine-cuda/src/kernels/mining.cu index bedcdeb..0d527aa 100644 --- a/crates/engine-cuda/src/kernels/mining.cu +++ b/crates/engine-cuda/src/kernels/mining.cu @@ -353,6 +353,11 @@ __device__ __forceinline__ u64 gf_canon(u64 a) { #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. @@ -876,7 +881,7 @@ __device__ __forceinline__ void permute(u64 st[12], bool skip_last_ext = false, // 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) { +__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++) { @@ -888,7 +893,12 @@ __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]); +#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) { @@ -919,14 +929,15 @@ __device__ __forceinline__ void permute_loop(u64 st[12], bool skip_last_ext, boo if (!skip_last_ext) ext_layer(st); } #else -__device__ __forceinline__ void permute_loop(u64 st[12], bool skip_last_ext, bool skip_first_ext) { +__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 -__device__ __forceinline__ void run_permute(u64 st[12], bool skip_last_ext = false, bool skip_first_ext = false) { - if (LOOP) permute_loop(st, skip_last_ext, skip_first_ext); +__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); } @@ -954,8 +965,12 @@ struct MiningUniforms { // 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) == 320, "MiningUniforms must match the host layout"); +static_assert(sizeof(MiningUniforms) == 416, "MiningUniforms must match the host layout"); #if LAIR_NONCE_DIR // Column 7 of the external matrix circ(2*M4, M4, M4): what the layer adds per @@ -1020,13 +1035,13 @@ __device__ __forceinline__ void mine_nonces(u32* __restrict__ results, 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(st, false, true); + run_permute(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])); - run_permute(st); + run_permute(st, uni.diag); #endif st[0] = gf_add(st[0], 1ull); st[1] = gf_add(st[1], 1ull); @@ -1034,7 +1049,7 @@ __device__ __forceinline__ void mine_nonces(u32* __restrict__ results, // 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(st, true); + run_permute(st, uni.diag, true); { u64 c0 = gf_canon(ext_layer_out0(st)); u32 h0 = bswap32((u32)(c0 & EPS)); @@ -1044,7 +1059,7 @@ __device__ __forceinline__ void mine_nonces(u32* __restrict__ results, } ext_layer(st); #else - run_permute(st); + run_permute(st, uni.diag); #endif // First squeeze: most significant 256 bits of the hash. @@ -1067,7 +1082,7 @@ __device__ __forceinline__ void mine_nonces(u32* __restrict__ results, u32 hash_le[16]; #pragma unroll for (int i = 0; i < 8; i++) hash_le[15 - i] = bswap32(first[i]); - run_permute(st); + run_permute(st, uni.diag); #pragma unroll for (int i = 0; i < 4; i++) { u64 c = gf_canon(st[i]); diff --git a/crates/engine-cuda/src/lib.rs b/crates/engine-cuda/src/lib.rs index d6fa1a5..b5f7d5b 100644 --- a/crates/engine-cuda/src/lib.rs +++ b/crates/engine-cuda/src/lib.rs @@ -173,9 +173,13 @@ struct MiningUniforms { /// 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, 32 x u32, 12 x u64), +// 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 {} @@ -183,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::() == 320); +const _: () = assert!(std::mem::size_of::() == 416); pub struct CudaEngine { engine_id: usize, @@ -539,6 +543,7 @@ impl CudaEngine { start_nonce: start_u32s, target: res.target_u32s, layer0_base, + diag: qp_poseidon_constants::POSEIDON2_MATRIX_DIAG_12_RAW, }; let stream = res.stream.clone(); -- 2.52.0