From 0f3b23d7a7082303eeecf04901fbe8a3ee776dcb Mon Sep 17 00:00:00 2001 From: igneum-labs <337424239+igneum-labs@users.noreply.github.com> Date: Tue, 6 Oct 2026 07:40:41 +0000 Subject: [PATCH] Counter ASIC 3.0 items 6 and 7: proposed reserve order and spec text (docs/plans/counter-asic-3-reserve.md), the vendor-share section on the benchmark page, the observer README row --- docs/benchmarks/repro.md | 46 ++++++++- docs/plans/counter-asic-3-reserve.md | 137 +++++++++++++++++++++++++++ tools/observer/README.md | 1 + 3 files changed, 183 insertions(+), 1 deletion(-) create mode 100644 docs/plans/counter-asic-3-reserve.md diff --git a/docs/benchmarks/repro.md b/docs/benchmarks/repro.md index 4c4cc72e0..2527fd7aa 100644 --- a/docs/benchmarks/repro.md +++ b/docs/benchmarks/repro.md @@ -206,7 +206,51 @@ key of 8.1 is the operator's own step, on the file. | The prover on PC 2 | the first script read the prover's state from `api/state` too, saw nothing, and left the live prover off from 00:49 to the restore job (run-prover-on-pc2-20261006); the script now reads settings.json and restores unconditionally | done tonight; the rule in the job | | The reward | the amount and payer are `docs/plans/funding.md`, not funded; the terms are quoted unchanged | the project lead's decision | -## 8. Files +## 8. The vendor-share metric (Counter ASIC 3.0 item 7) + +6 October 2026, branch `ca3-reserve`. The history audit's rank 7 (`docs/analysis/asic-resistance-history.md` section 4.3): a one-vendor fleet is a softer form of chip capture (lesson 8: Equihash's NVIDIA tilt, Ethash's balance), so the share of hash rate by vendor is published beside the benchmark and the AMD gap gets a 3.0 target. The observer computes it every 60 s into `live_state.vendor_share` (`tools/observer/vendor-share.mjs`; the one hook in `observer.mjs`; `node --test tools/observer/vendor-share.test.mjs` fires on a known-finished and a known-failed case; `node tools/observer/vendor-share.mjs --dry` prints the live reading and writes nothing). + +Definition. Hash rate by vendor (nvidia, amd, apple, intel, unknown), two readings, each named for what it is: + +| Reading | What it is | What it covers | What it misses | +|---|---|---|---| +| fleet-reported | the sum of `now=X MH/s wall` from the newest STATUS line of every miner worker that uploaded to the log intake in the last 10 minutes, by the vendor in the worker's label (`miner-nvidia-`, `miner-amd-`, `miner-mac-`; `miner-other-` by the app's `cards:` line) | only machines that upload to the intake: the project's own fleet and any app install carrying the intake key | the rest of the network entirely; a worker that stopped uploading | +| chain-attributed | each miner id's (`vote_key_hash`) share of the BLUE blocks of the last 10 minutes times the node's network hash-rate estimate, attributed to the vendor of the fleet worker that logged that vote key (`identity N '...' vote_key_hash=` in its first uploads), else "unknown" | the whole network's blocks | the vendor of any miner that is not a reporting worker (that is the "unknown" row, the honest hole); pending and red blocks are left out | + +Today's devnet, read-only from the live tables (dry mode, 07:36 UTC, 6 October 2026; nothing written, the live observer not restarted; the hook goes live with the merge): + +| Vendor | fleet-reported MH/s (workers) | chain-attributed share of blue blocks (10 min) | chain-attributed MH/s at the 135.8 MH/s estimate | Note | +|---|---|---|---|---| +| NVIDIA | 120.61 (1: PC 2's RTX 5090, with the prover on the card) | 0.986 (552 of 560) | 133.86 | PC 1's RTX 5090 was off the network at the reading (the project lead's desk) | +| AMD | 0 (0) | 0 | 0 | the RX 9070 XT is on PC 1, off at the reading | +| Apple | 0 (0) | 0 | 0 | the M5 Max is paused for measurements | +| Intel | 1.69 (1: the Windows laptop's UHD Graphics) | 0.014 (8 of 560) | 1.94 | the first outside machine, which carries the intake key | +| unknown | 0 | 0 | 0 | coverage 1.00: every blue block of the window came from a reporting worker | + +At 07:21 UTC, before PC 1 went off, the same tables read 243 MH/s and 18 miner ids; the window then was NVIDIA about 0.9 (PC 1's and PC 2's 5090s) and AMD about 0.01 (the 9070 XT's first 3 blocks after its restart). The devnet today is the project's own machines, so the two readings agree; on a public testnet the chain-attributed "unknown" row is the number that matters, and the fleet-reported row is only the project's own share. + +Rate per watt and per pound, from the bench rows that exist (every price approximate, UK list prices from memory, 6 October 2026): + +| Card | MH/s (class v3) | Power | MH/s per W | Price (approximate) | MH/s per £ | Source | +|---|---|---|---|---|---|---| +| RTX 5090 | 136.1 | about 326 W (328.6 W peak with the prover, 5 October; the app's stability line today: PC 1 p95 307 W mining alone, PC 2 p95 338 W with the prover) | 0.42 (at 326 W); 0.44 at 307 W | about £2,000 | 0.068 | bench-log "Counter ASIC 2.0, the numbers"; miner_logs stability lines | +| RX 9070 XT | 18.6 to 19.2 | OWED: no measured figure (the app logs no stability line for the Radeon; PC 1 not available today); the board's 304 W TBP is a list figure, approximate | about 0.06 at the list TBP (approximate) | about £600 | 0.031 | bench-log; price approximate | +| Apple M5 Max | 27.9 | OWED: no powermetrics figure in the bench-log | not computed | the laptop, about £3,500 and up; not a GPU purchase | 0.008 (approximate, against the laptop price) | bench-log; price approximate | +| Intel UHD (laptop iGPU) | 1.85 (the US laptop's STATUS line) | not measured | | | | miner_logs | + +The 2.0 record put AMD at 2.2x worse per pound and 4.9x worse per watt than the 5090 (approximate); the rows above read 2.2x per pound and about 7x per watt at the list TBP, so the per-watt figure is the one the owed measurement must settle. + +Consequence per tier: + +| Tier | What a 90% NVIDIA share means | What the project does | +|---|---|---| +| An AMD owner (RX 9070 XT, 16 GB) | at equal network difficulty the card earns 19 / 136 = 0.14 of a 5090's income at about 0.3 of its price, so about half the income per pound; the share itself does not change that, but a 90% NVIDIA network sets the difficulty by NVIDIA's rate, and a chip built against one vendor's memory system would hit AMD owners first (lesson 8) | the line-width question stays open per the plan (`docs/plans/counter-asic-3.md`: w16 closes nothing, w64 makes the 5090 bandwidth-bound); the 3.0 target for the AMD gap is the per-read gap (2.4 G against 17.5 G dependent reads per second) measured per driver release, with the first target a 2x closing by the public testnet or a stated reason it cannot close; the vendor share is published so the tilt is visible | +| An Apple owner (M-series) | 27.9 MH/s is 0.2 of a 5090 on a machine bought for other reasons; the share means the Mac is a minority that no family in the reserve may cost more than 5% (1.13.2) | the reserve order of item 6 puts the two families Apple pays for (perm, shfla) at R1 and R3 with the costs measured; the 5% rule is checked live before each unlock | +| An NVIDIA owner | the majority sets the difficulty; a chip that beats NVIDIA is the only chip that matters | the chip model and the bounty (bench-log "Counter ASIC 2.0, the numbers") | +| A rig or a pool | a pool's share by vendor is what the chain-attributed reading cannot see (one vote key per pool) | pools publish their own vendor mix or appear as "unknown"; the detector (item 4) reads the per-program spread per card model | +| The project | a fleet-reported reading above the chain-attributed one means a reporting worker is not finding blocks (a stuck worker, the 4 October class); below it means miners outside the intake | the daily report carries both readings from the merge on | + +## 9. Files | File | What | |---|---| diff --git a/docs/plans/counter-asic-3-reserve.md b/docs/plans/counter-asic-3-reserve.md new file mode 100644 index 000000000..3e8035342 --- /dev/null +++ b/docs/plans/counter-asic-3-reserve.md @@ -0,0 +1,137 @@ +# 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) | RX 9070 XT, OpenCL (PC 1) | +|---|---|---|---|---| +| reference: add-xor-rotate | alu | 1.00 (4.873; 881) | PENDING the PC 2 job | OWED (PC 1 not released today) | +| reference: live `rotr` | rotr | 1.13 (5.492; 782) | PENDING | OWED | +| reference: live `shfl` (lane XOR mask) | shflx | 0.86 (4.196; 1,024) | PENDING | OWED | +| variable left shift | shl | 0.85 (4.119; 1,043) | PENDING | OWED | +| variable logical right shift | shr | 0.86 (4.194; 1,024) | PENDING | OWED | +| bit-field extract, immediate offset and width | bfe (`extract_bits`) 0.77; bfec (C form) 0.75 | 0.77 (3.731; 1,151) | PENDING (`bfe.u32` and the C form) | OWED | +| andn | andn | 0.75 (3.676; 1,168) | PENDING | OWED | +| byte permute, immediate selector | perm | 1.13 (5.500; 781), EMULATED (shifts and masks) | PENDING (`prmt`) | OWED (`v_perm_b32`) | +| popcount folded by add | popc | 0.87 (4.223; 1,017) | PENDING | OWED | +| clz folded by add | clz | 1.01 (4.914; 874) | PENDING | OWED | +| three-register select | sel | 0.76 (3.718; 1,155) | PENDING | OWED | +| second shuffle form, lane + delta mod 32 | shfla | 1.91 (9.287; 462) | PENDING (`shfl.sync.idx`) | OWED (`ds_bpermute_b32`) | +| comparison: dp4a | dot4u 1.60 (emulated), dot4s 4.73 (emulated) | as the 5 October rows (1.6x, 4.7x) | PENDING (`__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) | PENDING (`mma.sync.m8n8k16.u8`, inline PTX) | OWED (WMMA) | + +The three Mac runs agree within 4% on every row; the 5090 column is filled from the job's closing report the moment it is read (bench-log entry). 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). + +## 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+); whether SASS keeps it as one instruction on Blackwell is what the `bfe` against `bfec` rows measure | `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); native PTX `prmt`, 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); native PTX; 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; the 5090 rows (pending the PC 2 job) give the per-op cost; `bfe` is the one to watch (SASS has no single bfe on recent architectures by the measured `bfe` against `bfec` rows) | +| 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) | PENDING: the PC 2 job, after `/tmp/igneum-devnet/pc2-ca3.clear` and under the mkdir lock; one job carries every family three times | +| 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 | diff --git a/tools/observer/README.md b/tools/observer/README.md index 466fe10ed..152112e9e 100644 --- a/tools/observer/README.md +++ b/tools/observer/README.md @@ -49,6 +49,7 @@ Created on start if missing. | `live_checkpoints` | one per checkpoint index, kept 7 days | `index`, `hash`, `blue_score`, `daa_score`, `state` (proposed, certified, locked), `signed_weight`, `active_weight`, `total_weight`, `fraction_active`, `fraction_total`, `votes_seen`, `voters`, `aggregators` (key hashes whose sortition proof made them aggregators), `locked_at`, `first_seen_at`, `updated_at`. | | `live_proofs` | one per planned shard of a chain block, kept `LIVE_RETAIN_HOURS` | `block_hash`, `shard`, `shards` (in the plan), `block_number`, `block_daa`, `block_ts`, `pgas`, `state` (planned, proving, verified, paid), `prover` (id8), `verified`, `carried_by`, `carrier_number`, `carrier_daa`, `lag_daa`, `payout_wei`, `received_at`, `updated_at`. Primary key `(block_hash, shard)`, index on `received_at`. | | `live_state.proving` | jsonb, updated every 2 s | `supported` (false with `reason` when the node has no proving RPCs or the endpoint is unreachable), `active`, `activation_daa`, `tip_daa`, `verifier`, `pool` (entries, pending, verified, failed), `paid_shards_total`, `shard_budget_pgas`, `blocks_10m`, `blocks_fully_proven_10m`, `shards_proven_10m`, `shards_paid_10m`, `median_proof_lag_s` (median `lag_daa` of the last 10 minutes; the devnet targets one DAA step per second), `provers_10m`, `open_blocks`, `pending_plans`, `evm_rpc`. | +| `live_state.vendor_share` | jsonb, updated every 60 s (`tools/observer/vendor-share.mjs`, Counter ASIC 3.0 item 7) | `at`, `window_s` (600), `network_hps`, `fleet_reported` (`by_vendor` {nvidia, amd, apple, intel, unknown} with `mhs`, `workers`, `share`; `total_mhs`, `workers`, `machines`, `rows` per worker with label, machine id8, card, mhs, ages), `chain_attributed` (`by_vendor` with `blue_blocks`, `miners`, `share`, `mhs`; `blue_blocks_total`, `mapped_keys`, `unknown_miner_ids`), `by_vendor` (the two readings side by side), `coverage` (1 minus the unknown share). Fleet-reported = the sum of `now=MH/s wall` from each intake-reporting worker's newest STATUS line in the window (our fleet only); chain-attributed = each vote key's share of the window's blue blocks times the network estimate, vendor from the worker that logged the key, else unknown. `docs/benchmarks/repro.md` section 8. | | `live_state.finality` | jsonb, updated every 2 s | `params`, `chain_id`, `next_index`, `finality_active`, `latest_locked_index`, `latest_locked_hash`, `latest_locked_blue_score`, `weights` (`total_weight`, `active_weight`, `voters`, `keys[]` with `id`, `blocks`, `voter`, `participation`, `stripped_until_daa`, `revealed`). | ## Keeping it current (autosync)