# ProxiChorizo — platform notes and measured optimisations

The work per attempt: 1 SHA-512 + try-and-increment hash-to-curve (≈2 tries,
each a decompression = one field exponentiation), one 254-bit scalar
multiplication on Ed25519 (extended coordinates, a = −1), ×8 cofactor, one
inversion for compression, 1 SHA-512. ≈3,600 field multiplications per attempt.

## Build-time variants (macros in proxichorizo.cu)

| Macro | Values | Meaning |
|---|---|---|
| `PX_FE32` | unset / defined | field element as 5×51-bit limbs with 64-bit multiplies (unset) or 8×32-bit limbs with PTX `mad.lo.cc/madc.hi.cc` carry chains and a dedicated squaring (defined) |
| `PX_WINDOW` | 0, 2, 4 | scalar multiplication: plain double-and-add (0), 2-bit window with a 3-point table, 4-bit window with a 16-point table in local memory |
| `PX_LB_THREADS`, `PX_LB_BLOCKS` | 256, 1/2 | `__launch_bounds__` for mine_kernel; 2 blocks/SM caps registers at 128 and raises occupancy |
| `PX_SQ_MUL` | defined | debugging: squaring falls back to multiplication |
| `PX_SHA_OLD` | defined | debugging: SHA-512 with the 80-word schedule instead of the rolling 16-word one |

## Measurements, 2026-10-19 (German test rig, driver 615.71, CUDA 13.4; the 3070s sit on x1/x4 gen1 risers)

**Correction:** the numbers in the first table were taken with `--bench` before the
fix of 10:45 UTC, when a bench without a key used an all-zero scalar, so the scalar
multiplication skipped every point addition and rates read ~1.45× too high. Relative
comparisons between field implementations still hold; comparisons that change the
number of additions (window size) do not. The second table is with a realistic scalar.

### Honest numbers (realistic scalar), RTX 5080

| Build | MH/s |
|---|---|
| fe51, window 4, lb1 (first version) | 6.69 |
| fe51, window 0, lb2 | 6.37 |
| fe32 + dedicated squaring, window 0, lb2 | 8.67 |
| fe32 + sq, window 2, lb2 | 6.78 |
| **fe32 + sq, window 4, lb2** | **9.39** |
| fe32 + sq, window 4, lb1 | 7.85 |
| fe32 + sq, window 2, lb1 | 6.17 |

### Honest numbers (realistic scalar), RTX 3070 (sm_86, the other 3070 mining)

| Build | MH/s |
|---|---|
| **fe32 + sq, window 0, lb2** | **2.85** |
| fe32 + sq, window 4, lb2 | 2.77 |
| fe32 + sq, window 2, lb2 | 1.97 |
| fe32 + sq, window 2, lb1 | 1.95 |

So the best window differs by generation: Blackwell (sm_120) wants the 4-bit table
(+8 % over plain), Ampere (sm_86) plain double-and-add (+3 %). The 2-bit window's
three-way branch is bad everywhere. This is the per-platform table the single
binary must carry.

### First-round table (zero-scalar bench, inflated ~1.45×)

| Build | RTX 5080 (sm_120) | RTX 3070 (sm_86, other 3070 mining) |
|---|---|---|
| fe51, window 4, lb1 (first version) | 7.95 MH/s | 2.21 (idle rig: 2.4–2.6) |
| fe51, window 0 | 8.55 | |
| fe51, window 2 | 8.37 | |
| fe51, window 0, lb2 | 9.26 | 1.70 |
| fe51, window 2, lb2 | 9.33 | |
| fe32 (mul-based squaring), window 0, lb2 | 8.97 | 2.21 |
| **fe32 + dedicated squaring, window 0, lb2** | **12.50** | **4.25** |
| same, threads 256 iters 4 | 12.59 | |
| same, threads 128 iters 4 | 12.12 | |
| same, threads 192 iters 2 | 11.64 | |

Register/stack (ptxas, sm_120): first version 255 regs, 3168 B stack, spills;
fe32+sq+lb2: 128 regs, 368 B stack, no spills.

