21 KiB
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
nbecomes live at the start of eran(1.13.1), or earlier by the 90% signalling path of section 5.7; never by a release. Every entry takesW_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 bytei= byte(sel4 >> 2i) AND 3ofdst,sel4an 8-bit immediate drawn per instruction. Native: PTXprmt.b32, RDNAv_perm_b32; Apple byuchar4swizzle (emulation, 1.13x per op measured 6 October 2026). Edge vectors:dst= 0x00FF00FF with everysel4of 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 withsel4= (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: PTXpopc.b32andclz.b32, RDNAv_bcnt_u32_b32andv_clz_i32_u32, MSLpopcountandclz. Edge vectors:src= 0 (popcount 0, clz 32);src= 0xFFFFFFFF (32, 0);src= 1 (1, 31);src= 0x80000000 (1, 0);dst= 0xFFFFFFFF withsrc= 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),deltain 1..31 drawn per instruction, within the lane's own 32-lane warp. Native: PTXshfl.sync.idx.b32, RDNAds_bpermute_b32, MSLsimd_shuffle. Edge vectors:delta= 1 and 31 on a warp whosesrcis the lane index (the wrap at lane 31 and lane 0);delta= 16 against the liveshflmask 16 (equal results);srcall equal (dst unchanged except by the xor with itself: 0 on every lane); one lane'ssrc= 0xFFFFFFFF, the rest 0, atdelta= 7.R4,
bfe(bit-field extract). Semantics:dst = (src >> off) AND (2^w - 1),offin 0..31 andwin 1..32 drawn per instruction,off + wclamped to 32 (bits past 31 read as 0). Native: PTXbfe.u32, RDNAv_bfe_u32, MSLextract_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 (PTXshl.b32,shr.b32; RDNAv_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 = (bitbitof src2) ? src : dst,bitin 0..31 drawn per instruction. Native everywhere (PTXselp.b32, RDNAv_cndmask_b32, MSLselect). Edge vectors:src2= 0 and 0xFFFFFFFF atbit= 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 withsrc= 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 |