From 2bf3246634c6ecdc4a701780204095645b4d706c Mon Sep 17 00:00:00 2001 From: rob thijssen Date: Thu, 3 Sep 2026 15:51:11 +0300 Subject: [PATCH 1/2] engine-cuda: deferred carries, carry-flag arithmetic, one nonce per thread quantus/miner#3 step 5, measured on benjy (RTX 4090, 250 W, miner paused, interleaved rounds, parity verified on every variant): variant MH/s straight port (PR #16) 297-305 + grid cap 32768 threads/SM 304 (1 nonce per thread up to 4M) + deferred-carry linear layers 344 (+13%) + carry-flag gf_add / gf_reduce (PTX) 385 (+12%) + carry-flag accumulators 414 (+7%) 414 MH/s is 2.86x origin's wgpu build (144.4) on the same card and limit. Fewer instructions per hash also let the power-capped card clock higher (1920 -> 1995 MHz), which is part of the gain. Rejected by measurement: __noinline__ permute (-13%: code size was not the limit, the call convention was), -maxrregcount 80/64 (flat), internal-round unroll 1 or 22 (flat, within noise of 2). The batch sweep found batch size irrelevant once the grid keeps one nonce per thread (304-306 MH/s from 1M to 16M); batch stays at 1M for the smallest stale cost. All knobs remain as -D macros and MINER_NVCC_FLAGS so the harness can keep testing variants. Co-Authored-By: Claude Fable 5.1 Claude-Session: https://claude.ai/code/session_01CBgs2nSi4H2mdh8kD8vMX5 --- crates/engine-cuda/build.rs | 7 + crates/engine-cuda/src/kernels/mining.cu | 160 ++++++++++++++++++++++- crates/engine-cuda/src/lib.rs | 5 +- 3 files changed, 170 insertions(+), 2 deletions(-) diff --git a/crates/engine-cuda/build.rs b/crates/engine-cuda/build.rs index c3d6337..a305469 100644 --- a/crates/engine-cuda/build.rs +++ b/crates/engine-cuda/build.rs @@ -18,6 +18,7 @@ //! MINER_NVCC_CCBIN host compiler for nvcc (-ccbin) //! MINER_CUDA_ALLOW_UNSUPPORTED_COMPILER=1 pass -allow-unsupported-compiler //! MINER_CUDA_REQUIRE=1 fail the build instead of warning when nvcc is missing +//! MINER_NVCC_FLAGS extra nvcc arguments (whitespace-separated), e.g. -DLAIR_LAZY_ADD=0 use std::env; use std::fs; @@ -118,6 +119,7 @@ fn main() { "MINER_NVCC_CCBIN", "MINER_CUDA_ALLOW_UNSUPPORTED_COMPILER", "MINER_CUDA_REQUIRE", + "MINER_NVCC_FLAGS", ] { println!("cargo:rerun-if-env-changed={v}"); } @@ -163,6 +165,11 @@ fn main() { { cmd.arg("-allow-unsupported-compiler"); } + if let Ok(extra) = env::var("MINER_NVCC_FLAGS") { + for a in extra.split_whitespace() { + cmd.arg(a); + } + } for sm in &sms { cmd.arg("--generate-code") .arg(format!("arch=compute_{sm},code=sm_{sm}")); diff --git a/crates/engine-cuda/src/kernels/mining.cu b/crates/engine-cuda/src/kernels/mining.cu index 178f9a8..356b396 100644 --- a/crates/engine-cuda/src/kernels/mining.cu +++ b/crates/engine-cuda/src/kernels/mining.cu @@ -28,6 +28,36 @@ typedef unsigned int u32; // 2^64 mod P = 2^32 - 1 #define EPS 0xFFFFFFFFull +#ifndef LAIR_PTX_CARRY +#define LAIR_PTX_CARRY 1 +#endif +#ifndef LAIR_PTX_ACC +#define LAIR_PTX_ACC LAIR_PTX_CARRY +#endif + +#if LAIR_PTX_CARRY +// Carry-flag arithmetic: the wrap is read from the condition code instead of +// a compare, and folded with a multiply-add (c * EPS is exact for c in {0,1}). +__device__ __forceinline__ u64 gf_add(u64 a, u64 b) { + u64 s0, c1, s1, c2; + asm("add.cc.u64 %0, %2, %3;\n\taddc.u64 %1, 0, 0;" : "=l"(s0), "=l"(c1) : "l"(a), "l"(b)); + asm("add.cc.u64 %0, %2, %3;\n\taddc.u64 %1, 0, 0;" : "=l"(s1), "=l"(c2) : "l"(s0), "l"(c1 * EPS)); + return s1 + c2 * EPS; +} + +__device__ __forceinline__ u64 gf_reduce(u64 lo, u64 hi) { + u64 hi_hi = hi >> 32; + u64 hi_lo = hi & EPS; + u64 t0, bm; + // bm = -borrow: all ones when lo < hi_hi. + asm("sub.cc.u64 %0, %2, %3;\n\tsubc.u64 %1, 0, 0;" : "=l"(t0), "=l"(bm) : "l"(lo), "l"(hi_hi)); + t0 -= (bm & EPS); + u64 t1 = hi_lo * EPS; + u64 t2, c; + asm("add.cc.u64 %0, %2, %3;\n\taddc.u64 %1, 0, 0;" : "=l"(t2), "=l"(c) : "l"(t0), "l"(t1)); + return t2 + c * EPS; +} +#else // a + b mod P, lazy. A wrapped carry folds back as 2^64 = EPS (mod P); the // second fold can only be needed when the first one wrapped. __device__ __forceinline__ u64 gf_add(u64 a, u64 b) { @@ -49,6 +79,7 @@ __device__ __forceinline__ u64 gf_reduce(u64 lo, u64 hi) { if (t2 < t0) t2 += EPS; return t2; } +#endif __device__ __forceinline__ u64 gf_mul(u64 a, u64 b) { return gf_reduce(a * b, __umul64hi(a, b)); @@ -70,6 +101,126 @@ __device__ __forceinline__ u64 gf_canon(u64 a) { return a - (a >= P64 ? P64 : 0ull); } +#ifndef LAIR_LAZY_ADD +#define LAIR_LAZY_ADD 1 +#endif +#ifndef LAIR_INT_UNROLL +#define LAIR_INT_UNROLL 2 +#endif +// nvcc does not macro-expand `#pragma unroll N`; stringise through _Pragma. +#define LAIR_PRAGMA(x) _Pragma(#x) +#define LAIR_UNROLL(n) LAIR_PRAGMA(unroll n) +#ifndef LAIR_PTX_CARRY +#define LAIR_PTX_CARRY 1 +#endif +#ifndef LAIR_PTX_ACC +#define LAIR_PTX_ACC LAIR_PTX_CARRY +#endif +#ifndef LAIR_NOINLINE_PERMUTE +#define LAIR_NOINLINE_PERMUTE 0 +#endif + +// Deferred-carry accumulator: the value is lo + hi * 2^64 with hi small. Adds +// cost an add-with-carry instead of two compare-and-fold steps; the folds are +// paid once per output when the accumulator is reduced. Sound as long as +// hi < 2^32, which a linear layer's sums (at most a few dozen terms) satisfy. +struct Acc { + u64 lo; + u32 hi; +}; + +__device__ __forceinline__ Acc acc_of(u64 a) { + Acc r; r.lo = a; r.hi = 0u; return r; +} + +#if 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)); + a.lo = s; + a.hi += c; + return a; +} + +__device__ __forceinline__ Acc acc_add2(Acc a, Acc b) { + u64 s; u32 c; + asm("add.cc.u64 %0, %2, %3;\n\taddc.u32 %1, %4, %5;" : "=l"(s), "=r"(c) : "l"(a.lo), "l"(b.lo), "r"(a.hi), "r"(b.hi)); + a.lo = s; + a.hi = c; + return a; +} +#else +__device__ __forceinline__ Acc acc_add(Acc a, u64 b) { + u64 s = a.lo + b; + a.hi += (s < a.lo) ? 1u : 0u; + a.lo = s; + return a; +} + +__device__ __forceinline__ Acc acc_add2(Acc a, Acc b) { + u64 s = a.lo + b.lo; + a.hi += b.hi + ((s < a.lo) ? 1u : 0u); + a.lo = s; + return a; +} +#endif + +// 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) { + return gf_add(a.lo, (u64)a.hi * EPS); +} + +// a * b + c mod P: the addend is folded into the 128-bit product before the +// single reduction, saving a gf_add per term. +__device__ __forceinline__ u64 gf_mul_add(u64 a, u64 b, u64 c) { + u64 lo = a * b; + u64 hi = __umul64hi(a, b); + u64 s = lo + c; + hi += (s < lo) ? 1ull : 0ull; + return gf_reduce(s, hi); +} + +#if LAIR_LAZY_ADD +// External linear layer with deferred carries: every output is a sum of a +// handful of inputs, reduced once. +__device__ __forceinline__ void ext_layer(u64 st[12]) { + Acc acc[12]; +#pragma unroll + for (int chunk = 0; chunk < 3; chunk++) { + int o = chunk * 4; + u64 x0 = st[o], x1 = st[o + 1], x2 = st[o + 2], x3 = st[o + 3]; + Acc t01 = acc_add(acc_of(x0), x1); + Acc t23 = acc_add(acc_of(x2), x3); + Acc t0123 = acc_add2(t01, t23); + Acc t01123 = acc_add(t0123, x1); + Acc t01233 = acc_add(t0123, x3); + acc[o + 3] = acc_add(acc_add(t01233, x0), x0); + acc[o + 1] = acc_add(acc_add(t01123, x2), x2); + acc[o] = acc_add2(t01123, t01); + acc[o + 2] = acc_add2(t01233, t23); + } + Acc sums[4]; +#pragma unroll + for (int k = 0; k < 4; k++) { + sums[k] = acc_add2(acc_add2(acc[k], acc[k + 4]), acc[k + 8]); + } +#pragma unroll + for (int i = 0; i < 12; i++) { + st[i] = acc_reduce(acc_add2(acc[i], sums[i & 3])); + } +} + +// Internal linear layer: one deferred sum, folded into each diagonal multiply. +__device__ __forceinline__ void int_layer(u64 st[12]) { + Acc s = acc_of(st[0]); +#pragma unroll + for (int i = 1; i < 12; i++) s = acc_add(s, st[i]); + u64 sum = acc_reduce(s); +#pragma unroll + for (int i = 0; i < 12; i++) st[i] = gf_mul_add(st[i], MDS_DIAG[i], sum); +} +#else // External linear layer: 4x4 MDS on each chunk, then circulant sums. __device__ __forceinline__ void ext_layer(u64 st[12]) { #pragma unroll @@ -105,8 +256,15 @@ __device__ __forceinline__ void int_layer(u64 st[12]) { #pragma unroll for (int i = 0; i < 12; i++) st[i] = gf_add(gf_mul(st[i], MDS_DIAG[i]), sum); } +#endif +#if LAIR_NOINLINE_PERMUTE +// One copy of the (fully unrolled) permutation instead of three inlined ones: +// 17k instructions instead of 51k, for instruction-cache pressure. +__device__ __noinline__ void permute(u64 st[12]) { +#else __device__ __forceinline__ void permute(u64 st[12]) { +#endif ext_layer(st); #pragma unroll for (int r = 0; r < N_EXTERNAL_HALF; r++) { @@ -114,7 +272,7 @@ __device__ __forceinline__ void permute(u64 st[12]) { for (int i = 0; i < 12; i++) st[i] = gf_sbox(gf_add(st[i], RC_INITIAL[r][i])); ext_layer(st); } -#pragma unroll 2 + LAIR_UNROLL(LAIR_INT_UNROLL) for (int r = 0; r < N_INTERNAL; r++) { st[0] = gf_sbox(gf_add(st[0], RC_INTERNAL[r])); int_layer(st); diff --git a/crates/engine-cuda/src/lib.rs b/crates/engine-cuda/src/lib.rs index b7b7aca..5e5ffd7 100644 --- a/crates/engine-cuda/src/lib.rs +++ b/crates/engine-cuda/src/lib.rs @@ -29,7 +29,10 @@ const KERNEL_ID: &str = "cuda"; const THREADS_PER_BLOCK: u32 = 256; /// Threads per SM the grid is sized for (override: MINER_CUDA_THREADS_PER_SM). /// Sets how many nonces each thread loops over for a given batch; tuned with #2. -const THREADS_PER_SM_DEFAULT: u32 = 8192; +// Measured on the 4090 (#3): 8192 caps the grid at ~1M threads, so batches +// above 1M loop nonces per thread and lose up to 4%; 32768 keeps one nonce per +// thread up to 4M and is flat across batch sizes at ~305 MH/s. +const THREADS_PER_SM_DEFAULT: u32 = 32768; fn threads_per_sm() -> u32 { std::env::var("MINER_CUDA_THREADS_PER_SM") From b640d3a711070c75850ab17afb15b5d37039e411 Mon Sep 17 00:00:00 2001 From: rob thijssen Date: Thu, 3 Sep 2026 15:56:23 +0300 Subject: [PATCH 2/2] engine-cuda: fold round constants into the preceding linear layer; bench on cuda paths Each linear layer now folds the constant of the round that follows it into its final deferred-carry reduction (or, for the internal rounds, into the element-0 multiply-add), removing the per-element gf_add before every S-box. Parity verified. On benjy the effect is within run-to-run noise: three interleaved rounds gave medians of 418.5 vs 409.7 MH/s against the previous default, with SM clock swinging 1935-2025 MHz across rounds. Kept on as LAIR_FOLD_RC=1 for the instruction count; not claimed as a measured gain. bench.yaml: crates/engine-cuda/** now triggers the harness on PRs. The previous PR's kernel changes did not run it. Co-Authored-By: Claude Fable 5.1 Claude-Session: https://claude.ai/code/session_01CBgs2nSi4H2mdh8kD8vMX5 --- .gitea/workflows/bench.yaml | 1 + crates/engine-cuda/src/kernels/mining.cu | 81 ++++++++++++++++++++++-- 2 files changed, 78 insertions(+), 4 deletions(-) diff --git a/.gitea/workflows/bench.yaml b/.gitea/workflows/bench.yaml index dfbe1d6..772e851 100644 --- a/.gitea/workflows/bench.yaml +++ b/.gitea/workflows/bench.yaml @@ -38,6 +38,7 @@ on: pull_request: paths: - crates/engine-gpu/** + - crates/engine-cuda/** - crates/engine-cpu/** - crates/pow-core/** - crates/miner-service/** diff --git a/crates/engine-cuda/src/kernels/mining.cu b/crates/engine-cuda/src/kernels/mining.cu index 356b396..5242041 100644 --- a/crates/engine-cuda/src/kernels/mining.cu +++ b/crates/engine-cuda/src/kernels/mining.cu @@ -34,6 +34,11 @@ typedef unsigned int u32; #ifndef LAIR_PTX_ACC #define LAIR_PTX_ACC LAIR_PTX_CARRY #endif +// Fold the next round's constant into the linear layer's final reduction +// instead of a separate gf_add per element before the S-box. +#ifndef LAIR_FOLD_RC +#define LAIR_FOLD_RC 1 +#endif #if LAIR_PTX_CARRY // Carry-flag arithmetic: the wrap is read from the condition code instead of @@ -116,6 +121,11 @@ __device__ __forceinline__ u64 gf_canon(u64 a) { #ifndef LAIR_PTX_ACC #define LAIR_PTX_ACC LAIR_PTX_CARRY #endif +// Fold the next round's constant into the linear layer's final reduction +// instead of a separate gf_add per element before the S-box. +#ifndef LAIR_FOLD_RC +#define LAIR_FOLD_RC 1 +#endif #ifndef LAIR_NOINLINE_PERMUTE #define LAIR_NOINLINE_PERMUTE 0 #endif @@ -181,10 +191,20 @@ __device__ __forceinline__ u64 gf_mul_add(u64 a, u64 b, u64 c) { return gf_reduce(s, hi); } +// a * b + (c.lo + c.hi * 2^64) mod P, for a deferred-carry addend. +__device__ __forceinline__ u64 gf_mul_add_acc(u64 a, u64 b, Acc c) { + u64 lo = a * b; + u64 hi = __umul64hi(a, b); + u64 s = lo + c.lo; + hi += ((s < lo) ? 1ull : 0ull) + (u64)c.hi; + return gf_reduce(s, hi); +} + #if LAIR_LAZY_ADD // External linear layer with deferred carries: every output is a sum of a // handful of inputs, reduced once. -__device__ __forceinline__ void ext_layer(u64 st[12]) { +// `rc` is the next round's constant row to fold into the outputs, or nullptr. +__device__ __forceinline__ void ext_layer_rc(u64 st[12], const u64* rc) { Acc acc[12]; #pragma unroll for (int chunk = 0; chunk < 3; chunk++) { @@ -207,19 +227,37 @@ __device__ __forceinline__ void ext_layer(u64 st[12]) { } #pragma unroll for (int i = 0; i < 12; i++) { - st[i] = acc_reduce(acc_add2(acc[i], sums[i & 3])); + Acc o = acc_add2(acc[i], sums[i & 3]); + if (rc != nullptr) o = acc_add(o, rc[i]); + st[i] = acc_reduce(o); } } +__device__ __forceinline__ void ext_layer(u64 st[12]) { ext_layer_rc(st, nullptr); } + // Internal linear layer: one deferred sum, folded into each diagonal multiply. -__device__ __forceinline__ void int_layer(u64 st[12]) { +// `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(u64 st[12], u64 rc0, bool has_rc0, const u64* rc_row) { Acc s = acc_of(st[0]); #pragma unroll for (int i = 1; i < 12; i++) s = acc_add(s, st[i]); u64 sum = acc_reduce(s); + if (rc_row != nullptr) { #pragma unroll - for (int i = 0; i < 12; i++) st[i] = gf_mul_add(st[i], MDS_DIAG[i], sum); + 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])); + } else { + if (has_rc0) { + st[0] = gf_mul_add_acc(st[0], MDS_DIAG[0], acc_add(acc_of(sum), rc0)); + } else { + st[0] = gf_mul_add(st[0], MDS_DIAG[0], sum); + } +#pragma unroll + for (int i = 1; i < 12; i++) st[i] = gf_mul_add(st[i], MDS_DIAG[i], sum); + } } + +__device__ __forceinline__ void int_layer(u64 st[12]) { int_layer_rc(st, 0ull, false, nullptr); } #else // External linear layer: 4x4 MDS on each chunk, then circulant sums. __device__ __forceinline__ void ext_layer(u64 st[12]) { @@ -265,6 +303,40 @@ __device__ __noinline__ void permute(u64 st[12]) { #else __device__ __forceinline__ void permute(u64 st[12]) { #endif +#if LAIR_LAZY_ADD && LAIR_FOLD_RC + // Each linear layer folds the constant of the round that follows it. + ext_layer_rc(st, RC_INITIAL[0]); +#pragma unroll + for (int r = 0; r < N_EXTERNAL_HALF; r++) { +#pragma unroll + for (int i = 0; i < 12; i++) st[i] = gf_sbox(st[i]); + if (r + 1 < N_EXTERNAL_HALF) { + ext_layer_rc(st, RC_INITIAL[r + 1]); + } else { + ext_layer(st); + st[0] = gf_add(st[0], RC_INTERNAL[0]); + } + } + LAIR_UNROLL(LAIR_INT_UNROLL) + for (int r = 0; r < N_INTERNAL; r++) { + st[0] = gf_sbox(st[0]); + if (r + 1 < N_INTERNAL) { + int_layer_rc(st, RC_INTERNAL[r + 1], true, nullptr); + } else { + int_layer_rc(st, 0ull, false, RC_TERMINAL[0]); + } + } +#pragma unroll + for (int r = 0; r < N_EXTERNAL_HALF; r++) { +#pragma unroll + for (int i = 0; i < 12; i++) st[i] = gf_sbox(st[i]); + if (r + 1 < N_EXTERNAL_HALF) { + ext_layer_rc(st, RC_TERMINAL[r + 1]); + } else { + ext_layer(st); + } + } +#else ext_layer(st); #pragma unroll for (int r = 0; r < N_EXTERNAL_HALF; r++) { @@ -283,6 +355,7 @@ __device__ __forceinline__ void permute(u64 st[12]) { for (int i = 0; i < 12; i++) st[i] = gf_sbox(gf_add(st[i], RC_TERMINAL[r][i])); ext_layer(st); } +#endif } __device__ __forceinline__ u32 bswap32(u32 v) {