igneum/docs/plans/counter-asic-3-reserve.md

# Counter ASIC 3.0 item 6: the reserve ordered by chip-unfriendliness (PROPOSED)

6 October 2026, branch `ca3-reserve`. The history audit's rank 6 (`docs/analysis/asic-resistance-history.md` section 4.3: families that force a full 32-bit datapath per lane first, `mm8` last, because int8 matrix blocks are licensable IP at every node and Least Authority's ProgPoW suggestion 5 was "watch ML hardware"). Everything here is PROPOSED text for spec 1.13.2: the order and the weights are a decision for the project lead, and nothing here touches the generator, a vector, the manifest or a consensus parameter. The measurements are in `docs/bench-log.md`, entry "6 October 2026, Counter ASIC 3.0 item 6"; this document carries the readings, the ISA facts, the chip structures and the proposed spec text.

## 1. The step cost per family per card

Method (the int8 analysis's, `docs/analysis/int8-matrix-family.md` section 4): a dependent chain of one op per step per lane, 1,048,576 lanes x 4,096 steps, best of 3 dispatches, three runs, bit-exact against a CPU reference on two whole 32-lane warps. The step cost is the chain's best time over the add-xor-rotate chain's (`alu`, the live reference, 5 ops per step). Every candidate chain is the family's one op plus the same four glue ops. Sources `proto-metal/family-probe.swift`, `proto-cuda/family-probe.cu`, job `tools/ca3-reserve/pc2-family-probe.ps1`.

| Family (1.13.2) | Probe row | M5 Max, Metal: step cost (best ms; G steps/s), load 7.64 | RTX 5090, CUDA (PC 2, the card to itself, best of 3 runs): step cost (best ms; G steps/s) | RX 9070 XT, OpenCL (PC 1) |
|---|---|---|---|---|
| reference: add-xor-rotate | alu | 1.00 (4.873; 881) | 1.00 (0.541; 7,941) | OWED (PC 1 not released today) |
| reference: live `rotr` | rotr | 1.13 (5.492; 782) | 1.32 (0.715; 6,005) | OWED |
| reference: live `shfl` (lane XOR mask) | shflx | 0.86 (4.196; 1,024) | 1.49 (0.808; 5,315) | OWED |
| variable left shift | shl | 0.85 (4.119; 1,043) | 1.27 (0.687; 6,253) | OWED |
| variable logical right shift | shr | 0.86 (4.194; 1,024) | 1.28 (0.691; 6,212) | OWED |
| bit-field extract, immediate offset and width | bfe (`extract_bits`) 0.77; bfec (C form) 0.75 | 0.77 (3.731; 1,151) | 1.54 (0.835; 5,142), `bfe.u32` and the C form equal to the nanosecond: a compiler sequence, no single instruction | OWED |
| andn | andn | 0.75 (3.676; 1,168) | 1.26 (0.680; 6,313) | OWED |
| byte permute, immediate selector | perm | 1.13 (5.500; 781), EMULATED (shifts and masks) | 1.30 (0.706; 6,084), `prmt` | OWED (`v_perm_b32`) |
| popcount folded by add | popc | 0.87 (4.223; 1,017) | 1.50 (0.813; 5,283) | OWED |
| clz folded by add | clz | 1.01 (4.914; 874) | 1.63 (0.884; 4,856) | OWED |
| three-register select | sel | 0.76 (3.718; 1,155) | 1.32 (0.713; 6,024) | OWED |
| second shuffle form, lane + delta mod 32 | shfla | 1.91 (9.287; 462) | 1.53 (0.829; 5,179), `shfl.sync.idx`: the same cost as the live xor shuffle | OWED (`ds_bpermute_b32`) |
| comparison: dp4a | dot4u 1.60 (emulated), dot4s 4.73 (emulated) | as the 5 October rows (1.6x, 4.7x) | 1.16 (0.625; 6,870), `__dp4a` (1.17x on 5 October) | 1.06x on 5 October (`v_dot4_i32_iu8`) |
| comparison: mm8 as a chain | mm8 | OWED (Metal 4 `matmul2d`; Swift 5.8 toolchain has no tensor API) | 2.43 (1.313; 3,272), one `mma.sync.m8n8k16.u8` per warp per step, bit-exact (the fragment layout confirmed) | OWED (WMMA) |

The three Mac runs agree within 4% on every row and the three 5090 runs within 3% (12% on `shl`, which caught the clock ramp); every row on both cards is bit-exact. The AMD column is the owed row: PC 1 is the project lead's desk today.

What the Mac rows say. Six of the seven candidates cost at most the live `rotr` on Apple. The two Apple pays for: `perm` at 1.13x (no byte-permute function in MSL, the `uchar4` swizzle compiles to shifts and masks) and `shfla` at 1.91x (`simd_shuffle` by a computed lane index against `simd_shuffle_xor`, 2.2x the live shuffle). Both are inside the 8x per-op bound of 1.13.2 by a wide margin, and at `W_new` = 4 points of the 64-instruction program a 1.91x op is under 1% of the program's ALU time on a hash that spends its time on 128 dependent DRAM reads (approximate: argued from the step cost and the weight, not measured; the 5% rule is checked on the vendor's card with the family live, as 1.13.2 says).