Lessons:
- 64-bit integer multiplies are emulated on NVIDIA; the 8×32 carry-chain schoolbook is ~1.5–2× faster per field multiply. Blackwell only benefits once the squaring is dedicated (the mul-based fe32 was 3 % slower there).
- A multi-instruction inline-asm statement with many `"+r"` operands that hold **equal values** (e.g. zero limbs) can get those operands coalesced into one register by the compiler, so writes collide. Keep statements small (one limb pair per statement) and carry explicitly; always use `"=&r"` for pure outputs. The random field test (`--fetest`) did not catch it because the identity checks passed on random inputs; the sqrt(−1) derivation in `init_kernel` did.
- The 16-point window table lives in local memory (2.5 KB/thread) and costs more in occupancy than it saves in additions; plain double-and-add wins on both generations.
- `__launch_bounds__(256, 2)` (≤128 registers) is a free 7–10 %.
- Divergence in the try-and-increment loop (geometric number of tries, warp waits for the slowest lane) is the next structural cost (~25 %): compaction into a separate scalar-multiplication pass is the planned fix.

## Round 2: compacted hash-to-curve (2026-10-19 11:30 UTC)

The try-and-increment loop made warps wait for their slowest lane (expected
≈6 decompressions per warp for ≈2 needed per lane). `PX_COMPACT` splits an
iteration into `h2c_kernel` passes (one SHA-512 + one decompression per item,
failures queued for the next pass, 12 passes max) and `mult_kernel` (×8, scalar
multiplication, ×8, compression, SHA-512) over the found points with no
divergence; both kernels also carry far fewer registers than the fused one.
`--verify` proves every variant finds identical hits over one nonce range.

| GPU | fe32_w4_lb2 | fe32_w0_lb2_c | **fe32_w4_lb2_c** (256×4) |
|---|---|---|---|
| RTX 5080 | 9.41 | 14.58 | **16.60 MH/s** |
| RTX 3070 (other 3070 mining) | 2.79 | 5.16 | **5.98 MH/s** |

Versus the first version (6.7 / ~2.1): **2.5× on Blackwell, ~2.8× on Ampere.**
With the divergence gone the 4-bit window wins on both generations, so the
platform table now defaults every architecture to `fe32_w4_lb2_c` 256×4.

## Round 3: wNAF scalar multiplication (2026-10-19 12:00 UTC)

`PX_WINDOW=5` = width-5 wNAF of the constant scalar (digits computed on the host in
`v_set_scalar`, kept in constant memory), 8 odd multiples cached as (Y+X, Y−X, 2dT, 2Z)
for mixed additions, negation by swapping the sums. ~43 additions instead of 63 + 14
table operations. Variant `fe32_wnaf_lb2_c`: RTX 5080 16.6 → **17.3 MH/s**, RTX 3070
6.0 → **6.2 MH/s** (+4 %). Now the platform default everywhere.

## CMP 170HX (GA100, sm_80) — 2026-10-19 12:50 UTC

Autotune on a CMP 170HX (PCIe gen2 x1 riser): `fe32_w4_lb2_c` 256×2 **4.7 MH/s** (another
unit of the same card: 6.6 MH/s — the cards differ in clocks/power state), `fe32_wnaf_lb2_c`
4.1, `fe32_w0_lb2_c` 3.6, fused kernels 1.3–1.7. So GA100 prefers the 4-bit table over
wNAF, the opposite of GA10x/Blackwell; the platform table has its own row for 8.0.
HBM2 bandwidth plays no role — the work is pure integer ALU; the x1 link does not
matter either (a job is 200 bytes).

## Round 4: fewer multiplications per doubling (2026-10-19 14:30 UTC)

A doubling reads X, Y, Z only; T is needed solely as the input of the next
addition. In the wNAF loop the digit is known (constant memory, uniform), so a
doubling followed by another doubling skips T (3M + 4S instead of 4M + 4S); the
same in the 4-bit and plain variants and in the cofactor ×8. RTX 3070:
`fe32_wnaf_lb2_c` 6.55 → **6.87 MH/s** (+4.9 %), hits identical under `--verify`.

