From 3763996ae698c96be111c8711a17bc0914bf5be5 Mon Sep 17 00:00:00 2001 From: rob thijssen Date: Sun, 13 Sep 2026 22:29:46 +0300 Subject: [PATCH] engine-cuda: fused-PTX field arithmetic, +54..66% on every card The closed-source qpow-cuda kernel runs 2x our hashrate on the same silicon with the same 1472 multiplies per hash. Its PTX shows why: every field operation is one self-contained inline-PTX block, the 128-bit product is reduced with mad.lo.cc/madc.hi.cc (one IMAD.WIDE with a carry-out) plus three ALU ops, the internal layer's row sum rides in the mad.wide partial products for free, and accumulators are 32-bit limb chains that ptxas turns into IADD/IADD.X with merged carries. Ours paid a compare-and-select reduction (~14 ops), a 64-bit add-with-carry per diagonal element, and five instructions per accumulator add on sm_120. This ports those primitives behind LAIR_FUSED_MUL (default 1): gf_mul, gf_sqr (three-product squaring), gf_mul_add with the addend folded into the partial products, acc_add/acc_add2 as limb chains, and a four-op acc_reduce. The round constant now joins the unreduced sum in the internal layer. LAIR_EXACT_REDUCE (default 0) keeps the reduction bit-exact for three more ops per multiply. With 0 the final borrow is dropped, as in their kernel: wrong only when the product's bits 64..95 are zero and its low 64 bits are below its top 32 bits, about 2^-64 per multiply (emulated against a reference: 0/3M random mismatches, 2^48*2^48 fails as predicted). Exactness measured at -15% on beast, so it is off. Measured, three interleaved rounds, parity 40/40 on each host: beast 2x5090 1356 -> 2146 MH/s (+58%) benjy 4090 460 -> 711 MH/s (+54%) quadbrat 3060 84 -> 140 MH/s (+66%) Static sm_120 31,280 -> 20,888 instructions, sm_86/89 39,208 -> 22,576, no spills. ncu: 22,913 instructions per nonce against their 22,233. Co-Authored-By: Claude Fable 5.1 Claude-Session: https://claude.ai/code/session_01Ue5ZZm1Hiv5zPnucykKKuF --- crates/engine-cuda/src/kernels/mining.cu | 228 ++++++++++++++++++++++- 1 file changed, 225 insertions(+), 3 deletions(-) diff --git a/crates/engine-cuda/src/kernels/mining.cu b/crates/engine-cuda/src/kernels/mining.cu index 2bcb56e..45d8352 100644 --- a/crates/engine-cuda/src/kernels/mining.cu +++ b/crates/engine-cuda/src/kernels/mining.cu @@ -118,6 +118,128 @@ __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,6 +261,7 @@ __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 __device__ __forceinline__ u64 gf_sbox(u64 x) { @@ -201,7 +324,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 +394,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 +476,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 @@ -302,11 +522,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); }