research: Goldilocks multiplies on the INT8 tensor cores #24

New Issue

grenade · 2026-09-13T19:55:57Z

grenade commented

2026-09-13 19:55:57 +00:00

Why

After #23 the CUDA kernel is at instruction parity with the closed-source qpow-cuda kernel (22.9k vs 22.2k instructions per nonce on sm_120) and both sit at ~15 instructions per field multiply, 1472 multiplies per nonce, power-bound at the 400 W cap. ncu on beast: FMA pipe 114 SM-cycles/nonce, ALU pipe 238, measured 354, i.e. the pipes barely overlap and math-pipe throttle is the top stall in both kernels. Every remaining scalar lever is a few percent (launch shape, nonce-direction precompute, felt-0-only rejection, whole-S-box asm blocks). Nothing scalar gets a 5090 much past ~1.3 GH/s.

The one idea with a large ceiling is to take the 64x64 products off the IMAD pipe. Blackwell's INT8 tensor throughput is tens of times the scalar IMAD.WIDE rate, and the multiplies are about a third of the instructions and probably more than half the energy per hash.

Idea

Batch the products of a warp (or of several nonces per thread) as small integer matrix multiplies on 8-bit limbs: a 64-bit operand is 8 limbs, a 64x64 product is an 8x8 outer product of limbs accumulated into 32-bit columns with carries resolved afterwards. Published schemes exist for big-integer and modular multiplication on INT8 tensor cores (search "tensor core big integer multiplication", "modular multiplication tensor cores"); none has been shown for Goldilocks at this granularity.

Points to settle before building anything:

Mapping: which products can be batched together. The S-box chain x, x^2, x^3, x^4, x^7 is serial per element but 12 elements per external round (and 32 lanes) are independent; the internal-round diagonal multiplies are 12 independent constant multiplies. That is 48 or 12 products per round per thread to tile.
The scalar remainder: carry resolution of the 8-bit column sums and the Goldilocks reduction stay on the ALU pipe, which is already the busier one (15.2k ALU instructions per nonce). The win exists only if the limb accumulation plus reduction costs less ALU than today's ~10 ALU per multiply.
Fragment shuffling: moving operands into and out of mma fragments each round may cost as much as the multiply it replaces. Register pressure (110 today, 128 for the rival) leaves little room.
Power: the cards are power-bound, so the metric is energy per hash, not instructions. Tensor ops are efficient per MAC but only if the tiles are full.

How to evaluate cheaply

Micro-benchmark: 12 independent 64x64 products per thread, scalar IMAD.WIDE vs INT8 mma, measured in SM-cycles and watts on beast, before touching the kernel. If the mma path is not at least ~2x cheaper per product including limb packing and carry resolution, stop.
If it passes, prototype the internal-round diagonal multiply first (constants are fixed, so one operand's fragments are precomputed), behind a knob, and check with the bench harness's parity jobs.

Same method as the rest of the perf work: screen statically with cuobjdump -sass counts, confirm on hardware with interleaved rounds, never trust an isolated probe with warp-uniform inputs.

Related: #3 (engine-cuda), #23 (fused-PTX arithmetic, the current state of the art).

🤖 Generated with Claude Code

https://claude.ai/code/session_01Ue5ZZm1Hiv5zPnucykKKuF

## Why After #23 the CUDA kernel is at instruction parity with the closed-source qpow-cuda kernel (22.9k vs 22.2k instructions per nonce on sm_120) and both sit at ~15 instructions per field multiply, 1472 multiplies per nonce, power-bound at the 400 W cap. ncu on beast: FMA pipe 114 SM-cycles/nonce, ALU pipe 238, measured 354, i.e. the pipes barely overlap and math-pipe throttle is the top stall in both kernels. Every remaining scalar lever is a few percent (launch shape, nonce-direction precompute, felt-0-only rejection, whole-S-box asm blocks). Nothing scalar gets a 5090 much past ~1.3 GH/s. The one idea with a large ceiling is to take the 64x64 products off the IMAD pipe. Blackwell's INT8 tensor throughput is tens of times the scalar IMAD.WIDE rate, and the multiplies are about a third of the instructions and probably more than half the energy per hash. ## Idea Batch the products of a warp (or of several nonces per thread) as small integer matrix multiplies on 8-bit limbs: a 64-bit operand is 8 limbs, a 64x64 product is an 8x8 outer product of limbs accumulated into 32-bit columns with carries resolved afterwards. Published schemes exist for big-integer and modular multiplication on INT8 tensor cores (search "tensor core big integer multiplication", "modular multiplication tensor cores"); none has been shown for Goldilocks at this granularity. Points to settle before building anything: - Mapping: which products can be batched together. The S-box chain x, x^2, x^3, x^4, x^7 is serial per element but 12 elements per external round (and 32 lanes) are independent; the internal-round diagonal multiplies are 12 independent constant multiplies. That is 48 or 12 products per round per thread to tile. - The scalar remainder: carry resolution of the 8-bit column sums and the Goldilocks reduction stay on the ALU pipe, which is already the busier one (15.2k ALU instructions per nonce). The win exists only if the limb accumulation plus reduction costs less ALU than today's ~10 ALU per multiply. - Fragment shuffling: moving operands into and out of mma fragments each round may cost as much as the multiply it replaces. Register pressure (110 today, 128 for the rival) leaves little room. - Power: the cards are power-bound, so the metric is energy per hash, not instructions. Tensor ops are efficient per MAC but only if the tiles are full. ## How to evaluate cheaply 1. Micro-benchmark: 12 independent 64x64 products per thread, scalar IMAD.WIDE vs INT8 mma, measured in SM-cycles and watts on beast, before touching the kernel. If the mma path is not at least ~2x cheaper per product including limb packing and carry resolution, stop. 2. If it passes, prototype the internal-round diagonal multiply first (constants are fixed, so one operand's fragments are precomputed), behind a knob, and check with the bench harness's parity jobs. Same method as the rest of the perf work: screen statically with `cuobjdump -sass` counts, confirm on hardware with interleaved rounds, never trust an isolated probe with warp-uniform inputs. Related: #3 (engine-cuda), #23 (fused-PTX arithmetic, the current state of the art). 🤖 Generated with [Claude Code](https://claude.com/claude-code) https://claude.ai/code/session_01Ue5ZZm1Hiv5zPnucykKKuF

grenade commented

2026-09-13 20:18:08 +00:00

Go/no-go micro-benchmark: NO-GO

Ran 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.

## Go/no-go micro-benchmark: NO-GO Ran 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.

grenade closed this issue

2026-09-13 20:18:21 +00:00

grenade referenced this issue

2026-09-13 20:18:38 +00:00

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

grenade referenced this issue

2026-09-14 12:28:53 +00:00

Architecture: loose coupling to origin for a performance-tuned miner #1

grenade referenced this issue

2026-09-14 19:15:20 +00:00

quanpool-miner 6.2.4 measured on beast: 1234 MH/s per 5090, what it does differently, what was tried #27

Sign in to join this conversation.