diff --git a/docs/bench-log.md b/docs/bench-log.md
index 754a7c6f..39d053b9 100644
--- a/docs/bench-log.md
+++ b/docs/bench-log.md
@@ -1988,7 +1988,27 @@ Branch `ca3-reserve`, worker "reserve" (`docs/plans/counter-asic-3-reserve.md` c
 
 Reading of the Mac rows. The run-to-run spread is under 4% on every row. The dot4 rows reproduce the 5 October figures (1.6x unsigned, 4.7x signed), which is the check on the method. A step cost under 1.00 means the family's op plus the glue is cheaper than the five-op reference chain: the reference's two registers are a tighter dependency than the three-register candidate chains, and Apple's shifts, extract, andn and select each cost about what an xor costs. Three rows cost more than the reference: `perm` (1.13: no byte-permute function in MSL; the `uchar4` swizzle compiles to shifts and masks, so a byte permute is emulated on Apple at about the price of the live `rotr`), `clz` (1.01) and `shfla` (1.91: a shuffle by a computed lane index costs 2.2x the live xor shuffle on Apple, `simd_shuffle` against `simd_shuffle_xor`; the second shuffle form is the one candidate Apple pays for). `mm8` as a chain on Apple is owed (Metal 4 `matmul2d`; this toolchain is Swift 5.8 without the tensor API).
 
-**RTX 5090 (PC 2, 1ccfe586), CUDA**: PENDING the PC 2 job (`run-ca3-family-pc2-20261006`, published only after `/tmp/igneum-devnet/pc2-ca3.clear` and under the mkdir lock; the card to itself: `--stop-miners`, prover off for the run). The rows land in this entry when the closing report is read.
+**RTX 5090 (PC 2, 1ccfe586), CUDA** (job `run-ca3-family-pc2-20261006`, a signed `run` job with `--stop-miners`, published 08:41:45Z after `/tmp/igneum-devnet/pc2-ca3.clear` (08:24:27Z) under the mkdir lock (taken 08:41:26Z, released 08:43:24Z the moment the closing report was read); ran 08:42:27Z to 08:43:03Z, done, exit 0, 36 s; `node tools/jobs.mjs run-ca3-family-pc2-20261006 --all`). The card to itself: the app had stopped the miner before the script started (`workers_before`: no CUDA compute app, the card at 847 MHz SM and 72 W), the script posted the card off through `api/cards` in both key forms (settings.json carries two NVIDIA keys, `nvidia:0:NVIDIA GeForce RTX 5090` with 8 identities and the older `nvidia:NVIDIA GeForce RTX 5090` with 2) and read the card quiet by nvidia-smi's compute-apps list and the process list after 30 s; prover off at 08:42:27Z and back on at 08:43:02Z (`{"ok":true}`, in the finally block); the cards restored to their settings. nvcc 12.8 in WSL2, `-arch=sm_120`, the source sha256 `1a3d07b8...0cf90d` equal on the Mac, the Windows side and inside WSL. Three runs of `./family-probe --reps 3`, CUDA event time; the SM clock ramped from 862 MHz at run 1 to 2,572 MHz at run 3 (gpu_before per run; 129 W at the end), so the best of the three runs is the card's warm figure and the table carries it; the run-to-run spread of the best values is under 3% on every row except `shl` (12%: run 2 caught the ramp). Driver 610.47:
+
+| kernel | best ms, runs 1 / 2 / 3 | best of the three, ms | G steps/s (best) | ns per step (best) | ops per step counted | step cost (ratio to `alu`, best) | bit-exact, 3 runs |
+|---|---|---|---|---|---|---|---|
+| alu | 0.553 / 0.541 / 0.553 | 0.541 | 7,941 | 132 | 5 | 1.00 | yes |
+| rotr (live) | 0.729 / 0.719 / 0.715 | 0.715 | 6,005 | 175 | 5 | 1.32 | yes |
+| shflx (live) | 0.808 / 0.817 / 0.817 | 0.808 | 5,315 | 197 | 6 | 1.49 | yes |
+| shl | 0.687 / 0.770 / 0.696 | 0.687 | 6,253 | 168 | 5 | 1.27 | yes |
+| shr | 0.698 / 0.697 / 0.691 | 0.691 | 6,212 | 169 | 5 | 1.28 | yes |
+| bfe (`bfe.u32`, inline PTX) | 0.837 / 0.851 / 0.835 | 0.835 | 5,142 | 204 | 5 | 1.54 | yes |
+| bfec (C form) | 0.836 / 0.837 / 0.835 | 0.835 | 5,144 | 204 | 5 | 1.54 | yes |
+| andn | 0.680 / 0.700 / 0.693 | 0.680 | 6,313 | 166 | 5 | 1.26 | yes |
+| perm (`__byte_perm`) | 0.706 / 0.715 / 0.718 | 0.706 | 6,084 | 172 | 5 | 1.30 | yes |
+| popc | 0.819 / 0.813 / 0.835 | 0.813 | 5,283 | 198 | 5 | 1.50 | yes |
+| clz | 0.897 / 0.898 / 0.884 | 0.884 | 4,856 | 216 | 5 | 1.63 | yes |
+| sel | 0.723 / 0.713 / 0.729 | 0.713 | 6,024 | 174 | 5 | 1.32 | yes |
+| shfla (lane + 3) | 0.845 / 0.848 / 0.829 | 0.829 | 5,179 | 202 | 6 | 1.53 | yes |
+| dot4i (`__dp4a`) | 0.677 / 0.656 / 0.625 | 0.625 | 6,870 | 153 | 5 | 1.16 | yes |
+| mm8 (`mma.sync.m8n8k16.u8`, one per warp per step) | 1.313 / 1.314 / 1.350 | 1.313 | 3,272 | 321 | 1 mma + 4 | 2.43 | yes |
+
+Reading of the 5090 rows. Every row is bit-exact, `mm8` included, so the m8n8k16 fragment layout of the CPU reference (PTX ISA 9.4 section 9.7.16.5.3) is the layout the hardware uses. The `alu` chain reads 7,941 G steps/s here against 8,754 through OpenCL event time on 5 October: a CUDA event pair around a 0.54 ms kernel carries about 0.05 ms of launch, which also compresses every ratio toward 1 (approximate: the ratios are the card's at the 2% level, not better). On this card every candidate costs more than the reference chain, unlike Apple: the 5090 runs the two-register add-xor-rotate chain at one IMAD and one funnel shift per step, and the three-register candidate chains pay their glue. Against the live `rotr` (1.32), the candidates read: `andn` 0.95x, `shl` 0.96x, `shr` 0.97x, `perm` 0.98x, `sel` 1.00x, `popc` 1.14x, `shfla` 1.16x (the same as the live xor shuffle, 1.49: the indexed shuffle costs NVIDIA nothing extra), `bfe` 1.17x, `clz` 1.23x. `bfe.u32` and the C form cost the same to the nanosecond (0.835 ms), so the compiler emits the same code for both and no single-instruction bit-field extract is in play on this architecture (not checked by cuobjdump; the equal times are the evidence). `dp4a` reads 1.16x (1.17x on 5 October). `mm8` is the most expensive row on NVIDIA too (2.43x the reference: one tensor-core mma per warp per dependent step, latency-bound), which supports its place at the end of the reserve on the honest-card side as well as on the chip side.
 
 **RX 9070 XT (PC 1, ae432dc7), OpenCL**: OWED. PC 1 is the project lead's desk and not released today (the brief's rule); the OpenCL twin of the probe (`__builtin_amdgcn_*` paths for `v_bfe_u32`, `v_perm_b32`, `v_bcnt_u32_b32`, `v_cndmask_b32`, `ds_bpermute_b32`) is the next job on that card.
 
@@ -1997,7 +2017,7 @@ Consequences per tier, Mac rows (the hash is latency-bound by 128 dependent DRAM
 | Tier | What the rows mean | What is being done |
 |---|---|---|
 | Apple user (M-series laptop or desktop, the app's Metal worker) | six of the seven candidates (`shl`, `shr`, `bfe`, `andn`, `popc`, `sel`) cost at most the live `rotr` step; `clz` the same as the reference; `perm` 1.13x (emulated); `shfla` 1.91x, the only candidate over the live `shfl`'s cost by more than 2x on this card. At 4 points of 64 a 1.91x op costs under 1% of the program's ALU time, itself a small share of a latency-bound hash (approximate: argued, measured when live) | the proposed order puts `shfla` after the plain datapath families (R6), so Apple pays it last; `mm8` stays last |
-| NVIDIA user (8 to 32 GB card) | pending the 5090 rows above | the PC 2 job |
+| NVIDIA user (8 to 32 GB card) | every candidate is native and costs 0.95x to 1.23x the live `rotr` step (`andn`, the shifts, `perm`, `sel` under 1.0x; `popc` 1.14x; `shfla` 1.16x; `bfe` 1.17x; `clz` 1.23x); nothing on this card is emulated above a compiler sequence; at 4 points of 64 no family moves the ALU time by over 1% (argued) on a hash bound by DRAM reads | the family-live measurement at each unlock rehearsal; nothing to change in the order for NVIDIA |
 | AMD user (RX 9070 XT, 16 GB) | owed: no row today | the PC 1 job when the desk is free |
 | A rig or a pool user | the same per-card figures; no family changes the dependent-read bound | nothing until a family is live |
 | A chip | every candidate but `mm8` is a 32-bit datapath structure (barrel shifter, byte crossbar, popcount tree, 32-lane shuffle crossbar: `docs/plans/counter-asic-3-reserve.md` section 3 names them with approximate areas); none is licensable as a block the way an int8 matrix unit is | the reserve order of that document |
diff --git a/docs/plans/counter-asic-3-reserve.md b/docs/plans/counter-asic-3-reserve.md
index 3e803534..fdf99aa1 100644
--- a/docs/plans/counter-asic-3-reserve.md
+++ b/docs/plans/counter-asic-3-reserve.md
@@ -6,27 +6,30 @@
 
 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) |
+| 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) | 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) |
+| 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; 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.
+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.
@@ -34,7 +37,7 @@ Sources read 6 October 2026: PTX ISA 9.4 (https://docs.nvidia.com/cuda/parallel-
 | 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) |
+| 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 |
@@ -66,10 +69,10 @@ A 12-op chip (the ledger's M1: a sequencer over the live families and 8 register
 | 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 |
+| 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); 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 |
+| 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.
@@ -118,7 +121,7 @@ The `andn` and `sel` entries carry a note for the generator: both are injecting
 | 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) |
+| 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) |
@@ -128,7 +131,7 @@ The `andn` and `sel` entries carry a note for the generator: both are injecting
 
 | 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 |
+| 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 |