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

perf: fused-PTX field arithmetic, +54..66% on every card #23

Merged
grenade merged 1 commits from perf/fused-field-ptx into main 2026-09-13 19:40:28 +00:00
Conversation 1 Commits 1 Files Changed 1
+225 -3
1 changed files with 225 additions and 3 deletions
Download Patch File Download Diff File Expand all files Collapse all files
crates/engine-cuda/src/kernels/mining.cu
+225 -3
View File
@@ -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);
}