Tried and dropped the same day (RTX 3070, `fe32_wnaf_lb2_c` 256×4):
- two points per thread sharing one inversion (Montgomery's trick): 6.83 vs 6.86 —
  the first point parked in local memory while the second is multiplied costs what
  the saved ~265-operation inversion brings;
- the multiply kernel at 1 block/SM (`PX_LB_MULT_BLOCKS 1`, 148 registers, no
  spills): 5.54 at 256 threads, 6.38 at 128 — occupancy beats registers here.
  The macro stays for other platforms.

ptxas, sm_120, `fe32_wnaf_lb2_c`: h2c_kernel 118 registers, 128 B stack; mult_kernel
128 registers (the cap of 2 blocks/SM), 1 KB stack = the 8 cached odd multiples, 12 B
spills. The table in local memory is the remaining structural cost of wNAF.

## One binary, all variants (since 2026-10-19 11:00 UTC)

`proxichorizo.cu` is the host; `proxichorizo_device.cuh` is the whole device side and is included
once per variant inside a namespace with the macros set. Variants compiled in:

| name | field | window | launch bounds |
|---|---|---|---|
| `fe32_w0_lb2` | 8×32 carry chains + dedicated squaring | plain double-and-add | 256 threads, 2 blocks/SM |
| `fe32_w4_lb2` | same | 4-bit table | same |
| `fe51_w4_lb1` | 5×51 64-bit (the first version) | 4-bit table | 256, 1 |
| `fe32_w0_lb2_c` | 8×32 + sq, compacted pipeline | plain | 256, 2 |
| `fe32_w4_lb2_c` | 8×32 + sq, compacted pipeline | 4-bit table | 256, 2 |
| `fe32_wnaf_lb2_c` | 8×32 + sq, compacted pipeline | width-5 wNAF, cached mixed adds | 256, 2 |

Selection: `PLATFORMS[]` in proxichorizo.cu maps compute capability → variant and
geometry (12.x → fe32_w4_lb2; 8.x/7.5 → fe32_w0_lb2); `--variant NAME` forces one
for all GPUs; `--autotune [secs]` measures every variant × {256×2, 128×4, 256×4} on
each GPU and takes the best; `--list-variants`, `--selftest` and `--fetest` run all.
Distribution build: `-gencode` for sm_75, sm_80, sm_86, sm_89, sm_120 plus
compute_120 PTX for newer cards. Adding a platform = a row in `PLATFORMS[]`;
adding a kernel idea = a macro in proxichorizo_device.cuh and a namespace block.

## AMD / OpenCL port (2026-10-19, `opencl/`)

Same algorithm in OpenCL C 1.2 (`opencl/proxichorizo.cl`), compiled per GPU and variant by the
installed driver at start-up (45–50 s per variant on ROCm; cached in the process). The host
(`opencl/proxichorizo_cl.cpp`) shares the HTTP/job/submit/stats code with the CUDA miner
through `host_common.inc`, generated from `proxichorizo.cu` by `gen_host.py`. Variants:
`fe32_*` = 8×32-bit limbs with 64-bit accumulators (comba), `fe51_*` = 5×51 with `mul_hi`;
`_c` = compacted pipeline; `wnaf`/`w4`/`w0` as on CUDA. Every selected variant must reproduce
the 16 embedded test vectors on the device before it is used (`prog_validated`): ROCm 6.4's
compiler miscompiles the fe51 layout (last byte of the compressed point wrong), ROCm 7.14 does
not; if the platform default fails validation the miner falls back to `fe32_wnaf_c`.

### First round (fe32 / fe51)

| GPU | driver | best variant | MH/s |
|---|---|---|---:|
| RX 7900 XTX (gfx1100, 48 WGP) | ROCm 7.14 OpenCL | fe51_wnaf_c ×8 5.49 / fe32_wnaf_c ×8 5.36 | ≈ 5.4 |
| RX 7600 XT (gfx1102) | ROCm 6.4 | fe32_wnaf_c ×8 | 1.9 |
| MI50 (gfx906, 60 CU) | ROCm 6.4 | fe32_wnaf_c ×8 | 1.9 |

ISA (gfx1100, fe32_wnaf_c): mult_kernel 248 VGPRs, 1.4 KB scratch/lane (the cached wNAF table
and byte arrays), occupancy 5/16, the field multiply compiles to `v_mad_u64_u32` (4127 of them).
RDNA executes 32×32→64 multiply-adds at a quarter of the ALU rate, so the realistic ceiling of a
7900 XTX is well below a comparable NVIDIA part; `-O3` changes nothing. Word-based SHA-512
message assembly (no byte arrays in scratch) brought +3 %.

### Second round: 9×29-bit limbs (`fe29_*`, now the AMD default)

The 8×32 comba needs a carry chain after every 64-bit accumulate because the column sums
exceed 64 bits. With 9 limbs of 29 bits every one of the 81 partial products is < 2^58, so a
whole column (≤ 9 products, or 2^30-limb operands after one lazy addition: 9 × 2^60 < 2^64)
fits a 64-bit accumulator and the schoolbook is 81 bare `v_mad_u64_u32` with no carry handling
at all; the 17 columns are folded once at the end (2^261 ≡ 1216 mod p). Subtraction adds 1024p
limb-wise before subtracting so operands with 30/31-bit limbs stay non-negative (8p was not
enough: its top limb is 2^26 and a lazily added top limb wrapped, which made `fe_eq` lie — found
with the CPU reference build in `opencl/hostref/`, which compiles the very same `.cl` as C and
asserts the limb invariants).

| GPU | driver | fe32_wnaf_c | **fe29_wnaf_c** | fe29_w4_c | fe29_w0_c | fe29_wnaf (fused) |
|---|---|---:|---:|---:|---:|---:|
| RX 7900 XTX (gfx1100) | ROCm 7.14 | 5.43 | **10.81** | 9.64 | 8.68 | 7.99 |
| RX 7600 XT (gfx1102) | ROCm 6.4 | 1.91 | **3.58** | | | |
| MI50 (gfx906, 60 CU) | ROCm 6.4 | 1.95 | **4.52** | | | |

2× across three generations (RDNA3, CDNA-era Vega 20), `--iters 8` best everywhere; live on
the pool the XTX reports 10.3–10.5 MH/s with shares accepted. ISA (gfx1100, fe29_wnaf_c):
mult_kernel 177 VGPRs (was 248), 1.16 KB scratch, occupancy 8/16, 8524 `v_mad_u64_u32`
(2× the fe32 count — all 81 products are mads now, but each is a single instruction instead of
mad + carry chain). The 1 T mad64/s measured ceiling argument below now puts the XTX at
≈ 2.5 T mad64-equivalents/s of the 2.31 T peak: the multiply is at the ALU limit; what is left is
the non-multiply work (carries, SHA-512, the scratch-resident wNAF table).

Measured instruction throughput on the RX 7900 XTX (`tools/clmulbench.cpp`, 8 independent
accumulators, the loop also carries 8 cheap updates per 8 multiplies):

| op | Gop/s |
|---|---:|
| 32-bit add/xor | 3681 |
| `mad24` (24×24+32) | 4489 |
| `mul_lo` 32×32→32 | 2834 |
| `mul_hi` 32×32→32 | 2453 |
| `(ulong)a*b + acc` (`v_mad_u64_u32`) | 2310 |

An attempt costs ≈ 3,600 field multiplications × 64 `v_mad_u64_u32` ≈ 230k of them plus
adds/carries/SHA, so 5.4 MH/s is ≈ 1.25 T mad64/s = 55 % of the measured peak before the other
work is counted — the ceiling is ≈ 10 MH/s theoretical, 7–7.5 MH/s realistic. Radix 2^24 limbs
(11 limbs, 121 lo + 121 hi products at the faster `mad24` rate) cost about twice as many
cycle-equivalents as the 8×32 schoolbook and were not pursued. For comparison an RTX 5080 does
~8 T IMAD32/s in the CUDA kernel: the hardware ratio, not the port, explains most of the gap.

## Round 5: launch size (2026-10-20 evening, 8× RTX PRO 6000 Blackwell Server Edition cloud box, CUDA 13.4, driver 580)

The compacted pipeline runs its h2c passes and the mult pass over the whole launch, so a
bigger launch amortises the per-pass overhead and the tail of each pass. One nonce per
thread with many more blocks beats looping inside the thread at the same work per launch.
`--bench 5`, realistic scalar, `fe32_wnaf_lb2_c`, 256 threads (RTX PRO 6000 = 188 SMs, RTX 5080 = 84 SMs):

| blocks (per SM) × nonces/thread | RTX PRO 6000 | RTX 5080 |
|---|---|---|
| 1504 (8) × 4 — the 0.5.0 default | 32.46 MH/s | 17.38 (672 × 4) |
| 3008 (16) × 2 | 34.03 | 17.93 (1344 × 2) |
| 6016 (32) × 1 | 35.38 | 18.47 (2688 × 1) |
| 12032 (64) × 1 | 36.31 | 18.92 (5376 × 1) |
| 24064 (128) × 1 | 36.75 | 19.14 (10752 × 1) |
| 48128 (256) × 1 | 36.94 | 19.24 (21504 × 1) |
| 96256 (512) × 1 | 37.04 | |

Iterating inside the thread is always worse at equal work (12032 × 4: 34.7; 24064 × 2: 35.8).
**New Blackwell default: 128 blocks per SM × 256 threads × 1 nonce** (`Platform.blocksPerSM`,
+13 % on the PRO 6000, +10 % on the 5080); a launch is then ~150 ms on a 5080, ~170 ms on a PRO 6000,
still small against the 10 s slot. `--autotune` now tries 64 and 128 blocks/SM as well.
Ampere measured the same way on the RTX 3070 (46 SMs, sm_86 build, the German test rig): 368 × 4 = 6.78 MH/s,
736 × 2 = 6.92, 1472 × 1 = 6.91, 2944 × 1 = 7.01, 5888 (128/SM) × 1 = 7.07, 11776 × 1 = 7.08 → +4.5 %,
so GA10x, Ada, Turing and the default row use 128 × 1 as well. GA100 (CMP 170HX, `fe32_w4_lb2_c`)
keeps 8 × 2 until measured.

Launch bounds re-checked on the PRO 6000 (dev build, not shipped): `__launch_bounds__(256,1)`
33.1 MH/s, `(256,3)` 36.2, `(128,4)` 33.7, `(128,5)` 32.9, `(128,6)` 32.3, all at the large
launch — `(256,2)` at 36.7 stays the best, so the register budget of 128 is right for GB202 too.
The card draws ~485 W of its 600 W cap at 2430 MHz: not power-limited, pure ALU.

### Batched inversion (2026-10-20, RTX PRO 6000, not shipped)

`PX_INV_BATCH k` (proxichorizo_device.cuh) lets one thread multiply k points, park them and
share one `fe_invert` through Montgomery's trick (3 multiplications per point). nsys puts the
mult kernel at 86 % of GPU time and the inversion at ~10 % of it, so the ceiling was ~8 %.
Measured at 24064 × 256 × 1: k = 1 36.75 MH/s, k = 2 37.30 (+1.5 %), k = 4 36.87, k = 8 36.14 —
the parked points (local memory) eat the saving, as the k = 2 trial on the 3070 had shown.
Left at k = 1; the code stays for GPUs with more registers.

### Round 6 candidates (2026-10-20 late, RTX PRO 6000)

ncu on `mult_kernel` (86 % of GPU time): IMAD pipe 76 % busy, issue slots 57 %, IPC 2.3,
occupancy 33 % (128 registers); stalls per issue: wait 1.46 (fixed-latency dependencies, the
carry chains), no_instruction 1.32 (instruction fetch), math_pipe_throttle 1.26, long_scoreboard
0.01 — the wNAF table in local memory (1 KB/thread, L1 hit 30 %) does not stall anything, so
prefetching the next entry (`PX_PREFETCH`, kept in the source, off) gained nothing: 36.68 vs 36.75.

A 10×26-bit signed-limb representation with 64-bit column products (`PX_FE26`, ref10 layout;
no carry chains, one IMAD.WIDE per product) multiplies 1.4× faster in isolation (microbench:
175 vs 125 G field-mul/s per GPU) but the ref10 carry pass over 64-bit columns costs as much as
the products: 30.6 MH/s at 128 registers (1.4 KB spills), 31.0 at 255, against 36.7 for fe32.
Correct (vectors, field test), kept in the source for GPUs where 64-bit integer adds are cheap.

**Round 6 result: one doubling site + one addition site in the wNAF loop** (`ge_dbl_opt` with a
uniform "T wanted" flag, `ge_add_cached` called once with the sign as an argument, the table
loop not unrolled): the loop body halves, registers 128 → 122 with no spills, and the
`no_instruction` stalls drop 1.32 → 0.60: **36.75 → 40.4 MH/s (+10 %) on the RTX PRO 6000**,
SM throughput 76 → 82 %, issue slots 57 → 63 %. What remains is `wait` 1.5 (the carry chains:
ptxas never overlaps two `.cc` chains, measured: two independent multiplies interleaved in
source run exactly as fast as in sequence) and the IMAD pipe itself. `__noinline__` field
multiply/square (to shrink the loop further) collapses to 11.6 MH/s: call overhead and stack.

### Round 6b: shipped in 0.5.3 (2026-10-20 22:30 UTC)

| GPU | 0.5.2 | merged loop | + batched inversion (2 points) | gain |
|---|---:|---:|---:|---:|
| RTX PRO 6000 | 36.75 | 40.38 | **41.54** | +13 % |
| RTX 5080 | 18.3 | 20.9 | **21.46** | +17 % |
| RTX 3070 | 7.06 | 7.27 | **7.37** | +4.4 % |
| CMP 170HX (GA100) | 5.99 (w4, 8×2) | 6.89 (wNAF, 128×1) | **6.91** | +15 % |

With the smaller loop the wNAF variant beats the 4-bit window on GA100 as well (6.89 vs 6.38
at the same launch; the old 4.7-vs-4.1 note was an instruction-cache effect), so every NVIDIA
row now uses `fe32_wnaf_lb2_c`, 128 blocks/SM × 256 × 1, two points per inversion
(`PX_INV_BATCH 2`, the default for compacted variants).
Larger batches on the PRO 6000 after the loop merge: 3 points 41.66, 4 points 41.96 MH/s (vs 41.54
for 2; no spills up to 4). Left at 2: the smaller cards park less happily (3070: 2 points +1.4 %).

## Round 7: AMD gets the same loop, first-start tuning (2026-10-20/21 night)

The OpenCL kernel got the merged doubling/addition sites and two points per inversion
(`PX_INV_BATCH`, default 2, `--build-opts -DPX_INV_BATCH=1` to compare): RX 7900 XTX
`--bench 8`, clean GPU: 0.5.3 10.72 MH/s, merged loop 10.92, + batching **11.13 (+3.8 %)**.
(An earlier pair of numbers was taken with a forgotten `--selftest` still running on the card;
always `ps` before benchmarking a shared GPU.)

Both miners now tune unmeasured generations on first start (quick autotune: two kernels × two
shapes, 3 s per trial) and cache the choice per card and version in `proxichorizo-autotune.cache`;
measured rows start immediately. Verified on the PRO 6000 (CUDA, row temporarily marked unmeasured:
picks fe32_wnaf_lb2_c 24064×256×1 in 12 s, second start reads the cache) and on the 7900 XTX (OpenCL).

## Round 8: RTX A6000 (GA102, sm_86, 84 SMs, 300 W) — 2026-10-21 23:00 UTC

A 4-card A6000 box (CUDA 13.2) got the full autotune and an Nsight Compute profile of
`mult_kernel` on its spare card.

- Autotune: `fe32_wnaf_lb2_c` wins at 11.6 MH/s (10752×256×1); `fe32_w4_lb2_c` 10.8,
  `fe32_w0_lb2_c` 8.6, `fe51_w4_lb1` 3.1. The platform row's geometry (blocksPerSM 128,
  256 threads, 1 nonce) does 11.8–12.0 MH/s mining, so the Ampere row stays as it is.
  Per SM and GHz the A6000 (12.0 / 84 SMs / 1.69 GHz) matches the 3070 (7.4 / 46 / 1.9):
  the GA10x kernel is not leaving anything card-specific on the table.
