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

ptxas 13.0 -O1+ miscompiles unrolled inline-PTX carry chains in a probe kernel; deployed kernel unaffected by parity, root cause open #25

New Issue
Open
opened 2026-09-13 20:18:38 +00:00 by grenade · 2 comments
grenade commented 2026-09-13 20:18:38 +00:00
Owner
Copy Link

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.ai/code/session_01Ue5ZZm1Hiv5zPnucykKKuF

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
grenade commented 2026-09-13 20:19:11 +00:00
Author
Owner
Copy Link

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.
grenade commented 2026-09-13 20:19:25 +00:00
Author
Owner
Copy Link

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; } ```
Sign in to join this conversation.
No Branch/Tag Specified
Branches Tags
Labels
Clear labels
needs-benchmark
Gain is estimated, not measured; needs a run on a 5090 host
 needs-chain
May require a matching change in quantus/chain
 origin-coupling
Touches how the lair fork tracks Quantus-Network/quantus-miner
 perf
Mining throughput workstream
No Label
Milestone
No items
No Milestone
Projects
Clear projects
No project
Assignees
Clear assignees
grenade
No Assignees
1 Participants
Notifications
Due Date
No due date set.
Dependencies

No dependencies set.

Reference: blackbeard/miner#25
Delete Branch "%!s()"

Deleting a branch is permanent. Although the deleted branch may continue to exist for a short time before it actually gets removed, it CANNOT be undone in most cases. Continue?