High-performance external mining service for Quantus Network with support for CPU, GPU, and hybrid CPU+GPU mining.
Like the official miner, but fast: a native CUDA engine for NVIDIA cards, kept bit-exact with the reference hash and measured against every other miner we can get hold of. No dev fee, no licence server, no telemetry, solo or pool. Apache-2.0, like upstream.
Measured hashrate, three interleaved 30-second rounds per card at the enforced power cap, memory clock locked at 810 MHz, CUDA 13.0 driver 580:
| card | power cap | this fork | upstream --cuda-gpu | quantusminer.com qpow-cuda 1.0.7 | quanpool-miner 6.2.4 (5% fee) |
|---|---|---|---|---|---|
| RTX 5090 | 400 W | 1104 MH/s | 944 | 1171 | 1234 (1172 after fee) |
| RTX 4090 | 250 W | 727 MH/s (loop kernel) | |||
| RTX 3060 | 130 W | 141 MH/s |
The 5090 column was measured on the same card in the same session (2026-09-13); the closed pool binaries were run in benchmark mode only. How the pool kernels get their remaining edge, and what has been tried, is written up in issue #27.
Since v4.0.2-lair.2 the binary carries two kernels. The default (unrolled)
is fastest on the RTX 5090 and 3060; RTX 40-series cards measure 1.7% better
on the rolled-loop one, selected with MINER_CUDA_KERNEL=loop in the
environment (it reports as kernel cuda-loop on the metrics). Both are
bit-exact with the CPU reference; try each with quantus-bench on your card.
Binaries are on the releases page:
one Linux x86_64 quantus-miner carrying cubins for sm_86 (RTX 30), sm_89
(RTX 40) and sm_120 (RTX 50) plus PTX for anything newer, built on Fedora 43
(glibc 2.42 or newer), and quantus-bench, the measurement harness. Nothing
is compiled at run time; the driver needs CUDA 13.0 support (580.x).
Reproduce the numbers. quantus-bench pins the power limit it expects,
warms up, reports the median of timed windows with their spread, and checks
GPU hashes against the CPU reference for random jobs:
# hashrate: median of 5 x 30 s windows, one worker thread per card
./quantus-bench --duration-secs 30 --runs 5 --workers 1 --expect-power-limit 400
# parity: 2000 random jobs, every found hash recomputed on the CPU
./quantus-bench --duration-secs 3 --runs 1 --parity-jobs 2000
Compare builds with interleaved rounds (A, B, A, B, ...) rather than one block of runs each: on a power-capped card, clock drift between blocks looks like a 0.5% code change.
Solo mining against your own node is the serve command below with the
node's auth token and certificate fingerprint; nothing else is needed.
# CPU-only build (default)
cargo build -p miner-cli --release
# With GPU support (recommended)
cargo build -p miner-cli --release
The binary will be available at target/release/quantus-miner.
The node requires a shared auth token and TLS cert pin. Both live under the
node's chain config dir (<base-path>/chains/<chain>/):
miner-auth-token — shared secret (not logged by the node; read this file)miner-tls-cert-sha256 — SHA-256 of the miner TLS cert (also printed in node logs)# Preferred: mount/read the node's chain config files
./target/release/quantus-miner serve \
--node-addr 127.0.0.1:9833 \
--auth-token-file /path/to/miner-auth-token \
--tls-cert-sha256-file /path/to/miner-tls-cert-sha256 \
--cpu-workers 4
# Or pass values directly (token from miner-auth-token; fingerprint also in node logs)
./target/release/quantus-miner serve \
--node-addr 127.0.0.1:9833 \
--auth-token <TOKEN> \
--tls-cert-sha256 <FINGERPRINT> \
--gpu-devices 1
# Hybrid CPU+GPU mining
./target/release/quantus-miner serve \
--node-addr 127.0.0.1:9833 \
--auth-token-file /path/to/miner-auth-token \
--tls-cert-sha256-file /path/to/miner-tls-cert-sha256 \
--cpu-workers 4 \
--gpu-devices 1
| Argument | Environment Variable | Description | Default |
|---|---|---|---|
--node-addr <ADDR> | MINER_NODE_ADDR | Node address to connect to | 127.0.0.1:9833 |
--auth-token <TOKEN> | MINER_AUTH_TOKEN | Shared secret from the node's miner-auth-token file (not logged) | required |
--auth-token-file <PATH> | MINER_AUTH_TOKEN_FILE | Read the shared secret from a file (preferred) | — |
--tls-cert-sha256 <HEX> | MINER_TLS_CERT_SHA256 | SHA-256 of the node's miner TLS cert (miner-tls-cert-sha256 / node logs) | required |
--tls-cert-sha256-file <PATH> | MINER_TLS_CERT_SHA256_FILE | Read the TLS cert fingerprint from a file | — |
--cpu-workers <N> | MINER_CPU_WORKERS | Number of CPU worker threads | Auto-detect |
--gpu-devices <N> | MINER_GPU_DEVICES | Number of GPU devices | Auto-detect |
--gpu-batch-size <N> | MINER_GPU_BATCH_SIZE | GPU batch size in nonces | 1000000 |
--cpu-batch-size <N> | MINER_CPU_BATCH_SIZE | CPU batch size in hashes | 10000 |
--gpu-throttle-ms <MS> | MINER_GPU_THROTTLE_MS | Sleep duration (ms) between GPU batches | 0 |
--metrics-port <PORT> | MINER_METRICS_PORT | Prometheus metrics port | 9900 |
GPU support uses WGPU for cross-platform acceleration:
Build with GPU support:
cargo build -p miner-cli --release
Platform requirements:
nvidia-driver, mesa-vulkan-drivers)sudo powermetrics --samplers gpu_power -i 1000nvidia-smi (NVIDIA) or radeontop (AMD)All serve examples need the auth token and TLS pin (files or inline values).
# CPU mining with 8 workers
./target/release/quantus-miner serve \
--auth-token-file /path/to/miner-auth-token \
--tls-cert-sha256-file /path/to/miner-tls-cert-sha256 \
--cpu-workers 8
# Pure GPU mining
./target/release/quantus-miner serve \
--auth-token-file /path/to/miner-auth-token \
--tls-cert-sha256-file /path/to/miner-tls-cert-sha256 \
--gpu-devices 1
# GPU mining with throttle (reduce GPU utilization)
./target/release/quantus-miner serve \
--auth-token-file /path/to/miner-auth-token \
--tls-cert-sha256-file /path/to/miner-tls-cert-sha256 \
--gpu-devices 1 --gpu-throttle-ms 50
# Hybrid mining: 4 CPU + 1 GPU workers
./target/release/quantus-miner serve \
--auth-token-file /path/to/miner-auth-token \
--tls-cert-sha256-file /path/to/miner-tls-cert-sha256 \
--cpu-workers 4 --gpu-devices 1
# With verbose logging
RUST_LOG=debug ./target/release/quantus-miner serve \
--auth-token-file /path/to/miner-auth-token \
--tls-cert-sha256-file /path/to/miner-tls-cert-sha256 \
--cpu-workers 2 --gpu-devices 1
# Production setup with metrics
./target/release/quantus-miner serve \
--node-addr 127.0.0.1:9833 \
--auth-token-file /path/to/miner-auth-token \
--tls-cert-sha256-file /path/to/miner-tls-cert-sha256 \
--cpu-workers 6 \
--gpu-devices 1 \
--metrics-port 9900
The miner uses a QUIC-based protocol for communication with the node:
Ready { token } must match the node's miner-auth-tokenquantus-miner/2Ready (miner→node), NewJob (node→miner), JobResult (miner→node)For full protocol specification, see the node's MINING.md.
# Benchmark CPU performance
./target/release/quantus-miner benchmark --cpu-workers 8 --duration 30
# Benchmark GPU performance
./target/release/quantus-miner benchmark --gpu-devices 1 --duration 30
# Benchmark hybrid performance
./target/release/quantus-miner benchmark --cpu-workers 4 --gpu-devices 1 --duration 30
46 activities
Second commit (12bd36a) adds the per-host switch to the deploy matrix: cuda_kernel is rendered into miner.env as MINER_CUDA_KERNEL, kernel stays the id validate asserts. benjy runs loop / cuda-loop on the measurement above; quadbrat and beast stay unrolled / cuda. The earlier bench-workflow failure on this PR (run 67) was the harness refusing a card my own A/B was occupying at that minute, not a build problem.
12bd36a deploy: per-host CUDA kernel choice; benjy runs the loop kernelbenjy and quadbrat, same binary, three interleaved rounds, parity 200/200 on the loop kernel on each:
| host | unrolled | loop | | |---|---|---|---| | benjy RTX 4090 (sm_89) | 716-718 MH/s at 1935-1950 MHz | 727-736 MH/s at 2010-2055 MHz | +1.7% | | quadbrat RTX 3060 (sm_86) | 141.4-142.0 | 138.9-139.1 | -1.8% | | beast RTX 5090 (sm_120) | 2206 | 2130 | -3.4% |
On the 4090 the 88-register no-spill build runs at a higher clock and wins. So the loop kernel is the better default for benjy today; the switch is per host in the deploy matrix (MINER_CUDA_KERNEL=loop in its environment, kernel: cuda-loop for validate), a follow-up once this merges.
816b1ef engine-cuda: loop-structured mining_loop kernel behind MINER_CUDA_KERNELSeveral statements above are stale; this comment is the current reading.
The repository is public and publishes release binaries (#28, first tag v4.0.2-lair.1). The purpose has widened from a fleet advantage to an open miner that independent operators can run against the closed pool binaries: on 2026-09-13 one reward preimage authored ~62% of blocks in the exporter's window. The branch model and the no-PRs-to-origin rule are unchanged; origin's own CUDA engine (v4.2.0, --cuda-gpu) measured 944 MH/s on a 5090 against our 1104 (#27), so there is nothing to take from it yet.
All three hosts mine since the mainnet call on 2026-09-09 (#20); beast's inference is off. The "beast does not mine" policy above is history.
Perf, measured, per card at the power caps: 5090 283 (wgpu) -> 1104 MH/s, 4090 144 -> 716, 3060 37 -> 141. The steps: #16 native CUDA (2x), #17 deferred carries, #22 per-arch carry path (+20% sm_120), #23 fused-PTX arithmetic (+54..66%), #26 whole-grid batches, early reject, nonce direction (+3%). Memory clock lock on all hosts (lair/quantus#10; +2% on the 5090, +13% on the 4090).
Workstream table, resolved: #2, #3, #4, #6, #9 closed with numbers earlier; #5 and #7 closed today as superseded by the CUDA engine; #8 closed as complete. Open: #25 (ptxas anomaly, a standing warning: every kernel change passes exhaustive parity on all three architectures before merge), #27 (the remaining ~10% to the fastest closed binary: a loop-structured permute with one continuous carry chain per round; the tensor-core idea #24 was measured and rejected).
Correction to the audit's "no algorithmic slack": still true for the hash, but the closed kernels showed slack in how the arithmetic is scheduled, and the largest single gain of the whole effort (#23) came from there.
monday, september 14, 2026 — 12:28:54 utcClosing as superseded. The bubble this was meant to hide was measured at 1.9% of batch time on the 4090 and about 5% on a 5090, and #26 recovered what was recoverable on the engine that actually runs: batches are now whole grids, two per launch, which took beast from 2146 to 2205 MH/s (+2.8%) and benjy +1%; going to four or eight nonces per thread measured neutral to -1.2%, so the remaining gap between launches is not worth a second stream on the CUDA engine either. The quanpool kernel does run two streams, and its profile (#27) puts its advantage in instruction count and register footprint, not in launch overlap.
For the wgpu fallback this remains true in principle, but that engine now only runs on NVIDIA hosts without a CUDA driver, and touching origin's search_range for a fallback path is exactly the coupling #1 says to avoid.
Closing as superseded. Every arithmetic idea listed here (carry-free wide products, the row sum folded into the multiply-add, deferred-carry accumulation, constants as immediates) has been implemented and measured in the native CUDA engine instead, where it could be expressed exactly: #16, #17, #23 and #26 take the 4090 from 144 MH/s on the wgpu kernel to 716, and the 5090 from 283 to 1104 per card. On NVIDIA hosts the wgpu kernel is now only the fallback when no CUDA driver is present, and a 5 to 15% improvement to a path that runs at a quarter of the rate is not worth a 500-line duplicate of an origin file to maintain through merges.
If a non-NVIDIA card ever joins the fleet, the relevant question is what origin's kernel does on that vendor, not this port.
monday, september 14, 2026 — 12:28:34 utcClosing: everything this issue asked for is in place and has run in anger.
cuda-13.0 runner on every push to main, deployed with validate and rollback (#12, #13, #16); the validate step asserts the commit, the cuda kernel label and advancing per-device counters.main since, most recently run 65 on 8b01387.v4.0.2-lair.1).The one item never done, sccache against the fleet cache, stays undone on purpose: the miner builds in about 70 s and the bench harness in 30 s incremental, so there is nothing to recover. If runner time ever becomes the bottleneck, open a fresh issue for it.
8b01387 Merge pull request 'release: tag-triggered binaries and a README that states the measured rates' (#28) from release/binaries-and-readme into main7831323 release: tag-triggered binaries, README section with measured ratesc1cac91 engine-cuda: LAIR_MINBLOCKS launch-bounds knob, with the occupancy results7831323 release: tag-triggered binaries, README section with measured ratesUpstream's own engine measured for comparison. origin/main 0ba8ddf (v4.2.0, "CUDA: software-pipelined internal rounds, free-addend products, no early exit"), built on beast, quantus-miner benchmark --cuda-gpu --cpu-workers 0 --gpu-devices 1 --duration 30 on GPU 0 at the 400 W cap, miner paused, NVRTC from CUDA 13.0:
| engine | MH/s per 5090 |
|---|---|
| upstream wgpu (its default) | 278 |
| upstream --cuda-gpu | 944 |
| ours (#26), same window, one card | 1079 |
| quanpool-miner 6.2.4 | 1234 |
So upstream's CUDA engine is 12.5% slower than ours on this card, and quanpool's lead is their own work on top of it, not something we can merge from origin. Upstream CUDA commit titles are not performance evidence; measure before evaluating any of them for the fork.
sunday, september 13, 2026 — 22:03:50 utcThe LAIR_MINBLOCKS launch-bounds knob used for the register experiments is on branch perf/minblocks-knob (c1cac91), pushed but not merged: it changes no default, and merging would only trigger a rebuild-and-redeploy of the same kernel. Merge it with the next real kernel change.
c1cac91 engine-cuda: LAIR_MINBLOCKS launch-bounds knob, with the occupancy results17caeab Merge pull request 'perf: whole-grid batches, early reject on element 0, precomputed nonce direction (+2.9% on beast)' (#26) from perf/scheduling into main2e6fce1 engine-cuda: batches of whole grids, early reject on element 0, nonce directionProbe source for the record (also beast:/tmp/kregs/tcmul.cu). Build nvcc -O3 -arch=sm_120 [-DPROD_TMP|-DPROD_32|-DPROD_C|-DRUNTIME_DIAG] tcmul.cu; run ./tcmul t for the lane-0 trace, ./tcmul 1 5 for the checksum comparison, ./tcmul 1360 2000 for timing. CPU reference: the emulation in the #23 PR description reproduces the t trace for lane 0 (seed 17, splitmix as in the file).
// Go/no-go micro-benchmark for quantus/miner#24: 12 constant Goldilocks
// multiplies per thread per iteration (the internal-round diagonal layer),
// scalar IMAD.WIDE path (as shipped in #23) vs INT8 tensor-core path.
//
// Tensor mapping (the best found): mma.m16n8k16.u8 with A = 16 operands x 16
// limbs (8 used) from 16 lanes, B = Toeplitz of the constant's 8 limbs, two
// mma per 16 products give the 16 column sums c_k = sum_{i+j=k} a_i b_j, each
// < 2^19. Column sums land spread over the 4 lanes of a group; each lane forms
// its 128-bit partial and a 2-stage shuffle butterfly sums them, then the same
// fused reduction as the scalar path, then the result is shuffled back to the
// owning lane.
#include <cstdio>
#include <cstdint>
#include <cstdlib>
#include <cuda_runtime.h>
typedef unsigned long long u64; typedef unsigned u32;
__device__ __constant__ u64 DIAG[12] = {
0xc3b6c08e23ba9300ull, 0xd84b5de94a324fb6ull, 0x0d0c371c5b35b84full, 0x7964f570e7188037ull,
0x5daf18bbd996604bull, 0x6743bc47b9595257ull, 0x5528b9362c59bb70ull, 0xac45e25b7127b68bull,
0xa2077d7dfbb606b5ull, 0xf3faac6faee378aeull, 0x0c6388b51545e883ull, 0xd27dbb6944917b60ull };
__device__ __constant__ u32 EPSC = 0xFFFFFFFFu;
#define 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"
#ifdef PROD_TMP
// Same 32-bit chain, but every intermediate is a C temporary bound as an asm
// operand: no `.reg` declarations inside the block at all.
__device__ __forceinline__ u64 gf_mul(u64 a, u64 b, u32 e) {
u32 a0 = (u32)a, a1 = (u32)(a >> 32), b0 = (u32)b, b1 = (u32)(b >> 32);
u64 p0, m, m2, p3;
asm("mul.wide.u32 %0, %4, %6;\n\tmul.wide.u32 %1, %5, %6;\n\tmul.wide.u32 %2, %4, %7;\n\tmul.wide.u32 %3, %5, %7;"
: "=l"(p0), "=l"(m), "=l"(m2), "=l"(p3) : "r"(a0), "r"(a1), "r"(b0), "r"(b1));
u32 p0l = (u32)p0, p0h = (u32)(p0 >> 32), m0 = (u32)m, m1 = (u32)(m >> 32), m2l = (u32)m2, m2h = (u32)(m2 >> 32), p3l = (u32)p3, p3h = (u32)(p3 >> 32);
u32 cw, ll = p0l, lh, hl, hh, c;
asm("add.cc.u32 %0, %3, %5;\n\taddc.cc.u32 %1, %4, %6;\n\taddc.u32 %2, 0, 0;"
: "=r"(m0), "=r"(m1), "=r"(cw) : "r"(m0), "r"(m1), "r"(m2l), "r"(m2h));
asm("add.cc.u32 %0, %3, %4;\n\taddc.cc.u32 %1, %5, %6;\n\taddc.u32 %2, %7, %8;"
: "=r"(lh), "=r"(hl), "=r"(hh) : "r"(p0h), "r"(m0), "r"(p3l), "r"(m1), "r"(p3h), "r"(cw));
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"(c) : "r"(hl), "r"(hh), "r"(e));
return ((u64)lh << 32) | ll;
}
#elif defined(PROD_32)
// Product with 32-bit carry chains only: no addc.u64 after add.cc.u32 and no addc.u32 after add.cc.u64.
__device__ __forceinline__ u64 gf_mul(u64 a, u64 b, u32 e) {
u64 r;
asm("{\n\t"
".reg .b32 a0, a1, b0, b1, p0l, p0h, m0, m1, m2l, m2h, p3l, p3h, 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"
"mul.wide.u32 p0, a0, b0;\n\t" "mul.wide.u32 m, a1, b0;\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 {m0, m1}, m;\n\t" "mov.b64 {m2l, m2h}, m2;\n\t" "mov.b64 {p3l, p3h}, p3;\n\t"
"add.cc.u32 m0, m0, m2l;\n\t" "addc.cc.u32 m1, m1, m2h;\n\t" "addc.u32 cw, 0, 0;\n\t"
"add.cc.u32 lh, p0h, m0;\n\t" "addc.cc.u32 hl, p3l, m1;\n\t" "addc.u32 hh, p3h, cw;\n\t"
"mov.b32 ll, p0l;\n\t"
REDUCE_PTX("%3")
"mov.b64 %0, {ll, lh};\n\t" "}"
: "=l"(r) : "l"(a), "l"(b), "r"(e));
return r;
}
#elif defined(PROD_C)
__device__ __forceinline__ u64 gf_reduce4(u32 ll0, u32 lh0, u32 hl, u32 hh, u32 e);
__device__ __forceinline__ u64 gf_mul(u64 a, u64 b, u32 e) {
u64 lo = a * b, hi = __umul64hi(a, b);
return gf_reduce4((u32)lo, (u32)(lo >> 32), (u32)hi, (u32)(hi >> 32), e);
}
#else
__device__ __forceinline__ u64 gf_mul(u64 a, u64 b, u32 e) {
u64 r;
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"
REDUCE_PTX("%3")
"mov.b64 %0, {ll, lh};\n\t" "}"
: "=l"(r) : "l"(a), "l"(b), "r"(e));
return r;
}
#endif
__device__ __forceinline__ u64 gf_reduce4(u32 ll, u32 lh, u32 hl, u32 hh, u32 e) {
u32 c;
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"(c) : "r"(hl), "r"(hh), "r"(e));
return ((u64)lh << 32) | ll;
}
__device__ __forceinline__ u64 splitmix(u64& s) {
u64 z = (s += 0x9e3779b97f4a7c15ull);
z = (z ^ (z >> 30)) * 0xbf58476d1ce4e5b9ull;
z = (z ^ (z >> 27)) * 0x94d049bb133111ebull;
return z ^ (z >> 31);
}
__global__ void k_scalar(u64* out, int iters) {
u64 s = blockIdx.x * 1315423911ull + threadIdx.x * 2654435761ull + 17;
u64 v[12];
#pragma unroll
for (int i = 0; i < 12; i++) v[i] = splitmix(s);
u32 e = EPSC;
for (int it = 0; it < iters; it++) {
#pragma unroll
for (int i = 0; i < 12; i++) v[i] = gf_mul(v[i], DIAG[i], e);
}
u64 x = 0;
#pragma unroll
for (int i = 0; i < 12; i++) x ^= v[i] * (i + 1);
out[blockIdx.x * blockDim.x + threadIdx.x] = x;
}
__device__ __forceinline__ void mma_u8(int d[4], u32 a0, u32 a1, u32 b) {
asm volatile("mma.sync.aligned.m16n8k16.row.col.s32.u8.u8.s32 {%0,%1,%2,%3}, {%4,%5}, {%6}, {%0,%1,%2,%3};"
: "+r"(d[0]), "+r"(d[1]), "+r"(d[2]), "+r"(d[3]) : "r"(a0), "r"(a1), "r"(b));
}
__device__ __forceinline__ u32 limb(u64 x, int k) { return (k < 0 || k > 7) ? 0u : (u32)((x >> (8 * k)) & 0xFF); }
// One constant multiply of the 32 lanes' element v, returning the new value
// for this lane. b1/b2: this lane's B fragments for the constant.
__device__ __forceinline__ u64 tc_mul32(u64 v, u32 b1, u32 b2, u32 e, int lane) {
const int g = lane >> 2, t = lane & 3, s = 16 * t;
u64 res = 0;
#pragma unroll
for (int base = 0; base < 32; base += 16) {
// A fragment: rows g and g+8 are the operands of lanes base+g, base+g+8.
u64 xg = __shfl_sync(0xffffffffu, v, base + g);
u64 xh = __shfl_sync(0xffffffffu, v, base + g + 8);
u32 a0 = t == 0 ? (u32)xg : t == 1 ? (u32)(xg >> 32) : 0u;
u32 a1 = t == 0 ? (u32)xh : t == 1 ? (u32)(xh >> 32) : 0u;
int d1[4] = {0, 0, 0, 0}, d2[4] = {0, 0, 0, 0};
mma_u8(d1, a0, a1, b1); // columns c_0..c_7
mma_u8(d2, a0, a1, b2); // columns c_8..c_15
u64 r[2];
#pragma unroll
for (int p = 0; p < 2; p++) { // product g (p=0) and g+8 (p=1)
u32 lo = (u32)d1[2 * p] + ((u32)d1[2 * p + 1] << 8); // c_{2t} + c_{2t+1}*2^8 (< 2^28)
u32 hi = (u32)d2[2 * p] + ((u32)d2[2 * p + 1] << 8); // c_{8+2t} + c_{9+2t}*2^8
// 128-bit partial P = lo << s + hi << (64+s), as {plo, phi}
u64 plo = (u64)lo << s;
u64 phi = (s ? ((u64)lo >> (64 - s)) : 0ull) + ((u64)hi << s);
// butterfly over the 4 lanes of the group (xor 1, xor 2)
#pragma unroll
for (int m = 1; m <= 2; m <<= 1) {
u64 olo = __shfl_xor_sync(0xffffffffu, plo, m);
u64 ohi = __shfl_xor_sync(0xffffffffu, phi, m);
u64 nlo = plo + olo;
phi = phi + ohi + (nlo < plo ? 1ull : 0ull);
plo = nlo;
}
r[p] = gf_reduce4((u32)plo, (u32)(plo >> 32), (u32)phi, (u32)(phi >> 32), e);
}
// hand the results back to the owning lanes base..base+15
int me = lane - base; // 0..15 if in this batch
int src = 4 * (me & 7);
u64 y0 = __shfl_sync(0xffffffffu, r[0], src);
u64 y1 = __shfl_sync(0xffffffffu, r[1], src);
if (me >= 0 && me < 16) res = (me < 8) ? y0 : y1;
}
return res;
}
__global__ void k_tensor(u64* out, int iters) {
u64 s = blockIdx.x * 1315423911ull + threadIdx.x * 2654435761ull + 17;
u64 v[12];
#pragma unroll
for (int i = 0; i < 12; i++) v[i] = splitmix(s);
u32 e = EPSC;
const int lane = threadIdx.x & 31, g = lane >> 2, t = lane & 3;
// B fragments per constant: B1[k][n] = d_{n-k}, B2[k][n] = d_{n+8-k}; lane holds k=4t..4t+3, n=g.
u32 B1[12], B2[12];
#pragma unroll
for (int i = 0; i < 12; i++) {
u64 d = DIAG[i];
u32 b1 = 0, b2 = 0;
#pragma unroll
for (int j = 0; j < 4; j++) {
int k = 4 * t + j;
b1 |= limb(d, g - k) << (8 * j);
b2 |= limb(d, g + 8 - k) << (8 * j);
}
B1[i] = b1; B2[i] = b2;
}
for (int it = 0; it < iters; it++) {
#pragma unroll
for (int i = 0; i < 12; i++) v[i] = tc_mul32(v[i], B1[i], B2[i], e, lane);
}
u64 x = 0;
#pragma unroll
for (int i = 0; i < 12; i++) x ^= v[i] * (i + 1);
out[blockIdx.x * blockDim.x + threadIdx.x] = x;
}
__global__ void k_dbg2(int* out) {
u64 s = threadIdx.x * 2654435761ull + 17;
u64 v[12], w[12];
for (int i = 0; i < 12; i++) { v[i] = splitmix(s); w[i] = v[i]; }
u32 e = EPSC;
const int lane = threadIdx.x & 31, g = lane >> 2, t = lane & 3;
u32 B1[12], B2[12];
for (int i = 0; i < 12; i++) {
u64 d = DIAG[i]; u32 b1 = 0, b2 = 0;
for (int j = 0; j < 4; j++) { int k = 4 * t + j; b1 |= limb(d, g - k) << (8 * j); b2 |= limb(d, g + 8 - k) << (8 * j); }
B1[i] = b1; B2[i] = b2;
}
int first = -1;
for (int it = 0; it < 3 && first < 0; it++)
for (int i = 0; i < 12; i++) {
v[i] = gf_mul(v[i], DIAG[i], e);
w[i] = tc_mul32(w[i], B1[i], B2[i], e, lane);
if (first < 0 && v[i] != w[i]) first = it * 100 + i;
}
out[lane] = first;
}
__global__ void k_dbg3(u64* out) { // non-unrolled twin of the full kernels, block 0 seeds, 5 iterations, both paths
u64 s = threadIdx.x * 2654435761ull + 17;
u64 v[12], w[12];
for (int i = 0; i < 12; i++) { v[i] = splitmix(s); w[i] = v[i]; }
u32 e = EPSC;
const int lane = threadIdx.x & 31, g = lane >> 2, t = lane & 3;
u32 B1[12], B2[12];
for (int i = 0; i < 12; i++) {
u64 d = DIAG[i]; u32 b1 = 0, b2 = 0;
for (int j = 0; j < 4; j++) { int k = 4 * t + j; b1 |= limb(d, g - k) << (8 * j); b2 |= limb(d, g + 8 - k) << (8 * j); }
B1[i] = b1; B2[i] = b2;
}
for (int it = 0; it < 5; it++)
for (int i = 0; i < 12; i++) { v[i] = gf_mul(v[i], DIAG[i], e); w[i] = tc_mul32(w[i], B1[i], B2[i], e, lane); }
u64 x = 0, y = 0;
for (int i = 0; i < 12; i++) { x ^= v[i] * (i + 1); y ^= w[i] * (i + 1); }
out[threadIdx.x * 2] = x; out[threadIdx.x * 2 + 1] = y;
}
__global__ void k_trace(u64* out, int iters) { // unrolled like k_scalar; lane 0 of block 0 records v[i] after every multiply
u64 s = blockIdx.x * 1315423911ull + threadIdx.x * 2654435761ull + 17;
u64 v[12];
#pragma unroll
for (int i = 0; i < 12; i++) v[i] = splitmix(s);
u32 e = EPSC;
for (int it = 0; it < iters; it++) {
#pragma unroll
for (int i = 0; i < 12; i++) {
#ifdef RUNTIME_DIAG
v[i] = gf_mul(v[i], DIAG[i] ^ (u64)(threadIdx.x >> 31), e); // still DIAG[i] for every thread, but not a compile-time-visible operand
#else
v[i] = gf_mul(v[i], DIAG[i], e);
#endif
if (threadIdx.x == 0 && blockIdx.x == 0) out[it * 12 + i] = v[i]; }
}
}
__global__ void k_dbg(u64* out) {
u64 s = threadIdx.x * 2654435761ull + 17;
u64 v = splitmix(s);
u32 e = EPSC;
const int lane = threadIdx.x & 31, g = lane >> 2, t = lane & 3;
u32 b1 = 0, b2 = 0; u64 d = DIAG[0];
for (int j = 0; j < 4; j++) { int k = 4 * t + j; b1 |= limb(d, g - k) << (8 * j); b2 |= limb(d, g + 8 - k) << (8 * j); }
out[lane * 4 + 0] = v;
out[lane * 4 + 1] = gf_mul(v, d, e);
out[lane * 4 + 2] = tc_mul32(v, b1, b2, e, lane);
// reference: 128-bit product via __umul64hi then the same reduce
u64 lo = v * d, hi = __umul64hi(v, d);
out[lane * 4 + 3] = gf_reduce4((u32)lo, (u32)(lo >> 32), (u32)hi, (u32)(hi >> 32), e);
}
int main(int argc, char** argv) {
if (argc > 1 && argv[1][0] == 't') {
u64* d; cudaMalloc(&d, 5 * 12 * 8); k_trace<<<1, 256>>>(d, 5); u64 h[60]; cudaMemcpy(h, d, sizeof h, cudaMemcpyDeviceToHost);
for (int k = 0; k < 60; k++) printf("%d %d %016llx\n", k / 12, k % 12, h[k]);
return 0;
}
if (argc > 1 && argv[1][0] == 'f') {
u64* d; cudaMalloc(&d, 256 * 2 * 8); k_dbg3<<<1, 256>>>(d); u64 h[512]; cudaMemcpy(h, d, sizeof h, cudaMemcpyDeviceToHost);
for (int l = 0; l < 4; l++) printf("lane %d rolled scalar %016llx rolled tensor %016llx\n", l, h[2*l], h[2*l+1]);
return 0;
}
if (argc > 1 && argv[1][0] == 'e') {
int* d; cudaMalloc(&d, 32 * 4); k_dbg2<<<1, 32>>>(d); int h[32]; cudaMemcpy(h, d, sizeof h, cudaMemcpyDeviceToHost);
for (int l = 0; l < 32; l++) printf("lane %2d first mismatch (iter*100+elem): %d\n", l, h[l]);
return 0;
}
if (argc > 1 && argv[1][0] == 'd') {
u64* d; cudaMalloc(&d, 32 * 4 * 8); k_dbg<<<1, 32>>>(d); u64 h[128]; cudaMemcpy(h, d, sizeof h, cudaMemcpyDeviceToHost);
for (int l = 0; l < 32; l++) printf("lane %2d v=%016llx scalar=%016llx tensor=%016llx ref=%016llx %s\n", l, h[l*4], h[l*4+1], h[l*4+2], h[l*4+3], h[l*4+1]==h[l*4+2] ? "ok" : "DIFF");
return 0;
}
int blocks = argc > 1 ? atoi(argv[1]) : 170 * 8, threads = 256;
int iters = argc > 2 ? atoi(argv[2]) : 2000;
size_t n = (size_t)blocks * threads;
u64 *d_a, *d_b; cudaMalloc(&d_a, n * 8); cudaMalloc(&d_b, n * 8);
// correctness: few iterations, compare checksums lane by lane
k_scalar<<<blocks, threads>>>(d_a, 5); k_tensor<<<blocks, threads>>>(d_b, 5);
cudaDeviceSynchronize();
u64* h_a = (u64*)malloc(n * 8); u64* h_b = (u64*)malloc(n * 8);
cudaMemcpy(h_a, d_a, n * 8, cudaMemcpyDeviceToHost); cudaMemcpy(h_b, d_b, n * 8, cudaMemcpyDeviceToHost);
size_t bad = 0; for (size_t i = 0; i < n; i++) if (h_a[i] != h_b[i]) bad++;
printf("correctness: %zu/%zu lanes differ between scalar and tensor paths\n", bad, n);
for (size_t i = 0; i < 4; i++) printf(" lane %zu scalar %016llx tensor %016llx\n", i, h_a[i], h_b[i]);
cudaEvent_t e0, e1; cudaEventCreate(&e0); cudaEventCreate(&e1);
for (int rep = 0; rep < 3; rep++) {
float ms_s, ms_t;
cudaEventRecord(e0); k_scalar<<<blocks, threads>>>(d_a, iters); cudaEventRecord(e1); cudaEventSynchronize(e1);
cudaEventElapsedTime(&ms_s, e0, e1);
cudaEventRecord(e0); k_tensor<<<blocks, threads>>>(d_b, iters); cudaEventRecord(e1); cudaEventSynchronize(e1);
cudaEventElapsedTime(&ms_t, e0, e1);
double prods = (double)n * iters * 12;
printf("rep %d scalar %8.2f ms %7.1f Gmul/s tensor %8.2f ms %7.1f Gmul/s ratio tensor/scalar time %.2fx\n",
rep, ms_s, prods / ms_s / 1e6, ms_t, prods / ms_t / 1e6, ms_t / ms_s);
}
cudaError_t err = cudaGetLastError(); if (err != cudaSuccess) printf("cuda error: %s\n", cudaGetErrorString(err));
return 0;
}
sunday, september 13, 2026 — 20:19:26 utcPost-deploy guard run on beast with the #23 kernel (bench-harness, miner paused): parity OK, 2000/2000 jobs found solutions, all verified against CPU. Each job is a random 32-byte header and random 512-bit start nonce at difficulty 100,000, so that is 2000 full two-permutation hashes compared bit-for-bit with pow_core. The anomaly does not reach the shipped kernel's code shape.
sunday, september 13, 2026 — 20:19:11 utcRan step 1 on beast (one 5090, miner paused, 348,160 threads x 2000 iterations x 12 constant multiplies each, both paths ending in the same fused Goldilocks reduction).
Tensor mapping tested (the best I could construct): mma.sync.m16n8k16.u8 with A = 16 operands x 8 limbs from 16 lanes, B = Toeplitz of the constant's limbs, two mma per 16 products yielding the 16 column sums (each < 2^19, verified against a CPU reference on all 32 lanes); column sums land spread over the 4 lanes of a group, each lane forms its 128-bit partial, a 2-stage shuffle butterfly sums them, then the fused reduction, then a shuffle back to the owning lane. Verified bit-exact against the scalar path and a Python reference in the rolled build.
| path | time (2000 x 12 x 348k products) | rate | |---|---|---| | scalar IMAD.WIDE + fused reduce | 3.8-3.9 ms | ~2,150 Gmul/s | | INT8 mma path | 60.6-61.6 ms | ~137 Gmul/s |
15-16x slower, against a go threshold of 2x faster. Static SASS per kernel: tensor 2,832 instructions of which 576 SHFL, 413 SHF, 346 IADD, 322 IMAD; scalar 2,464 (995 IMAD, 756 IADD3). The structural reason: one m16n8k16 covers 8 useful products because a 64-bit operand only fills half of K and 15 of 16 output rows, so per warp-instruction it does no more products than one IMAD.WIDE does across 32 lanes, while the mma pipe issues slower than the IMAD pipe on this part, and the scatter of column sums across lanes costs about 40 shuffle-and-carry instructions per 16 products on top. Using more of the tile would need packing two operands into K, which mixes the products. FP16 with fp32 accumulate needs 10-bit limbs and runs at half the int8 rate, so it is worse.
Not pursuing further. Closing.
Probe source: beast:/tmp/kregs/tcmul.cu (build nvcc -O3 -arch=sm_120 -DPROD_TMP, modes t/e/f/d for traces, <blocks> <iters> for timing). While building it I hit a ptxas anomaly unrelated to tensor cores; filed separately.
dc6c58c Merge pull request 'perf: fused-PTX field arithmetic, +54..66% on every card' (#23) from perf/fused-field-ptx into main3763996 engine-cuda: fused-PTX field arithmetic, +54..66% on every card3763996 engine-cuda: fused-PTX field arithmetic, +54..66% on every card78dee04 Merge pull request 'perf: parameter-bank uniforms and per-arch carry path (+20.5% sm_120)' (#22) from perf/sm120-spills into main575c4d0 engine-cuda: per-arch carry path, +20.5% on sm_120206aa0d engine-cuda: record why gf_sqr uses the general product7804d9d engine-cuda: launch uniforms in the parameter bank, no spills575c4d0 engine-cuda: per-arch carry path, +20.5% on sm_120206aa0d engine-cuda: record why gf_sqr uses the general product7804d9d engine-cuda: launch uniforms in the parameter bank, no spills