- Profile (ncu, `mult_kernel` 200 ms launch): compute-bound, memory idle (10 % DRAM, 0.3 % LSU).
  Pipes: ALU 61 %, FMA-heavy (IMAD) 49 %, i.e. on Ampere the carry/add/logic side leads,
  not the multiplier as on Blackwell. Stalls per issued instruction: wait 1.34 (carry chains),
  math_pipe_throttle 1.24, no_instruction 1.18 (loop body vs. instruction cache), not_selected
  1.04, dispatch 0.96. Occupancy 33 % (128 registers/thread → 16 warps/SM); 43 % of issue
  cycles find no eligible warp. The card runs power-capped: 1.69 GHz under the 300 W limit
  (2.1 GHz max).
- Occupancy experiment (`__launch_bounds__(256,3)` = 85 regs, `(256,4)` = 64 regs, wNAF
  compact, same loop): lb3 11.29 MH/s (−4.7 %, no spills but the compiler's schedule gets
  worse), lb4 10.61 (−10 %, 1.7 KB spills). More warps do not pay for the register cap on
  GA102 either — same answer as lb3/lb5-6 gave on Blackwell. Not shipped.
- Conclusion: nothing cheap left for Ampere; the remaining ideas are the structural ones
  from round 6 (a multiplication without PTX carry chains, fewer field ops per doubling),
  which would have to be measured per generation.

