Found while running the #24 go/no-go probe (beast:/tmp/kregs/tcmul.cu). Recording it because the mining kernel from #23 uses the same inline-PTX idiom.
Symptom
A probe kernel does, per thread, 12 independent v[i] = gf_mul(v[i], DIAG[i]) chains in a #pragma unroll loop, iterated. gf_mul is the #23 asm block verbatim (four mul.wide.u32, carry assembly, mad.lo.cc/madc.hi.cc reduction). With nvcc -O3 -arch=sm_120 (CUDA 13.0):
3 of the 12 chains (elements 2, 6, 7) produce wrong values deterministically, from the first multiply, for every lane. The wrong values are not any simple perturbation (not a dropped carry, not a wrong operand, not a wrong EPS).
-Xptxas -O0: correct. -Xptxas -O1 and above: wrong. So the PTX is right and ptxas's optimiser changes the result.
#pragma unroll 1 on the element loop: correct. -G: correct.
Replacing only the product with C (a*b, __umul64hi) and keeping the PTX reduction: correct.
Rewriting the block with every intermediate as a C temporary bound as an asm operand (no .reg declarations inside {}): correct in one kernel context (k_trace, which stores each value) and still wrong in the twin kernel that only stores a checksum at the end. So it is code-shape dependent, and scoped .reg names are not the whole story.
Hypothesis, unproven: predicate pressure. Twelve interleaved carry chains need more live carries than the 7 predicate registers, and something in ptxas's predicate spilling or CC-flag tracking goes wrong at -O1+. The deployed kernel's internal layer is also 12 independent multiply-add chains, but interleaved with a long dependent S-box chain.
Why the deployed kernel is believed fine
bench-harness parity 40/40 on beast, benjy and quadbrat before #23 merged (each job compares a full GPU hash of a random header/nonce to pow_core; a per-multiply corruption of this kind would fail every hash).
A 2000-job parity run after the deploy (result recorded below).
The mining hosts are earning blocks on the #23 build.
To do
Reduce the probe to a minimal repro and check newer CUDA (13.1/13.2 are installed on benjy/beast) and -Xptxas -O2 vs -O3.
Consider adding an exhaustive hash-parity mode to bench-harness: an engine method or debug kernel that returns hashes for a nonce range so thousands of hashes can be compared per second, instead of one solution per job.
Until the cause is known: every kernel change must pass parity on all three architectures (already the bench.yaml flow), and any new asm block should be written in the operand-temporary style with fewer independent chains per unrolled region where possible.
Found while running the #24 go/no-go probe (beast:/tmp/kregs/tcmul.cu). Recording it because the mining kernel from #23 uses the same inline-PTX idiom.
## Symptom
A probe kernel does, per thread, 12 independent `v[i] = gf_mul(v[i], DIAG[i])` chains in a `#pragma unroll` loop, iterated. `gf_mul` is the #23 asm block verbatim (four `mul.wide.u32`, carry assembly, `mad.lo.cc`/`madc.hi.cc` reduction). With `nvcc -O3 -arch=sm_120` (CUDA 13.0):
- 3 of the 12 chains (elements 2, 6, 7) produce wrong values deterministically, from the first multiply, for every lane. The wrong values are not any simple perturbation (not a dropped carry, not a wrong operand, not a wrong EPS).
- `-Xptxas -O0`: correct. `-Xptxas -O1` and above: wrong. So the PTX is right and ptxas's optimiser changes the result.
- `#pragma unroll 1` on the element loop: correct. `-G`: correct.
- Replacing only the product with C (`a*b`, `__umul64hi`) and keeping the PTX reduction: correct.
- Rewriting the block with every intermediate as a C temporary bound as an asm operand (no `.reg` declarations inside `{}`): correct in one kernel context (`k_trace`, which stores each value) and still wrong in the twin kernel that only stores a checksum at the end. So it is code-shape dependent, and scoped `.reg` names are not the whole story.
Hypothesis, unproven: predicate pressure. Twelve interleaved carry chains need more live carries than the 7 predicate registers, and something in ptxas's predicate spilling or CC-flag tracking goes wrong at -O1+. The deployed kernel's internal layer is also 12 independent multiply-add chains, but interleaved with a long dependent S-box chain.
## Why the deployed kernel is believed fine
- bench-harness parity 40/40 on beast, benjy and quadbrat before #23 merged (each job compares a full GPU hash of a random header/nonce to pow_core; a per-multiply corruption of this kind would fail every hash).
- A 2000-job parity run after the deploy (result recorded below).
- The mining hosts are earning blocks on the #23 build.
## To do
- Reduce the probe to a minimal repro and check newer CUDA (13.1/13.2 are installed on benjy/beast) and `-Xptxas -O2` vs `-O3`.
- Consider adding an exhaustive hash-parity mode to bench-harness: an engine method or debug kernel that returns hashes for a nonce range so thousands of hashes can be compared per second, instead of one solution per job.
- Until the cause is known: every kernel change must pass parity on all three architectures (already the bench.yaml flow), and any new asm block should be written in the operand-temporary style with fewer independent chains per unrolled region where possible.
🤖 Generated with [Claude Code](https://claude.com/claude-code)
https://claude.ai/code/session_01Ue5ZZm1Hiv5zPnucykKKuF
Post-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.
Post-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.
Probe 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;
}
Probe 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).
```cuda
// 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;
}
```
Blocking a user prevents them from interacting with repositories, such as opening or commenting on pull requests or issues. Learn more about blocking a user.
Found while running the #24 go/no-go probe (beast:/tmp/kregs/tcmul.cu). Recording it because the mining kernel from #23 uses the same inline-PTX idiom.
Symptom
A probe kernel does, per thread, 12 independent
v[i] = gf_mul(v[i], DIAG[i])chains in a#pragma unrollloop, iterated.gf_mulis the #23 asm block verbatim (fourmul.wide.u32, carry assembly,mad.lo.cc/madc.hi.ccreduction). Withnvcc -O3 -arch=sm_120(CUDA 13.0):-Xptxas -O0: correct.-Xptxas -O1and above: wrong. So the PTX is right and ptxas's optimiser changes the result.#pragma unroll 1on the element loop: correct.-G: correct.a*b,__umul64hi) and keeping the PTX reduction: correct..regdeclarations inside{}): correct in one kernel context (k_trace, which stores each value) and still wrong in the twin kernel that only stores a checksum at the end. So it is code-shape dependent, and scoped.regnames are not the whole story.Hypothesis, unproven: predicate pressure. Twelve interleaved carry chains need more live carries than the 7 predicate registers, and something in ptxas's predicate spilling or CC-flag tracking goes wrong at -O1+. The deployed kernel's internal layer is also 12 independent multiply-add chains, but interleaved with a long dependent S-box chain.
Why the deployed kernel is believed fine
To do
-Xptxas -O2vs-O3.🤖 Generated with Claude Code
https://claude.ai/code/session_01Ue5ZZm1Hiv5zPnucykKKuF
Post-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.
Probe 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 tfor the lane-0 trace,./tcmul 1 5for the checksum comparison,./tcmul 1360 2000for timing. CPU reference: the emulation in the #23 PR description reproduces thettrace for lane 0 (seed 17, splitmix as in the file).