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

engine-cuda: attribute the loop kernel's internal-round excess; drop two no-op knobs

Browse Source 
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 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Ue5ZZm1Hiv5zPnucykKKuF
This commit is contained in:
grenade
2026-09-14 19:43:59 +03:00
co-authored by Claude Fable 5.1
parent 338b4a7d32
commit 4572a5e567
1 changed files with 25 additions and 9 deletions
Download Patch File Download Diff File Expand all files Collapse all files
crates/engine-cuda/src/kernels/mining.cu
+25 -9
View File
@@ -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