## Round 9: H100 SXM (GH100, sm_90, 132 SMs, 700 W) — 2026-10-22 01:40 UTC

First Hopper measurement, CUDA 13.4, one H100 80GB HBM3 (cloud spot box). Build `-gencode arch=compute_90,code=sm_90`.

| variant | launch | MH/s |
|---|---|---|
| **fe32_w4_lb2_c** | 16896 × 256 × 1 (128 blocks/SM) | **19.8** |
| fe32_w4_lb2_c | 33792 × 256 × 1 | 19.9 |
| fe32_wnaf_lb2_c | 16896 × 256 × 1 | 17.1 |
| fe32_w0_lb2_c | 16896 × 256 × 1 | 13.0 |
| fe32_w4_lb2 / w0_lb2 / fe51_w4_lb1 (non-compact) | best | 3.4 / 2.9 / 2.8 |

Hopper is the first generation where the 4-bit fixed window beats wNAF (+16 %); on Blackwell, GA10x and
(since Round 5) GA100 it is the other way round. The platform table got a measured sm_90 row
(`fe32_w4_lb2_c`, 256 × 1, 128 blocks/SM); the first-start quick-tune would have chosen the same.

Launch shape: 8448–67584 blocks all within 19.5–19.9 (plateau from 128 blocks/SM); 128 threads −20 %,
2 or 4 nonces per thread −5 to −20 %, 512 threads does not launch (128 regs × 512 = full register file).