What the 5090 rows say. Every candidate costs more than the two-register reference chain on NVIDIA (1.26x to 1.63x), so the live `rotr` (1.32x) is the fair bar: against it `andn`, the shifts, `perm` and `sel` read 0.95x to 1.00x, `popc` 1.14x, `shfla` 1.16x (the indexed shuffle costs NVIDIA the same as the xor one), `bfe` 1.17x and `clz` 1.23x. Nothing is emulated beyond a compiler sequence (`bfe` and `clz`), nothing is near the 8x bound, and `mm8` is the dearest row on NVIDIA too (2.43x: one tensor mma per warp per dependent step), so the honest-card side agrees with the chip side on ordering it last. The two cards disagree on `shfla`: 1.91x on Apple, free on NVIDIA; and on `perm`: emulated on Apple, native on NVIDIA. The order of section 5 stands on both columns; the AMD column (owed) is the one that can still move R3.

## 2. Native or emulated, per vendor, cited

Sources read 6 October 2026: PTX ISA 9.4 (https://docs.nvidia.com/cuda/parallel-thread-execution/index.html, section numbers below); the LLVM AMDGPU backend's instruction tables (https://github.com/llvm/llvm-project/tree/main/llvm/lib/Target/AMDGPU: `VOP1Instructions.td`, `VOP2Instructions.td`, `VOP3Instructions.td`, `DSInstructions.td`, `VOP3PInstructions.td`, main, line numbers below; AMD's CDN refused the RDNA 3 and RDNA 4 ISA PDFs again today, as on 5 October, so the mnemonics are the compiler's, not quoted from the ISA guide); Metal Shading Language Specification version 4.1 (https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf, section 6.4 Integer Functions and 6.10.2 SIMD-Group Functions). Apple's GPU instruction set is not published, so "native" on Apple means "an MSL function exists"; the step cost is the only measure of what the compiler emits.

| Family | NVIDIA (PTX) | AMD (RDNA, by the LLVM mnemonic) | Apple (MSL) | Emulated where |
|---|---|---|---|---|
| variable shifts | `shl.b32`, `shr.b32` (9.7.9.8, 9.7.9.9; "supported on all target architectures") | `v_lshlrev_b32`, `v_lshrrev_b32` (VOP2Instructions.td 899, 901) | `<<`, `>>` operators (6.4) | nowhere |
| bit-field extract | `bfe.u32` (9.7.1.20, sm_20+); on the 5090 the PTX form and the C form cost the same to the nanosecond, so it is a compiler sequence there, not one instruction | `v_bfe_u32` (VOP3Instructions.td 276) | `extract_bits` (6.4) | nowhere by the ISA; Apple's lowering unknown (0.77x measured, so no penalty) |
| andn | `and.b32` with `not.b32`, folded into `lop3.b32` on sm_50+ (9.7.9.1, 9.7.9.4, 9.7.9.6) | `v_and_b32` with `v_not_b32`, or one `v_bfi_b32` (VOP2 902, VOP1 378, VOP3 278) | `&`, `~` (6.4) | nowhere |
| byte permute | `prmt.b32` (9.7.10.7, sm_20+) | `v_perm_b32` (VOP3Instructions.td 426) | NO function; `uchar4` swizzle, 1.13x measured | Apple |
| popcount, clz | `popc.b32`, `clz.b32` (9.7.1.15, 9.7.1.16, sm_20+) | `v_bcnt_u32_b32` (VOP2 980: popcount plus an add in one instruction, the folded form exactly), `v_ffbh_u32`, renamed `v_clz_i32_u32` on gfx11+ (VOP1 380, 1296) | `popcount`, `clz` (6.4; `clz(0)` = 32 on all three) | nowhere |
| three-register select | `selp.b32` (9.7.7.3) | `v_cndmask_b32` (VOP2Instructions.td 878) | `select` (6.4) | nowhere |
| second shuffle form (lane + delta) | `shfl.sync.idx.b32` (9.7.10.6, sm_30+) | `ds_bpermute_b32` (DSInstructions.td 839: a backward permute through the LDS crossbar, not a VALU op; the live xor form has the same route or DPP) | `simd_shuffle` (6.10.2) | nowhere by the ISA; Apple 1.91x measured, AMD's LDS route owed |
| dp4a (comparison) | `dp4a` (9.7.1.24, sm_61+) | `v_dot4_i32_iu8` (VOP3PInstructions.td 859) | none; emulated 1.6x / 4.7x | Apple |
| mm8 (comparison) | `mma.sync.m8n8k16` `.u8` (9.7.16.5.3, sm_75+) | `v_wmma_i32_16x16x16_iu8` (VOP3PInstructions.td 1725, 2307) | Metal 4 `matmul2d` (a library path, cost owed) | nowhere by the ISA |

## 3. What a chip pays per family

A 12-op chip (the ledger's M1: a sequencer over the live families and 8 registers) carries per lane: a 32-bit adder (also the subtractor), a 32x32 multiplier with the high half (`mul`, `mulhi`, `mad`), xor, or, a 32-bit rotator (`rotl` by immediate, `rotr` by a register), the xor-shuffle butterfly per warp (5 stages, masks 1, 2, 4, 8, 16) and the load path. Every structure below is named with its approximate area relative to that lane's 32-bit adder (approximate, from memory of standard-cell estimates: a ripple adder is about 32 full-adder cells; a 2:1 mux is about a third of a full adder; these are order-of-magnitude figures to rank the families, not a layout).

| Family | Structure the chip must add | Approximate area, in 32-bit adders per lane | Note |
|---|---|---|---|
| variable shifts | a 32-bit barrel shifter: 5 stages x 32 2:1 muxes, plus the zero fill | 1 to 2; but the live `rotr` already needs a 32-bit barrel rotator, and a shifter is the rotator with a fill mask: about 0.2 more | the smallest addition to a chip that already runs class v3 |
| bit-field extract | the shifter plus a 32-bit mask generator (a decoder over the width) | 0.3 beyond the shifter | small |
| andn | 32 inverters on one input of the AND | under 0.1 | nothing a chip lacks |
| byte permute | a 4x4 byte crossbar: 4 output bytes x 4:1 mux x 8 bits = 128 mux bits, plus the selector decode | 1 to 2 | a structure no live family needs; a chip that rotates by multiples of 8 does half of it, the other half is new |
| popcount, clz | a 32-bit popcount tree (31 small adders, 5 levels) and a 32-bit priority encoder (leading-zero count) | 1.5 to 3 together | the tree is a new structure; neither is in the live set |
| three-register select | 32 2:1 muxes and a bit select | about 0.3 | nothing a chip lacks; every sequencer has operand muxes |
| second shuffle form | a full 32-lane x 32-bit crossbar per warp (32 outputs x 32:1 mux x 32 bits = 32,768 mux bits) in place of the 5-stage butterfly (5 x 32 x 32 = 5,120 mux bits) | about 6x the butterfly per warp, so about 3 to 6 adders per lane (approximate) | the largest 32-bit structure of the seven; on a GPU it is the existing shuffle network, so free for the honest card |
| dp4a (comparison) | four 8x8 multipliers and a 4-input adder tree | 4 to 6 | the chip's multiplier can be split into four 8x8 blocks; cheap for a chip |
| mm8 | a 8x8x16 u8 MAC tile per warp (1,024 MACs; 32 per lane) | about 100 per lane (approximate) | licensable IP at every node (rank 6); the structure the history says to order last |

## 4. The bounds of 1.13.2 per family

| Family | 8x per-op emulation bound | 5% hash-rate bound at `W_new` = 4 | Reading |
|---|---|---|---|
| variable shifts | holds everywhere: no emulation (Apple 0.85 / 0.86x, native PTX and RDNA) | holds, argued: a 0.86x op at 4 points of 64 moves the ALU time by under 1%, and the hash is read-bound (latency-bound share 0.95 to 1.06 on the three cards, bench-log "Counter ASIC 2.0, the numbers") | in |
| bit-field extract | holds: Apple 0.77x; native on PTX and RDNA | holds, argued | in |
| andn | holds: 0.75x; native everywhere | holds, argued | in |
| byte permute | holds: Apple emulated at 1.13x (inside 8x by 7x); NVIDIA 1.30x native `prmt` (0.98x the live `rotr`); RDNA `v_perm_b32` | holds, argued: a 1.13x op at 4 points is under 0.5% of the ALU time | in; Apple is the vendor that can only emulate, inside both bounds |
| popcount, clz | holds: 0.87x / 1.01x; native everywhere | holds, argued | in |
| three-register select | holds: 0.76x; native everywhere | holds, argued | in |
| second shuffle form | holds: Apple 1.91x (inside 8x by 4x); NVIDIA 1.53x, equal to the live xor shuffle; AMD's `ds_bpermute_b32` is an LDS op whose step cost is owed | holds, argued for Apple (under 1% at 4 points); AMD argued from the live `shfl` sharing the route, measured when the PC 1 row lands | in, pending the AMD row |
| mm8 | Metal 4 `matmul2d` is a path, cost owed; the per-lane dot4 is 1.6x emulated | holds, argued (`int8-matrix-family.md` section 4) | in, R8 |

Every bound above is argued from the step cost and the weight, not measured on a live program: no reserve family is in a live program today. The measurement is the family-live run of 1.13.2's own text, owed per family at its unlock rehearsal.

## 5. The proposed order

Ranked by the structure a chip must add beyond the class v3 datapath (section 3), then by what the honest vendors pay (section 1): the largest new 32-bit structure first, the trivial ones last, `mm8` last of all.

| Reserve slot | Family | Why here | `W_new` | Unlock by the rule of 1.13.2 (family n at era n; era = 180 days) |
|---|---|---|---|---|
| R1 | byte permute (`perm`) | a byte crossbar per lane that no live family needs; native on NVIDIA and AMD; Apple pays 1.13x, inside both bounds | 4 | era 1 (day 180) |
| R2 | popcount and clz folded by add (`popc`, `clz`) | a popcount tree and a priority encoder, both new; native on all three; costs no vendor anything | 4 | era 2 |
| R3 | second shuffle form (`shfla`) | the 32-lane crossbar is the largest 32-bit structure of the seven; native on all three; Apple pays 1.91x; the AMD LDS cost is the owed number that could move this row to R1 (cheap on AMD) or to R5 (dear) | 4 | era 3 |
| R4 | bit-field extract (`bfe`) | a mask generator on the shifter; native everywhere; 0.77x on Apple | 4 | era 4 |
| R5 | variable shifts (`shl`, `shr` by `src AND 31`, the direction an immediate bit) | the smallest addition to a chip that already rotates; native everywhere | 4 | era 5 |
| R6 | three-register select (`sel`) | operand muxes a sequencer has anyway; its value is the data-dependent path, not the silicon | 4 | era 6 |
| R7 | andn | an inverter; last of the datapath families | 4 | era 7 |
| R8 | `mm8` (integer matrix) | licensable IP at every node; the vendor-emulation rule written for it | 4 (as decided 5 October) | era 8 (day 1,440), or earlier by the 90% signal |

The decision this moves for the project lead: `mm8` was reserved on 5 October as R1 with an unlock at era 4. Under the rule "family n at era n" an eighth slot unlocks at era 8 (four years). Either the rule stays and `mm8` waits, or the entry keeps its era-4 unlock as a named exception (the text below keeps the era-4 date as an exception and says so).

## 6. PROPOSED spec text for 1.13.2 (replaces the paragraph from "The order and `W_new` are Open" to the end of the `mm8` entry)

> The reserve, in order. Reserve family `n` becomes live at the start of era `n` (1.13.1), or earlier by the 90% signalling path of section 5.7; never by a release. Every entry takes `W_new` = 4 points proportionally from the live non-load families (the load weight and count are untouched). Every entry's edge vectors are hand-built units run on every vendor of 1.15 before genesis. The ordering rule: the family whose silicon a class v3 chip lacks most comes first, the integer matrix family last (`docs/plans/counter-asic-3-reserve.md`).
>
> R1, `perm` (byte permute). Semantics: `dst = permute(dst, sel4)`, result byte `i` = byte `(sel4 >> 2i) AND 3` of `dst`, `sel4` an 8-bit immediate drawn per instruction. Native: PTX `prmt.b32`, RDNA `v_perm_b32`; Apple by `uchar4` swizzle (emulation, 1.13x per op measured 6 October 2026). Edge vectors: `dst` = 0x00FF00FF with every `sel4` of the form (3, 2, 1, 0), (0, 1, 2, 3), (0, 0, 0, 0), (3, 3, 3, 3); `dst` = 0x80808080 (the sign bits travel as bytes, no sign extension); `dst` = 0x01020304 with `sel4` = (1, 3, 0, 2) (the probe's selector: result 0x03010402); alternating 0x00 and 0xFF by lane.
>
> R2, `popc` (population count and count-leading-zeros). Semantics: `dst = dst + popcount(src)` when the immediate bit is 0, `dst = dst + clz(src)` when it is 1, `clz(0)` = 32, add modulo 2^32. Native: PTX `popc.b32` and `clz.b32`, RDNA `v_bcnt_u32_b32` and `v_clz_i32_u32`, MSL `popcount` and `clz`. Edge vectors: `src` = 0 (popcount 0, clz 32); `src` = 0xFFFFFFFF (32, 0); `src` = 1 (1, 31); `src` = 0x80000000 (1, 0); `dst` = 0xFFFFFFFF with `src` = 0xFFFFFFFF (the wrap to 31); both bits on the same operands.
>
> R3, `shfla` (shuffle by lane plus delta). Semantics: `dst = dst XOR src_of_lane((lane + delta) mod 32)`, `delta` in 1..31 drawn per instruction, within the lane's own 32-lane warp. Native: PTX `shfl.sync.idx.b32`, RDNA `ds_bpermute_b32`, MSL `simd_shuffle`. Edge vectors: `delta` = 1 and 31 on a warp whose `src` is the lane index (the wrap at lane 31 and lane 0); `delta` = 16 against the live `shfl` mask 16 (equal results); `src` all equal (dst unchanged except by the xor with itself: 0 on every lane); one lane's `src` = 0xFFFFFFFF, the rest 0, at `delta` = 7.
>
> R4, `bfe` (bit-field extract). Semantics: `dst = (src >> off) AND (2^w - 1)`, `off` in 0..31 and `w` in 1..32 drawn per instruction, `off + w` clamped to 32 (bits past 31 read as 0). Native: PTX `bfe.u32`, RDNA `v_bfe_u32`, MSL `extract_bits`. Edge vectors: `off` = 0, `w` = 32 (identity); `off` = 31, `w` = 1; `off` = 7, `w` = 13 (the probe); `off` = 24, `w` = 16 (the clamp: 8 bits come back); `src` = 0xFFFFFFFF at every pair above; `src` = 0.
>
> R5, `shl` (variable shifts). Semantics: `dst = dst << (src AND 31)` when the immediate bit is 0, `dst = dst >> (src AND 31)` (logical) when it is 1. Native everywhere (PTX `shl.b32`, `shr.b32`; RDNA `v_lshlrev_b32`, `v_lshrrev_b32`; MSL operators). Edge vectors: `src AND 31` = 0 (identity) and 31; `dst` = 0xFFFFFFFF at both; `dst` = 0x80000000 right by 31 (1) and left by 1 (0); `src` = 0xFFFFFFE0 (the mask reads 0, not 32); both bits on the same operands.
>
> R6, `sel` (three-register select). Semantics: `dst = (bit `bit` of src2) ? src : dst`, `bit` in 0..31 drawn per instruction. Native everywhere (PTX `selp.b32`, RDNA `v_cndmask_b32`, MSL `select`). Edge vectors: `src2` = 0 and 0xFFFFFFFF at `bit` = 0 and 31; `src` = `dst`; `src2` = `dst` (the selected bit reads the register being written: the value before the instruction); alternating by lane.
>
> R7, `andn`. Semantics: `dst = dst AND NOT src`. Native everywhere. Edge vectors: `src` = 0 (identity), 0xFFFFFFFF (0), `src` = `dst` (0); `dst` = 0xAAAAAAAA with `src` = 0x55555555 (unchanged).
>
> R8, `mm8` (integer matrix): the entry decided 5 October 2026, text unchanged (semantics, `W_new` = 4, the six edge vectors, native paths, the vendor-that-can-only-emulate rule), with one change: it is the eighth reserve family. Its unlock stays the start of era 4 (DAA 62,208,000) as a named exception to "family n at era n", or moves to era 8 if the project lead keeps the rule; one of the two is decided before the public testnet genesis.

The `andn` and `sel` entries carry a note for the generator: both are injecting in the sense of 1.4.1 only when the acceptance rule counts them so (`sel` writes `src` into `dst` on half the lanes; `andn` is not bijective in `dst`), so neither satisfies rule (b) of 1.4.6 alone; the weak-program census is re-run with each family at its unlock rehearsal.

## 7. Consequences per tier

| Tier | Today (the reserve is not live) | At each unlock, from the rows above |
|---|---|---|
| Apple user (M-series; the app's Metal worker) | nothing changes | R1 `perm` costs about 1.13x on 4 points of 64 (under 0.5% of ALU time, argued); R3 `shfla` 1.91x on 4 points (under 1%); every other family costs less than the live `rotr`. The measurement with the family live is owed at each unlock rehearsal; if a live run shows over 5% on Apple the family does not unlock (1.13.2) |
| NVIDIA user (8 to 32 GB) | nothing changes | every family native or a two-instruction compiler sequence; 0.95x to 1.23x the live `rotr` step on the 5090 (section 1); no family can cost over 1% of ALU time at 4 points (argued) |
| AMD user (RX 9070 XT) | nothing changes | every family native by the LLVM tables; the `shfla` cost through `ds_bpermute_b32` (an LDS op) is the owed number, and the one that could move R3 |
| Intel iGPU user | nothing changes | not measured anywhere; OpenCL C exposes every function (`popcount`, `clz`, `sub_group_shuffle`); owed with the first Intel measurement |
| A rig, a pool user | nothing changes | the same per-card figures; a pool verifies with the CPU reference, which gains one op per family and stays under the 10 ms gate (the verifier runs 2.08 ms per warp at x8 against a 10 ms gate) |
| A chip built against class v3 | nothing changes | R1 to R3 each add a structure it lacks (byte crossbar, popcount tree, 32-lane crossbar) within 18 months of genesis; R4 to R7 add little silicon but take weight from the families it was built for; `mm8` at R8 adds the licensable block last |

## 8. What is owed

| Item | State |
|---|---|
| RTX 5090 step costs (every row) | MEASURED (job `run-ca3-family-pc2-20261006`, 08:42Z to 08:43Z, the card to itself, every row bit-exact) |
| RX 9070 XT step costs (every row) | OWED: PC 1 is the project lead's desk and not used today; the OpenCL twin of the probe is the next job on that card |
| `mm8` as a chain on Apple (Metal 4 `matmul2d`) | OWED (toolchain) |
| AMD `ds_bpermute_b32` cost for `shfla` | OWED (the PC 1 row); decides whether R3 moves |
| The 5% hash-rate rule per family with the family live | argued, not measured, for every family (no reserve family is in a live program); measured at each unlock rehearsal |
| RDNA 3 and RDNA 4 ISA guides (the instruction text) | AMD's CDN refused the PDFs again; the mnemonics are from the LLVM tables |
| the project lead's decisions | the order R1 to R8; `W_new` = 4 per family; `mm8` at era 4 by exception or at era 8 by the rule |