From 338b4a7d3215279986c8d2dede102afe488fe8a5 Mon Sep 17 00:00:00 2001 From: rob thijssen Date: Mon, 14 Sep 2026 19:40:59 +0300 Subject: [PATCH] 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.