The card is compute-bound, not power-bound: 1,980 MHz at 400–440 W of the 700 W cap, 44 °C.
ncu on `mult_kernel` (w4): ALU pipe 60 %, IMAD-heavy 49 %, issue-active 56 %, LSU 0.2 %, occupancy 25 %
(128 registers, 2 blocks/SM); stalls per issue: no_instruction 1.62, wait 1.37, math-pipe throttle 1.22,
not_selected 0.94, dispatch 0.80. Same picture as the A6000 in Round 8 with slightly more instruction-fetch stall.

Tried and rejected (experimental build, not in the repo): launch bounds 3 blocks/SM (85 regs) 15.4 MH/s,
4 blocks/SM (64 regs, ~1 KB spills) 15.7; 2-bit window 14.8, 2-bit + 3 blocks/SM 12.6. More warps do not
pay for the spills and the shorter window costs more additions than it saves in table lookups.

Per SM·GHz the H100 (19.9 / (132 × 1.98) = 0.076) sits between GA100 (CMP 170HX 0.07) and the RTX PRO 6000
(0.084): Hopper has 64 INT32 lanes per SM like GA100, so a Hopper-specific kernel would need to move work onto
the FP32/FMA pipes, which the 32-bit-limb field arithmetic does not do. At €1.97/h spot the H100 is the worst
€/MH of the boxes tried (PRO 6000 box ≈ €0.027/MH·h, A6000 ≈ €0.030, H100 ≈ €0.099).

