blackbeard/miner
1
0
Fork 0
Code Issues 2 Pull Requests Actions Packages Projects Releases 2 Wiki Activity

engine-cuda: LAIR_MULADD screening variants (0 stays; 1-4 measured worse)

Browse Source 
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 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Ue5ZZm1Hiv5zPnucykKKuF
This commit is contained in:
rob thijssen
2026-09-14 19:40:59 +03:00
parent 4a6ed42d6b
commit 338b4a7d32
1 changed files with 123 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files

125
crates/engine-cuda/src/kernels/mining.cu
View File

@@ -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.