### Round 9b: branch-free carries, cached 4-bit table — 2026-10-22 02:00–03:10 UTC (H100, then every generation)

Pipe microbenchmark on the H100 (`mad.lo.cc`/`madc.hi.cc` chains, independent IMADs, dependent IMAD chain,
`mad.wide.u32`, DFMA, FFMA, and 4 IMAD + 4 DFMA / 4 IMAD + 4 FFMA mixes; 8 and 2 blocks per SM): every integer
form and DFMA run at 64 lane-ops per clock per SM, FFMA at 124, and the mixes at 58–61, i.e. the integer and
FP64/FP32 pipes do not add up through one dispatcher. A dependent IMAD chain already reaches the full 64 with
16 warps per SM. So the miner's ~420 W of the 700 W cap is what the arithmetic draws at full issue, there is no
idle pipe to move work to, and latency hiding is not the limit either.

ncu, wNAF vs 4-bit on the H100: the wNAF kernel executes 7 % fewer instructions (101.0 vs 108.6 G per bench)
but stalls on `no_instruction` 4.03 vs 1.66 per issue. Not the loop's skip branch (a run-length version with a
tight doubling loop measured 16.8): the stalls come from the small data-dependent branches inside every field
operation. `fe_add`/`fe_sub` branched on the carry/borrow out, which is set in about half the lanes, and
`fe_fold512` on the top word, which is non-zero in ~97 % of the lanes, so the warp diverged on nearly every one
of the ~4,000 such branches per point and ran the fix-up anyway. Made unconditional (a `mad.lo.cc` of c×38 plus
the carry chain, same instruction count as the old body):

| card | kernel | before | after |
|---|---|---|---|
| H100 (sm_90) | fe32_w4_lb2_c + cached table | 20.16 | **21.20** (+5.2 %; +7 % vs the shipped 19.8) |
| H100 | fe32_wnaf_lb2_c | 17.10 | 19.31 (+13 %) |
| RTX PRO 6000 (sm_120) | fe32_wnaf_lb2_c | 41.55 | **42.97** (+3.4 %) |
| RTX A6000 (sm_86) | fe32_wnaf_lb2_c | 11.60 | **12.01** (+3.5 %) |

Also in: the 4-bit window's table is now kept in cached form (`ge_add_cached`, 8 instead of 9 multiplications
per addition): H100 19.83 → 20.16 (+1.7 %). The 4-bit window remains Hopper-only; wNAF stays ahead everywhere
else (PRO 6000: w4 cached branch-free 38.1 vs wNAF branch-free 43.0).

Tried on the H100 and dropped: signed odd-digit 5-bit windows (Joye–Tunstall recoding of s−1, 51 additions, no
digit test, 16-entry odd table) 21.17 branch-free = the 4-bit kernel; `__noinline__` field functions 9.9;
squaring through the multiply 15.9; 96-register launch bounds (128×5, 192×3; 262 B spills) 13.4/12.5; one or
three points per inversion 19.3/19.6; the run-length wNAF loop 16.8. Compiled `mult_kernel` sizes are 700–930 KB
(45–59k instructions) per variant; the instruction count, not the size, is what the branch removal changed.