diff --git a/docs/analysis/class-v6/amd-intel-energy.md b/docs/analysis/class-v6/amd-intel-energy.md
new file mode 100644
index 000000000..431e2db5b
--- /dev/null
+++ b/docs/analysis/class-v6/amd-intel-energy.md
@@ -0,0 +1,280 @@
+# AMD and Intel energy on class v6: why the measured cards pay 3x to 6x a Blackwell card per hash, and what a kernel can change
+
+8 October 2026, written 19:4x to 20:3x UK, the AMD-and-Intel energy lane. Register rows: GPU-03 (the 64-register
+window's real cost) in its AMD and Intel cells; ECO-05's named cause (`docs/analysis/class-v6/eco-05-results.md` on
+counter-asic-4: "the measured RX 9070 XT (7.9 microjoules at its knee), the RX 7600 (6.2 estimated, 8.1 measured at stock)
+and the Arc B580 (10.4, watts estimated) cost 4x to 6x a Blackwell card per joule"); the P02 cohort's AMD and Intel cells.
+The one acceptance rule is P03: a change must cost the honest card no more than 5 percent per accepted work against the
+paired baseline (and no more than 2 percent of accepted throughput) and must not help the adversary.
+
+**Labels.** Every number carries one: **measured** (an instrument read it; the instrument is named), **modelled**
+(arithmetic on labelled inputs, or an offline compiler's report standing in for the driver's), **estimated** (no
+instrument behind it), **team-reported** (read by a person, not a job). Vendor and JEDEC figures are marked
+approximate. No row here is a wall reading: neither PC has a wall meter, so every watt is the device's own telemetry
+(ADLX `GPUPower` on AMD, `nvidia-smi power.draw` on NVIDIA, the Level Zero sysman energy counter on Intel, IOReport on
+Apple), and none meets P02's calibrated 2 percent wall meter. **Efficiency kind** (review B's rule, main's order 8 October
+2026): every efficiency row says which of two things it is, and the two are never mixed: **KT**, kernel throughput (the
+worker's `--bench-pack` or `--memprobe`, device event time, no pool, no shares) against device-reported watts; **SV**, the
+serving hash rate (the installed app's own `hash_now` for the card while it mines) against device-reported watts. No row
+here is the third kind, **end-to-end accepted work per wall joule in serving mode**: that needs accepted shares and a wall
+meter, and is owed (section 4.3). Nothing was built or run on the Mac: the offline compiles
+ran on igneum-build-1 (LLVM 18.1.3), the jobs on PC 1 and PC 2 through the signed jobs queue.
+
+## 0. The answer in one paragraph
+
+On this hash every card runs at 87 to 100 percent of its own dependent random-read ceiling divided by 128 (the unit of
+work is 128 dependent random 4-byte reads), so joules per hash = 128 x watts / random reads per second. Normalised per
+32 bits of memory bus, the AMD and Intel cards burn about what the RTX 5090 burns (18 to 19 W per 32 bits at their
+knees; the RX 7600 28 W at stock) but complete 0.21x to 0.47x its random reads per second (the 5090 1.08 G/s per 32
+bits; the 9070 XT 0.30; the 7600 0.44 to 0.51; the B580 0.23). That ratio is the whole gap: 3.4x for the 9070 XT, 3.1x
+to 3.5x for the 7600, 4.4x for the B580 against the 5090's floor, and 3.8x to 6.1x against the 5080 and the 5070 Ti,
+which is ECO-05's "4x to 6x". It sits in the memory system's random-access rate (GDDR7's four channels per device
+against GDDR6's two, and a per-channel rate that varies 2.2x between the three GDDR6 cards), not in the kernel: the
+ALU work is about 5 percent of an AMD card's joules, occupancy is 3x to 6x above what saturates the DRAM, the
+64-register window costs the RX 7600 and the Arc B580 no rate per unit of work (0 and -0.3 percent, measured), and a 16-byte read costs what a 4-byte read costs on every
+card measured. The kernel changes that keep every hash identical are worth 0 to about 10 percent on AMD, ranked in
+section 3; none closes the gap, and the one hash change that does (64-byte reads) halves the 5090 and fails P03.
+The levers that move ECO-05 tonight are the AMD operating point below the driver's floor (a baseline, not a candidate)
+and the RX 7600's metered knee (a PC 1 job queued at landing; section 4), which ECO-05 says adds sustainable worlds
+at 0.03 and 0.10 if it reads under 5 microjoules (modelled tonight at 5.0 to 5.4).
+
+## 1. The measured rows as they stand, with their instruments
+
+### 1.1 Rate and watts
+
+| Card | Class / point | MH/s (rate instrument) | Watts (watts instrument) | Microjoules per hash | Label | Kind | Source |
+|---|---|---|---|---|---|---|---|
+| RX 9070 XT (gfx1201, 16 GB GDDR6, 256-bit, PC 1 eGPU) | class v2, stock, app mining | 17.73 (the app's hash_now, mean of 24) | 198.9 (ADLX GPUPower, 12 samples, 193 to 212) | 11.2 | measured | SV | docs/bench-log.md, 5 Oct, job tele-measure-1 |
+| RX 9070 XT | class v5, stock (grid point 0/0) | 18.9 (the app's hash_now, median over a 75 s hold) | 202 (ADLX GPUPower, mean after a 30 s settle) | 10.7 | measured | SV | the ADLX grid, 8 Oct 12:13, `relay/playbooks/ca3-pc1-amd-grid.ps1`; floor/denominator.md |
+| RX 9070 XT | class v5, gmax -500 MHz, plimit -30 percent (the driver's floor; the grid's best of 24) | 18.9 to 19.0 (same) | 149.3 (same) | 7.9 | measured | KT | same |
+| RX 7600 (gfx1102, 8 GB GDDR6, 128-bit, PC 1) | class v4 sub-version 3, stock | 13.88 (card-in job, worker bench) | 113 (ADLX, card-in job) | 8.14 | measured | KT | floor/denominator.md, 15:50 UK |
+| RX 7600 | class v6 hl-v6-foldrw (the partner), stock, quiet | 15.79 (worker --bench-pack) | none | | measured rate only | KT | hash lane, 18:3x UK |
+| RX 7600 | class v6 hl-v6-all (window + fold + rw; 256 loads per hash = two units of work), stock, quiet | 7.92 = **15.84 per unit of work** | none | | measured rate only | KT | hash lane, 18:3x UK |
+| RX 7600 | class v6 sizes 1 / 2 / 4 / 5.5 GiB | 15.75 / 15.46 / 15.34 / 15.35 | none | | measured rate only | KT | reference-population.md section 6 |
+| RX 7600 | class v6 at stock with the class v4 run's 113 W | 15.79 | 113 (a different class's reading) | **7.2** | modelled | KT (modelled) | this file |
+| RX 7600 | knee | 13.9 | 86 (the 9070 XT's 24 percent applied) | 6.19 | estimated | KT (modelled) | reference-population.md |
+| Arc B580 (12 GB GDDR6, 192-bit, PC 2 eGPU and a PC 1 slot) | class v4, stock | 10.6 to 11.0 (worker bench) | about 110 (the board's class) | 10.4 | measured rate, estimated watts | KT | floor/denominator.md; the Intel lane, 7 Oct |
+| Arc B580 (PC 2, enclosure) | class v6 hl-v6-foldrw, stock, quiet (SIMD32, sub-group shuffle exchange) | 11.008 (worker --bench-pack, 60 dispatches of 2^24, device time) | not read: the Level Zero energy counter answered ZE_RESULT_ERROR_UNSUPPORTED_FEATURE (0x78000003) on all three power domains, unelevated, driver 32.0.101.6733 | | measured rate only | KT | job run-ae-pc2-b580-energy-20261008, 18:52Z |
+| Arc B580 | class v6 hl-v6-all, stock, quiet (the compiler drops to SIMD16: sub-group 16, local-memory exchange with barriers) | 5.489 = **10.98 per unit of work** | not read (same) | | measured rate only | KT | same, 18:54Z |
+| Arc B580 | class v6 at stock with the board-class watts | 11.0 | about 110 | 10.0 | measured rate, estimated watts | KT | this file |
+| RTX 5090 (32 GB GDDR7, 512-bit, PC 1) | class v5, the 1,300 MHz lock | 134.8 (worker) | 314 (nvidia-smi; 2.33 x 134.8) | 2.33 | measured | KT | reference-population.md (PC 1, tiers file) |
+| RTX 5080 (16 GB GDDR7, 256-bit) | class v5, 1,100 MHz | 71.2 | 146.6 | 2.06 | measured (rented) | KT | reference-population.md |
+| RTX 5070 Ti (16 GB GDDR7, 256-bit) | class v5 knee | 77.0 | 131 | 1.70 | stock measured (rented), knee modelled | KT | reference-population.md |
+| Apple M5 Max (36 GB LPDDR5X) | class v4 | 26.67 | 37.3 (IOReport GPU and DRAM channels) | 1.40 | measured | KT | floor/denominator.md |
+
+Tonight's jobs (section 4.1) added the B580's class v6 rows; its watts could not be read unelevated. The RX 7600's metered
+class v6 rows (ADLX at stock and two knob points) are owed to a PC 1 job still queued at landing.
+
+### 1.2 The random-read ceiling (the instrument that explains the rows)
+
+`igneum-worker-opencl --memprobe` (proto-opencl/host.c): dependent random 4-byte loads over a buffer, device event
+time, best over lanes in flight; `chase` is one dependent load per step per lane.
+
+| Card | Ceiling at 1024 MiB, G loads/s | Hash-implied (MH/s x 128) | Hash / ceiling | 64-byte read against 4-byte | Label | Source |
+|---|---|---|---|---|---|---|
+| RTX 5090 | 17.5 to 18.2 (card off in the app, CUDA) | 17.25 | 0.96 | 0.52x to 0.86x (two 32-byte sectors; bandwidth-bound at 584 GB/s) | measured | bench-log, read-width entry, 5 Oct |
+| RX 9070 XT | 2.42 to 2.66 (4,096 lanes already saturate it; eight independent chains per lane give the same) | 2.42 | 0.87 to 0.95 | 1.0x (2.47 to 2.87: the line is fetched either way) | measured | bench-log, 5 Oct |
+| RX 7600 | owed (the PC 1 job of section 4.1) | 2.02 (class v6) | | | hash-implied, modelled | |
+| Arc B580 | 1.407 to 1.409 (4,096 lanes and up; eight independent chains per lane 1.41, the same); at 256 MiB 1.55 to 1.57 | 1.409 (class v6) | 1.00 | not probed | measured | job run-ae-pc2-b580-energy-20261008 (tonight); the Intel lane 7 Oct read 1.41 |
+| Apple M5 Max | 3.50 | 3.47 | 1.0 | 1.0x | measured | bench-log, 5 Oct |
+
+## 2. The decomposition
+
+### 2.1 The identity
+
+The unit of work is 128 dependent random 4-byte dataset reads (program.json of both class v6 packs:
+`loads_per_hash` 128, `load_width_counts_4_16_64` [16, 0, 0], `bytes_per_hash` 512; hl-v6-all does two units per hash).
+Every card above runs at 0.87 to 1.0 of its probe ceiling / 128 (table 1.2), so:
+
+microjoules per unit = 128 x P / A, with P the card's watts at its operating point and A its random reads per second.
+
+Splitting both by the memory bus (32-bit units: the 5090 16, the 9070 XT 8, the 7600 4, the B580 6, the 5080 and 5070 Ti
+8; vendor figures, approximate):
+
+| Card | A per 32 bits (G reads/s) | P per 32 bits (W) | Nanojoules per random read (P / A) | Microjoules per hash | Against the 5090: watts ratio x reads ratio = joules ratio | Label |
+|---|---|---|---|---|---|---|
+| RTX 5090 at the lock | 1.08 | 19.6 | 18.2 | 2.33 | 1 | A measured, P measured |
+| RTX 5080 at 1,100 MHz | 1.14 | 18.3 | 16.1 | 2.06 | 0.94 x 0.95 = 0.88 | measured (rented) |
+| RX 9070 XT at the floor point | 0.30 | 18.7 | 61.7 | 7.9 | 0.95 x 3.6 = 3.4 | measured |
+| RX 9070 XT at stock | 0.30 | 25.3 | 84 | 10.7 | 1.29 x 3.6 = 4.6 | measured |
+| RX 7600 at stock, class v4 run | 0.44 | 28.3 | 63.5 | 8.14 | 1.44 x 2.4 = 3.5 | measured |
+| RX 7600 at stock, class v6 rate | 0.51 | 28.3 | 56 | 7.2 | 1.44 x 2.1 = 3.1 | modelled (watts from the v4 run) |
+| Arc B580 at stock | 0.23 | 18.3 | 78 to 80 | 10.4 | 0.93 x 4.7 = 4.4 | A measured, P estimated |
+| Apple M5 Max (for contrast) | 0.22 | 2.3 | 10.7 | 1.40 | 0.12 x 4.9 = 0.60 | measured |
+
+Reading: at their knees the AMD and Intel cards spend about the same watts per 32 bits of memory bus as the 5090
+(0.93x to 0.95x; the 7600 at stock and the 9070 XT at stock 1.3x to 1.4x because nothing has been taken off them), and
+the whole 3x to 4.5x is the random-read rate per 32 bits. Against the 5080 and the 5070 Ti (1.70 to 2.06) the same rows
+read 3.8x to 6.1x, which is the "4x to 6x" of ECO-05. The Mac shows the other way out: a memory system with a quarter of
+the 5090's random-read rate per bus width wins on joules because it spends an eighth of the watts per bus width.
+
+### 2.2 The random-read rate per 32 bits: where the 3.6x lives
+
+| Factor | 5090 against the 9070 XT | Label |
+|---|---|---|
+| Channels per 32-bit device: GDDR7 four, GDDR6 two | 2x | JEDEC device organisation, approximate |
+| Random reads per channel per second: 5090 about 270 M, 9070 XT about 150 M (the per-channel figure of floor/denominator.md) | 1.8x | modelled from the measured ceilings and the channel counts |
+| Product | 3.6x | matches the measured 1.08 / 0.30 |
+
+Within GDDR6 the per-channel rate varies 2.2x on the same DRAM family: the RX 7600 about 220 to 250 M per channel per
+second (8 channels; hash-implied, modelled), the 9070 XT about 150 M (16 channels; measured ceiling), the B580 about 115
+M (12 channels; measured ceiling). The 7600's rate says GDDR6 itself allows 1.5x to 2.2x more random reads per channel
+than the 9070 XT and the B580 obtain, so part of their deficit sits above the DRAM, in the memory controller, the fabric,
+the last-level cache path or address translation. That part is the only piece of the gap that software might reach
+(section 3, candidate 2); the GDDR7 against GDDR6 part is hardware. Tonight's B580 probe bounds the translation share
+on Xe2: its ceiling is 1.55 to 1.57 G/s at 256 MiB and 1.41 at 1024 MiB (measured), both far beyond its 18 MB L2, so
+the footprint (page reach) costs it about 10 percent and the other 90 percent of its per-channel shortfall is the memory
+controller and fabric's.
+
+Bandwidth is not the bound on the AMD and Intel cards: the 9070 XT moves 2.42 G x 64 bytes = 155 GB/s, 24 percent of its
+640 GB/s rating, and a 64-byte read costs it exactly what a 4-byte read costs (measured); the bound is the rate of random
+accesses, each opening a DRAM row. The 5090 at 64 bytes is the opposite case (it becomes bandwidth-bound and halves).
+
+### 2.3 The kernel's memory path on RDNA 3, RDNA 4 and Xe2
+
+Instrument: the class v6 OpenCL texts (`kernel.cl` of hl-v6-all, id 0x9d40978601a7df2a, and hl-v6-foldrw, id
+0x605d06cabc489f94, from build-1 `/srv/artefacts/packs/`) compiled offline with clang 18.1.3 for amdgcn-amd-amdhsa,
+-O3, OpenCL C 1.2, `IGNEUM_EXCHANGE 0` (the local-memory path the driver takes on the 9070 XT), the work-item and
+rotate built-ins shimmed to the target's own built-ins (no libclc on the box). The driver's compiler is AMD's own LLVM
+(the PAL,LC stack) of a different version, so these are **modelled** stand-ins for the driver's binary; the driver's own
+numbers come from the host's kernel line (private memory, local memory, sub-group) and, for the 9070 XT, the 5 October
+measurement (wave32, private memory 0).
+
+| Quantity | hl-v6-all, gfx1102 | hl-v6-all, gfx1201 | hl-v6-foldrw, gfx1102 | hl-v6-foldrw, gfx1201 | Label |
+|---|---|---|---|---|---|
+| Wave size | 32 | 32 | 32 | 32 | modelled; the driver reports wave32 on gfx1201 (measured, 5 Oct) |
+| VGPRs | 160 | 160 | 96 | 96 | modelled |
+| Waves per SIMD (LLVM 18's model) | 6 | 6 (LLVM 18 models 1,024 VGPRs for gfx1201; its gfx1100 model gives 9 at 156) | 10 | 10 | modelled |
+| Spills (scratch bytes) | 0 | 0 | 0 | 0 | modelled; the driver reports private memory 0 on the 9070 XT (measured, class v2) |
+| LDS per work-group | 256 B (the exchange's two buffers) | 256 B | 256 B | 256 B | modelled |
+| Barriers emitted | 0 (work-group = one wave32: elided) | 0 | 0 | 0 | modelled; matches WAVEFRONT.md's measured "barriers elided" |
+| Global loads in the loop body | 32 x global_load_b32 | 32 | 16 x global_load_b32 | 16 | modelled |
+| Loads with a full wait right after them | all 32 (51 vmcnt(0) waits) | all 32 | 13 of 16 (three overlap) | | modelled |
+| Hash-kernel code size | 33.6 KB | 34.0 KB | 6.3 KB | 6.6 KB | modelled |
+| Integer multiplies (mul_lo, mad, mul_hi) per body | 99 | 99 | 67 | 67 | modelled |
+
+The five questions the brief asks, answered on these numbers:
+
+1. **Wave size.** Wave32 on RDNA 3 and 4 (the compiler's choice and the driver's report); a work-group of 32 is one wave,
+   so the local-memory exchange compiles to LDS writes and reads with no barrier. The exchange path and the work-group
+   shape cost nothing on the 9070 XT (`--group-warps 1, 2, 4, 8` = 18.02 to 18.07 MH/s, measured 5 Oct). Xe2: the
+   kernel takes the khr or Intel sub-group shuffle only at a queried sub-group of exactly 32; the Arc's sub-group and
+   SIMD width under the 64-register window, measured tonight on the B580 (the host's kernel line, job
+   run-ae-pc2-b580-energy-20261008): the partner hl-v6-foldrw compiles at sub-group 32 and takes `sub_group_shuffle_xor`;
+   the window pack hl-v6-all drops to sub-group 16 (SIMD16, the Intel compiler's answer to 64 live registers), so the
+   host takes the local-memory exchange with barriers (256 bytes); private memory 0 in both (no spill). The cost of that
+   drop: none in rate (10.98 MH/s per unit of work against 11.01, -0.3 percent, measured), because the card is at 100
+   percent of its random-read ceiling either way.
+2. **The 16-byte read's coalescing.** Class v6 reads 4 bytes per load (program.json above); there is no 16-byte read in
+   the frozen object. The read-width experiment's 16-byte loads (w16) cost every card what 4 bytes cost (5090 139.8
+   against 136.1 MH/s, 9070 XT 17.90 against 18.15, measured 5 Oct): a lane's random read fetches a whole line (64 bytes
+   on AMD and Apple, a 32-byte sector on NVIDIA) and the 32 lanes of a wave hit 32 different lines (a 1 GiB dataset has
+   16.8 M 64-byte lines; two lanes of a wave share one with probability about 3 x 10^-5), so there is nothing to
+   coalesce.
+3. **LDS use.** 256 bytes per wave, only the exchange: 8 exchanges per iteration in the body of hl-v6-all, 4 in
+   hl-v6-foldrw, plus 13 per shadow pass x 27 passes; about 2,840 LDS write-read pairs per unit of work. At a few
+   picojoules per lane-access that is a few nanojoules per hash against 7,900 (modelled, under 0.1 percent).
+4. **The dependent-load chain's latency hiding.** Under the full chain every load's address mixes all 64 registers,
+   including the previous load's result, so each wave has one load in flight (all 32 loads wait vmcnt(0)). It does not
+   matter on these cards: 4,096 lanes in flight already saturate the 9070 XT's DRAM (2.64 G/s at 4,096 lanes, measured),
+   and occupancy 6 on 64 CUs holds 24,576 lanes (the 7600 at 32 CUs: 12,288). The latency is hidden by waves, not by
+   loads in flight within a wave, with 3x to 6x to spare.
+5. **Occupancy under the 64-register window.** 160 VGPRs against 96: 6 waves per SIMD against 10 (modelled), no spill.
+   Measured consequence on the 7600: none in rate (15.84 per unit of work against 15.79). On NVIDIA the hash lane
+   measured the window at 96/88 registers (5090) and 104/87 (4090), occupancy 67 to 83 percent, within 5 percent per
+   load. **GPU-03's AMD cell, rate: 0 percent per unit of work on the RX 7600 (measured, quiet, rate only); energy:
+   section 4. GPU-03's Intel cell, rate: -0.3 percent per unit of work on the Arc B580 (measured, quiet, rate only),
+   with the class v6 fingerprints equal to the CUDA and Metal references on both packs (5a6ad122a71a888f and
+   59e6708e46f1e87c, self-test PASS): the B580 computes class v6 bit-exact.**
+
+### 2.4 Where an AMD card's joules go (modelled split at the 9070 XT's floor point, 149.3 W, 18.9 MH/s)
+
+| Term | Watts | Basis | Label |
+|---|---|---|---|
+| ALU work | 6 to 12 | 55.3 k lane-ops per unit of work for the partner (8 iterations x (64 + 27 x 256 shadow ops)), 42 k for the window pack (its shadow covers two units); at 5 to 10 pJ per lane-op (the M5 Max's measured marginal 6.9 pJ per counted op as the scale); 17 percent of the card's measured 6.2 T int op/s chain throughput in use | modelled |
+| DRAM array and I/O for 2.42 G random 64-byte reads per second | 10 to 20 | 4 to 8 nJ per activate-read-transfer of a 64-byte burst (public GDDR6 figures, approximate) | modelled |
+| LDS exchange | under 0.2 | section 2.3 | modelled |
+| The rest: the die and board kept awake at full memory data rate (clock trees, fabric and last-level cache, memory PHY, VRM losses, fans) | about 120 | the remainder | modelled |
+
+The kernel's own work (ALU, LDS, the loads it issues) is 15 to 30 W of 149; the remainder is the price of a whole
+graphics card serving 8 or 16 GDDR6 channels' worth of random reads. That is why the knobs that take the core down
+without touching memory (the 9070 XT's grid: 24 points, the rate flat at 18.9 MH/s, measured) are worth more than any
+kernel text, and why they stop at the driver's floor (-500 MHz, -30 percent), not at the hash's.
+
+## 3. Candidate kernel changes that keep every hash identical, ranked
+
+Common to all: none changes the hash function, so none changes the adversary's cost (the adversary's chip computes the
+function, not our kernel), and none can help the adversary; each moves only the GPU side. The equality proof for each:
+(a) the pack's self-test on the changed kernel (cache head and FNV, dataset samples, 96 of 96 vector lanes), (b) the 2^24
+batch fingerprint at base nonce 0 equal to the unchanged kernel's and to the CUDA and Metal references on the same pack
+(hl-v6-foldrw 5a6ad122a71a888f, hl-v6-all 59e6708e46f1e87c on the pre-review export; the re-export's pair when it lands),
+on every platform the change ships to, (c) the P01 vector file (`igneum-pow hash-bound --count 1000000`, staged at
+build-1 `/srv/artefacts/packs/p01-vectors/`) through the changed kernel. P03 is then a paired energy and rate run on
+every mandatory cell the change touches (an OpenCL-only change touches no CUDA card). No candidate reduces the
+specialist's advantage, so none is a G2 upgrade: they are efficiency work on the baseline.
+
+| Rank | Change | Where | Expected gain (all modelled) | The test that proves equality | Risk under P03 | Clock |
+|---|---|---|---|---|---|---|
+| 0 (a baseline, not a candidate) | The AMD operating point below the driver's floor: a lower absolute core clock or a voltage offset where ADLX exposes one (RDNA 3 and 4 Adrenalin carry a voltage offset; ADLX's manual tuning interface for it is unverified here) | the Ember tune's AMD path, not the kernel | the ALU needs about 17 percent of the 9070 XT's shader throughput and about 30 percent of the 7600's (muls at a quarter rate), so the core can lose half its clock before the rate moves: 149 W to 110 to 125 W on the 9070 XT, 7.9 to 5.8 to 6.6 microjoules; the 7600 similar in proportion | clocks never change outputs; the grid's own per-point self-test and fingerprint | none on correctness; P03 places clock savings in the baseline, so it lowers the AMD baseline (ECO-05's input) and scores nothing as a candidate | the update-return lane's AMD knob; a read of ADLX's absolute and voltage interfaces by 12:00 tomorrow |
+| 1 | Occupancy throttle: launch only the lanes that saturate the DRAM (about 4,096 to 8,192 on the 9070 XT, measured) on a fraction of the CUs, so idle CUs clock-gate and the driver's power manager sees a partly idle die | host (a `--max-groups` cap on the dispatch, about 20 lines in host.c; the persistent-warp loop already exists for variant-5 packs) | 0 to 10 percent of the card's watts, unknown until measured; the ALU budget bounds the throttle at about a quarter of the 9070 XT's CUs and half the 7600's | per-nonce outputs do not depend on the launch shape (WAVEFRONT.md: work-groups of 32 to 256 bit-exact; the 2^24 fingerprint at every cap) | a cap below the ALU need loses rate; the sweep finds the knee | host flag and a sweep job: tomorrow 15:00 |
+| 2 | Find the GDDR6 per-channel shortfall above the DRAM (the 9070 XT and the B580 at 0.45x to 0.65x the 7600's per-channel rate) and remove it if it is software's: page size and TLB reach of the 1 GiB dataset buffer, the buffer's placement, the driver's allocation flags | host allocation and driver flags | on the B580 at most about 10 percent (its probe falls 10 percent from 256 to 1024 MiB, measured tonight: that is the translation share); the 9070 XT unprobed at two sizes, the same bound expected (modelled) | allocation never changes outputs; the fingerprint | none on correctness; a different allocation could cost the NVIDIA path nothing because it would be vendor-gated | B580 read tonight (10 percent); the 9070 XT's two-size probe when it is back in the housing |
+| 3 | Non-temporal dataset loads on AMD (`__builtin_nontemporal_load` in the OpenCL text, which AMD's LLVM lowers to the non-temporal bit on RDNA 3 and the TH_LOAD_NT hint on RDNA 4): no last-level-cache or L2 allocation for lines that are not reused | OpenCL text, AMD only | 0 to 3 percent of the watts (no 64-byte fill into the 32 or 64 MB cache per read); risk of losing the 3 to 6 percent of reads that hit that cache (0 to -3 percent rate) | the hint changes no value; the fingerprint and the vectors | the rate risk; paired run decides | an emitter flag and a PC 1 pair: tomorrow 15:00 |
+| 4 | The window's address mix computed incrementally: m is linear over GF(2) in the 64 registers, so m(src) = A xor T_src xor P xor rotr(P, 1), with T_k = rotl(r_k, 63 - k), A the xor of all T_k kept up to date on each register write (two ops per write) and P the xor of T_k below src; the compiler today emits about 108 ops per load (52 rotates, 54 xors, 2 xor3) | emitter, all three dialects | removes about 20 to 25 percent of the window pack's ALU ops per unit of work, so 1 to 2 percent of an AMD card's joules and less on NVIDIA (its ALU share is smaller still) | the identity is exact; a property test in igneum-pow over random register files and every src, then the fingerprint and vectors | none expected; it must not move the 5090 or 4090 window rows by more than noise | an emitter change on the hash lane's line, after the post-review freeze: days two and three |
+| 5 | The exchange through DPP or ds_swizzle instead of LDS (RDNA), sub-group shuffles on Intel | OpenCL text, vendor-gated | under 0.1 percent (section 2.3, item 3) | the fingerprint across the exchange paths (WAVEFRONT.md's emulator table already proves path equality) | none | not worth a slot; recorded to close the question |
+
+Not a candidate (changes the hash; recorded so it is not proposed again): reads of 64 bytes per load (w64) close the
+5090 against 9070 XT gap from 7.5x to 4.1x in rate, entirely by halving the 5090 (measured, 5 Oct), so the honest NVIDIA
+cells lose about 50 percent of accepted throughput against P03's 2 percent: out. W = 8 and W = 32 are standing rejected
+knobs. The hot table and the scratch read-modify-write cost the 9070 XT 13 to 33 percent (measured): out.
+
+What this means for the vendor question: no kernel change that keeps the hash identical is worth more than about 10
+percent on AMD (candidate 1) plus whatever candidate 2 finds above the DRAM; the 3x to 4.5x is GDDR7 against GDDR6 and a
+graphics card's fixed power over few channels. Closing it by changing the hash would cost the honest NVIDIA cards more
+than P03 allows, which the standing rule forbids ("no ratio is bought with honest GPU energy").
+
+## 4. Rows for the 21:00 economics landing and the P02 cohort, with clocks
+
+### 4.1 Tonight's two jobs
+
+| Job | Machine | What it reads | Instrument | Status and clock |
+|---|---|---|---|---|
+| run-ae-pc1-7600-energy-20261008 | PC 1 (ae432dc7), RX 7600 | memprobe at 1024 MiB; hl-v6-foldrw and hl-v6-all as the load at stock, plimit -30, gmax -500 with plimit -30; reset and read-back | the kit worker's bench (rate), ADLX through igneum-gpu-telemetry.exe every 5 s, samples after 20 s (watts) | published 19:41 UK behind the hash lane's ds2g and l8off; at 20:1x UK PC 1's runner was still held by the withdrawn overnight 7600 grid job (its last upload 19:06 UTC), so none of the three had started; the follow-up reset job run-ae-pc1-7600-reset-20261008 is queued behind it |
+| fetch-ae-ze-power-20261008 + run-ae-pc2-b580-energy-20261008 | PC 2 (1ccfe586), Arc B580 | memprobe at 256 and 1024 MiB; both v6 packs at stock | the kit worker's bench (rate), ze-power.exe (Level Zero sysman energy counters, read only, built on build-1 with mingw, sha256 354abddde13a5341521eb298a5728e65675af7474823faf9b8d5613caf840d8f) | ran 19:50 to 19:54 UK (268 s, done): rates, probes and fingerprints read; the energy counter unsupported unelevated (rc 0x78000003 on all three domains, the device seen as `Intel(R) Graphics [0xe20b]`) |
+
+
+### 4.2 The rows, each with its clock
+
+| # | Row | For | Value and label | Clock |
+|---|---|---|---|---|
+| 1 | The RX 7600's class v6 rate replaces the class v4 card-in rate in the reference population | the research lane's 21:00 landing | 15.79 MH/s (measured, quiet, rate only); the stock cell 7.2 microjoules (modelled on the v4 run's 113 W) until row 2 | now |
+| 2 | The RX 7600 metered at stock and at two knob points on class v6 | the 21:00 landing; ECO-05's section 5 condition (a 7600 knee under 5 microjoules adds the worlds at 0.03 and 0.10) | OWED: the PC 1 runner had not reached the job at landing; until it reports, the cells stay 7.2 at stock (modelled on the class v4 run's 113 W) and 5.0 to 5.4 at the floor point (modelled: the 9070 XT's 24 percent, or a flat rate at -30 percent power) | when PC 1's runner reaches it, about 20 minutes after the hash lane's ds2g and l8off; the rows go to the research lane's ECO-05 batch 02 the minute they land |
+| 3 | The Arc B580 on class v6 | the 21:00 landing; ECO-05's section 5 condition (a B580 under 8 adds worlds) | rate 11.01 MH/s (measured, KT); watts NOT READ: the Level Zero sysman energy counter answers unsupported (0x78000003) unelevated on driver 32.0.101.6733, so the cell stays about 110 W estimated, 10.0 microjoules (estimated); under 8 would need under 88 W | rate landed 19:54 UK; watts owed: an unelevated IGCL telemetry read (unverified) or a wall meter, tomorrow 15:00; an elevated read needs main's exception to the nothing-elevated rule |
+| 4 | GPU-03's AMD cell: the window's cost per unit of work on the RX 7600 | the register (row 11 and GPU-03) | rate 0 percent (15.84 against 15.79, measured, KT); energy OWED with row 2 (the same job reads both packs at each point) | rate landed; energy with row 2 |
+| 5 | GPU-03's Intel cell | the register | rate -0.3 percent per unit of work (measured, KT); SIMD16 under the window, no spill; fingerprints equal on both packs; energy not read | landed 19:54 UK |
+| 6 | The named cause of ECO-05's two-vendor failure | the 21:00 landing's text and eco-05-results.md | section 0 and 2.1 of this file (per 32 bits of bus: the same watts, 0.21x to 0.47x the random reads) | landed with this file |
+| 7 | The 9070 XT | the cohort | unchanged: 7.9 at the floor point (measured); no v6 row tonight (out of the housing since 14:2x) | when it is back on the bus |
+
+### 4.3 The P02 cohort, AMD and Intel cells
+
+| P02 requirement | State tonight | Owed |
+|---|---|---|
+| At least 3 AMD retail configurations | 2 measured (RX 9070 XT 16 GB, RX 7600 8 GB); the 9060 XT and 7600 XT are in the card-in queue, not on a PC | a third AMD card on PC 1; the hash lane's card-in job runs it as it arrives |
+| An advertised 8 GB mining tier | the RX 7600 8 GB (measured; the 5.5 GiB dataset fits at 6,732 MiB of 8,176) | |
+| Usable, not nominal, memory | the 7600's 8,176 MiB (measured) | the B580's and the 9070 XT's usable figure from the host's device line |
+| Every advertised vendor in the competitive core | Intel: the Arc B580 only; its class v6 rate and fingerprints measured tonight, its watts not readable unelevated | a second Intel configuration if Intel stays advertised |
+| Calibrated wall meter, at most 2 percent | none: every AMD and Intel watt is ADLX or Level Zero device telemetry | a wall meter on each PC (a purchase and a hand at the socket); until then every row reads "device-reported" |
+| 5 paired 30-minute runs per primary cell after 15 minutes of warm-up | none: tonight's points are about 1 minute each after 20 s | the paired protocol on PC 1 and PC 2 once the 2.0 devnet mines there (GPU-02 and GPU-03 under the standard) |
+| Three unaffiliated operators | none | the served kit (D1) |
+
+## 5. Unverified and owed
+
+- The offline compiles are LLVM 18 standing in for AMD's driver compiler; the driver's VGPR count and wave size on the
+  7600 and the B580 come from tonight's host kernel lines.
+- The channel counts per device and the GDDR6 and GDDR7 energy per access are JEDEC and public figures, approximate.
+- The split of section 2.4 is modelled; only a per-rail reading (not available through ADLX) would measure it.
+- The B580's watts: the Level Zero sysman energy counter is unsupported unelevated on driver 32.0.101.6733 (measured
+  tonight); the 110 W stays an estimate until an unelevated IGCL read, a wall meter, or main's exception for one
+  elevated read.
+- The RX 7600's metered class v6 rows (job run-ae-pc1-7600-energy-20261008, queued on PC 1 behind a held runner).
+- The ALU energy per op on RDNA is taken at the M5 Max's measured scale, not measured on AMD.
+- Candidate 0's ADLX absolute-clock and voltage interfaces are unverified on the tool; candidates 1 to 4 are unmeasured.
diff --git a/docs/analysis/class-v6/amd-intel-energy/offline-compile-prelude.h b/docs/analysis/class-v6/amd-intel-energy/offline-compile-prelude.h
new file mode 100644
index 000000000..ae5258efb
--- /dev/null
+++ b/docs/analysis/class-v6/amd-intel-energy/offline-compile-prelude.h
@@ -0,0 +1,10 @@
+// shims for an offline AMDGPU compile without libclc (analysis only; never a mining kernel)
+#define get_global_id(d) ((size_t)(__builtin_amdgcn_workgroup_id_x() * 32u + __builtin_amdgcn_workitem_id_x()))
+#define get_local_id(d) ((size_t)__builtin_amdgcn_workitem_id_x())
+#define get_global_size(d) ((size_t)(1u<<20))
+#define get_sub_group_size() 32u
+#define rotate(x, n) __builtin_rotateleft32((x), (n))
+#define mul_hi(a, b) ((uint)(((ulong)(a) * (ulong)(b)) >> 32))
+#define barrier(f) do { __builtin_amdgcn_fence(__ATOMIC_RELEASE, "workgroup"); __builtin_amdgcn_s_barrier(); __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "workgroup"); } while (0)
+#define vload4(o, p) (*(const __global uint4*)((p) + 4u*(o)))
+#define vstore4(v, o, p) (*(__global uint4*)((p) + 4u*(o)) = (v))
diff --git a/docs/analysis/class-v6/amd-intel-energy/offline-isa-summary.txt b/docs/analysis/class-v6/amd-intel-energy/offline-isa-summary.txt
new file mode 100644
index 000000000..26b42b197
--- /dev/null
+++ b/docs/analysis/class-v6/amd-intel-energy/offline-isa-summary.txt
@@ -0,0 +1,34 @@
+== hl-v6-all-gfx1102 (clang 18.1.3, -O3, OpenCL C 1.2, IGNEUM_EXCHANGE 0)
+; codeLenInByte = 33576
+; NumSgprs: 18
+; NumVgprs: 160
+; ScratchSize: 0
+; LDSByteSize: 256 bytes/workgroup (compile time only)
+; Occupancy: 6
+global_load_b32 32; s_barrier 0; full waits 51; VALU 4080; v_alignbit 1657; v_xor_b32 1724; v_xor3 72; int mul 99; ds_ 42; scratch 0
+== hl-v6-all-gfx1201 (clang 18.1.3, -O3, OpenCL C 1.2, IGNEUM_EXCHANGE 0)
+; codeLenInByte = 34004
+; NumSgprs: 18
+; NumVgprs: 160
+; ScratchSize: 0
+; LDSByteSize: 256 bytes/workgroup (compile time only)
+; Occupancy: 6
+global_load_b32 32; s_barrier 0; full waits 49; VALU 4077; v_alignbit 1657; v_xor_b32 1724; v_xor3 72; int mul 99; ds_ 42; scratch 0
+== hl-v6-foldrw-gfx1102 (clang 18.1.3, -O3, OpenCL C 1.2, IGNEUM_EXCHANGE 0)
+; codeLenInByte = 6268
+; NumSgprs: 20
+; NumVgprs: 96
+; ScratchSize: 0
+; LDSByteSize: 256 bytes/workgroup (compile time only)
+; Occupancy: 10
+global_load_b32 16; s_barrier 0; full waits 25; VALU 675; v_alignbit 84; v_xor_b32 104; v_xor3 20; int mul 67; ds_ 34; scratch 0
+== hl-v6-foldrw-gfx1201 (clang 18.1.3, -O3, OpenCL C 1.2, IGNEUM_EXCHANGE 0)
+; codeLenInByte = 6584
+; NumSgprs: 20
+; NumVgprs: 96
+; ScratchSize: 0
+; LDSByteSize: 256 bytes/workgroup (compile time only)
+; Occupancy: 10
+global_load_b32 16; s_barrier 0; full waits 23; VALU 675; v_alignbit 84; v_xor_b32 104; v_xor3 20; int mul 67; ds_ 34; scratch 0
+b3314b187c414925e52febb3683120cf7f8db2c8f4b802816b964d705c33c571  ../hl-v6-all/hl-v6-all/kernel.cl
+b66c22e859244f86f164e05841858ed1d39b42f1ab6611504872c652fb02ee23  ../hl-v6-foldrw/hl-v6-foldrw/kernel.cl
diff --git a/docs/analysis/class-v6/amd-intel-energy/run-ae-pc1-7600-energy-20261008.ps1 b/docs/analysis/class-v6/amd-intel-energy/run-ae-pc1-7600-energy-20261008.ps1
new file mode 100644
index 000000000..79cf1506e
--- /dev/null
+++ b/docs/analysis/class-v6/amd-intel-energy/run-ae-pc1-7600-energy-20261008.ps1
@@ -0,0 +1,107 @@
+# AMD-and-Intel energy lane (Igneum 2.0 register: GPU-03's AMD cells, ECO-05's owed RX 7600 knee), 8 October 2026.
+# The three-card Windows rig's RX 7600 (OpenCL gfx1102), measure only, unelevated: the class v6 kit's own igneum-worker-opencl.exe is the load
+# (--bench-pack, no app involvement: the app is never quit, paused or resumed, no api/cards, no app tune), the rebuilt
+# igneum-gpu-telemetry.exe (ADLX GPUPower, else GPUTotalBoardPower) samples every 5 s beside it, at three knob points:
+# stock, plimit -30, gmax -500 + plimit -30 (the 9070 XT's best). ADLX manual tuning is reset at the end and a read-back
+# sample is printed. First: --memprobe at 1024 MiB (the card's dependent random-read ceiling, the 9070 XT's and B580's
+# instrument). Every result line starts with RESULT; SUMMARY {json} ends it. The sampler is a child process whose pid is
+# written to this job's folder and stopped by that pid only.
+$ErrorActionPreference = 'Continue'
+$jobName = 'pc1-7600-energy'
+$started = Get-Date
+function Stamp { (Get-Date).ToUniversalTime().ToString('yyyy-MM-ddTHH:mm:ssZ') }
+function Summary([string] $status, [hashtable] $extra) {
+  $o = [ordered]@{ job = $jobName; status = $status; duration_s = [int]((Get-Date) - $started).TotalSeconds; finished_at = (Stamp) }
+  foreach ($k in $extra.Keys) { $o[$k] = $extra[$k] }
+  'SUMMARY ' + ($o | ConvertTo-Json -Compress -Depth 5)
+}
+"RESULT start $(Stamp) job=$jobName machine=$env:COMPUTERNAME app_version=$env:IGNEUM_APP_VERSION"
+$jobs = Split-Path $env:IGNEUM_JOB_DIR
+$kitId = $env:IGNEUM_V6_KIT_ID; if (-not $kitId) { $kitId = 'fetch-class-v6-kit-20261008' }
+$kit = Join-Path $jobs $kitId
+$exe = Join-Path $kit 'bin\windows\igneum-worker-opencl.exe'
+$packs = Join-Path $kit 'packs'
+if (-not (Test-Path $exe)) { "RESULT error worker missing at $exe (fetch job $kitId first)"; Summary 'failed' @{ error = 'worker missing' }; exit 2 }
+"RESULT worker $exe sha256 $((Get-FileHash -Algorithm SHA256 $exe).Hash.ToLower()) kit=$kitId"
+# the telemetry tool: the newest job-folder copy that prints a tune line, else the installed one (the grid playbook's rule)
+$install = Join-Path $env:LOCALAPPDATA 'Programs\Igneum Miner'
+$tool = $null; $cands = @()
+$jobsDir = Join-Path $env:LOCALAPPDATA 'igneum\app\jobs'
+if (Test-Path -LiteralPath $jobsDir) { $cands += @(Get-ChildItem -LiteralPath $jobsDir -Recurse -Filter 'igneum-gpu-telemetry*.exe' -ErrorAction SilentlyContinue | Sort-Object LastWriteTime -Descending | ForEach-Object { $_.FullName }) }
+$cands += Join-Path $install 'igneum-gpu-telemetry.exe'
+foreach ($c in $cands) { if (-not (Test-Path -LiteralPath $c)) { continue }; $pr = @(& $c --tune 2>&1 | ForEach-Object { "$_" }) | Where-Object { $_ -match '^tune \d+ ' }; if ($pr) { $tool = $c; break } }
+if (-not $tool) { "RESULT error no telemetry tool with a tune line"; Summary 'failed' @{ error = 'no tool' }; exit 2 }
+"RESULT tool $tool sha256 $((Get-FileHash -LiteralPath $tool -Algorithm SHA256).Hash.ToLower())"
+$tune = @(& $tool --tune 2>&1 | ForEach-Object { "$_" })
+$tune | ForEach-Object { "RESULT tune $_" }
+$line = $tune | Where-Object { $_ -match '^tune (\d+) name "([^"]*7600[^"]*)" .* ok\s*$' } | Select-Object -First 1
+if (-not $line) { "RESULT error no RX 7600 tune line"; Summary 'failed' @{ error = 'no 7600 tune line' }; exit 2 }
+$ord = [int]([regex]::Match($line, '^tune (\d+)').Groups[1].Value)
+$gr = [regex]::Match($line, 'gmax_range (-?\d+) (-?\d+)'); $prr = [regex]::Match($line, 'plimit_range (-?\d+) (-?\d+)')
+"RESULT card ordinal=$ord gmax_range=$($gr.Groups[1].Value)..$($gr.Groups[2].Value) plimit_range=$($prr.Groups[1].Value)..$($prr.Groups[2].Value)"
+# the OpenCL device of the 7600
+$list = @(& $exe --list 2>&1 | ForEach-Object { "$_" }); $list | ForEach-Object { "RESULT list $_" }
+$dev = $null; foreach ($l in $list) { if ($l -match '^\s*\[(\d+)\].*gfx1102' -and $l -notmatch 'dup') { $dev = [int]$Matches[1]; break } }
+if ($null -eq $dev) { "RESULT error no gfx1102 in --list"; Summary 'failed' @{ error = 'no gfx1102' }; exit 2 }
+$w = @(Get-CimInstance Win32_Process -Filter "Name = 'igneum-worker-opencl.exe'" -ErrorAction SilentlyContinue | ForEach-Object { "$($_.ProcessId):[$($_.CommandLine -replace '\s+', ' ')]" })
+$loaded = ($w | Where-Object { $_ -match "--device\s+$dev(\s|$)" }).Count -gt 0
+"RESULT device $dev gfx1102 card_state=$(if ($loaded) { 'loaded (another worker on the card: the rows are labelled loaded)' } else { 'quiet' }) workers=[$($w -join ' ')]"
+function SampleLine() { @(& $tool 2>&1 | ForEach-Object { "$_" }) | Where-Object { $_ -match ('^amd ' + $ord + ' ') } | Select-Object -First 1 }
+"RESULT idle_sample $(Stamp) $(SampleLine)"
+# 1. the random-read ceiling at the dataset size
+$mp = @(& $exe --memprobe --probe-mib 1024 --device $dev 2>&1 | ForEach-Object { "$_" })
+$mp | ForEach-Object { "RESULT memprobe $_" }
+# 2. the three knob points x the two packs, the bench as the load, the sampler beside it
+$sampler = Join-Path $env:IGNEUM_JOB_DIR 'sampler.ps1'
+$pidFile = Join-Path $env:IGNEUM_JOB_DIR 'sampler.pid'
+Set-Content -LiteralPath $sampler -Value @'
+param([string] $tool, [int] $ord, [string] $out)
+while ($true) { $l = @(& $tool 2>&1 | ForEach-Object { "$_" }) | Where-Object { $_ -match ('^amd ' + $ord + ' ') } | Select-Object -First 1; Add-Content -LiteralPath $out -Value ((Get-Date).ToUniversalTime().ToString('o') + ' ' + $l); Start-Sleep -Seconds 5 }
+'@
+function SetPoint([int] $g, [int] $p) {
+  $a = @(& $tool --card $ord --set-gmax $g 2>&1 | ForEach-Object { "$_" }) | Where-Object { $_ -match '^tune ' } | Select-Object -First 1
+  $b = @(& $tool --card $ord --set-plimit $p 2>&1 | ForEach-Object { "$_" }) | Where-Object { $_ -match '^tune ' } | Select-Object -First 1
+  return @{ ok = (($a -match ' ok\s*$') -and ($b -match ' ok\s*$')); lines = ($a + ' | ' + $b) }
+}
+$rows = @()
+$points = @(@(0, 0), @(0, -30), @(-500, -30))
+foreach ($pt in $points) {
+  $g = $pt[0]; $p = $pt[1]
+  if ($g -lt [int]$gr.Groups[1].Value -or $p -lt [int]$prr.Groups[1].Value) { "RESULT point gmax=$g plimit=$p skipped (outside the card's range)"; continue }
+  $set = SetPoint $g $p
+  "RESULT point gmax=$g plimit=$p set ok=$($set.ok) $($set.lines)"
+  if (-not $set.ok) { continue }
+  foreach ($pk in @('hl-v6-foldrw', 'hl-v6-all')) {
+    $d = Join-Path $packs $pk
+    $nb = if ($pk -eq 'hl-v6-all') { 40 } else { 80 }
+    $samp = Join-Path $env:IGNEUM_JOB_DIR ("samples-$pk-g$g-p$p.txt")
+    $sp = Start-Process -FilePath 'powershell.exe' -ArgumentList @('-NoProfile', '-ExecutionPolicy', 'Bypass', '-File', $sampler, '-tool', $tool, '-ord', $ord, '-out', $samp) -WindowStyle Hidden -PassThru
+    Set-Content -LiteralPath $pidFile -Value $sp.Id
+    $t0 = Get-Date
+    $out = @(& $exe --bench-pack --pack $d --device $dev --batch-log2 24 --batches $nb 2>&1 | ForEach-Object { "$_" })
+    $t1 = Get-Date
+    $spid = [int](Get-Content -LiteralPath $pidFile); Stop-Process -Id $spid -Force -ErrorAction SilentlyContinue; Remove-Item -LiteralPath $pidFile -ErrorAction SilentlyContinue
+    foreach ($l in $out) { if ($l -match '^(pack |class |RESULT |FAIL|error|warm-up|kernel|exchange)') { "RESULT bench $pk g=$g p=$p out $l" } }
+    $res = $out | Where-Object { $_ -match '^RESULT ' } | Select-Object -Last 1
+    $fp = ''; $mhs = 0; $check = ''
+    if ($res -match 'fingerprint=([0-9a-f]{16})') { $fp = $Matches[1] }
+    if ($res -match 'mhs=([0-9.]+)') { $mhs = [double]$Matches[1] }
+    if ($res -match 'check=(\w+)') { $check = $Matches[1] }
+    # the watts: samples from 20 s after the bench started (the dataset build and warm-up excluded) to its end
+    $ws = @(); $gc = @(); $raw = @()
+    if (Test-Path -LiteralPath $samp) { foreach ($s in Get-Content -LiteralPath $samp) { $raw += $s; $ts = [datetime]::Parse(($s -split ' ')[0]).ToUniversalTime(); if ($ts -ge $t0.ToUniversalTime().AddSeconds(20) -and $ts -le $t1.ToUniversalTime()) { $m = [regex]::Match($s, ' watts (-?[\d.]+)'); if ($m.Success -and [double]$m.Groups[1].Value -gt 0) { $ws += [double]$m.Groups[1].Value }; $c = [regex]::Match($s, ' gclk_mhz (-?[\d.]+)'); if ($c.Success) { $gc += [double]$c.Groups[1].Value } } } }
+    $raw | Select-Object -First 3 | ForEach-Object { "RESULT sample $pk g=$g p=$p $_" }
+    $wm = if ($ws.Count) { [math]::Round((($ws | Measure-Object -Average).Average), 1) } else { 0 }
+    $gm = if ($gc.Count) { [math]::Round((($gc | Measure-Object -Average).Average), 0) } else { 0 }
+    $uj = if ($mhs -gt 0 -and $wm -gt 0) { [math]::Round($wm / $mhs, 3) } else { 0 }
+    "RESULT ROW card=RX7600 pack=$pk gmax_off=$g plimit=$p mhs=$mhs watts=$wm uj_per_hash=$uj gclk=$gm samples=$($ws.Count) seconds=$([int]($t1 - $t0).TotalSeconds) fingerprint=$fp check=$check $(Stamp)"
+    $rows += [ordered]@{ pack = $pk; g = $g; p = $p; mhs = $mhs; watts = $wm; uj = $uj; gclk = $gm; samples = $ws.Count; fingerprint = $fp; check = $check }
+  }
+}
+$r = @(& $tool --card $ord --reset 2>&1 | ForEach-Object { "$_" }) | Where-Object { $_ -match '^tune ' } | Select-Object -First 1
+"RESULT reset $r"
+Start-Sleep -Seconds 3
+"RESULT after_reset_sample $(Stamp) $(SampleLine)"
+"RESULT after_reset_tune $((@(& $tool --tune 2>&1 | ForEach-Object { "$_" }) | Where-Object { $_ -match ('^tune ' + $ord + ' ') } | Select-Object -First 1))"
+Summary $(if ($rows.Count) { 'done' } else { 'failed' }) @{ rows = $rows; device = $dev; ordinal = $ord; kit = $kitId }
+exit $(if ($rows.Count) { 0 } else { 1 })
diff --git a/docs/analysis/class-v6/amd-intel-energy/run-ae-pc1-7600-reset-20261008.ps1 b/docs/analysis/class-v6/amd-intel-energy/run-ae-pc1-7600-reset-20261008.ps1
new file mode 100644
index 000000000..432b6aa60
--- /dev/null
+++ b/docs/analysis/class-v6/amd-intel-energy/run-ae-pc1-7600-reset-20261008.ps1
@@ -0,0 +1,23 @@
+# AMD-and-Intel energy lane, 8 October 2026: the guaranteed ADLX reset of the RX 7600 after run-ae-pc1-7600-energy-20261008
+# (a timeout, a killed job tree or a lost relay must never leave the card capped). Unelevated; the card's ordinal from the
+# tool's own tune line (the integrated Radeon's all-dash line is not a card); --reset, then a tune-line read-back.
+$ErrorActionPreference = 'Continue'
+$install = Join-Path $env:LOCALAPPDATA 'Programs\Igneum Miner'
+$tool = $null; $cands = @()
+$jobsDir = Join-Path $env:LOCALAPPDATA 'igneum\app\jobs'
+if (Test-Path -LiteralPath $jobsDir) { $cands += @(Get-ChildItem -LiteralPath $jobsDir -Recurse -Filter 'igneum-gpu-telemetry*.exe' -ErrorAction SilentlyContinue | Sort-Object LastWriteTime -Descending | ForEach-Object { $_.FullName }) }
+$cands += Join-Path $install 'igneum-gpu-telemetry.exe'
+foreach ($c in $cands) { if (-not (Test-Path -LiteralPath $c)) { continue }; $pr = @(& $c --tune 2>&1 | ForEach-Object { "$_" }) | Where-Object { $_ -match '^tune \d+ ' }; if ($pr) { $tool = $c; break } }
+if (-not $tool) { 'RESULT error no telemetry tool with a tune line'; 'SUMMARY {"job":"pc1-7600-reset","status":"failed"}'; exit 2 }
+$line = @(& $tool --tune 2>&1 | ForEach-Object { "$_" }) | Where-Object { $_ -match '^tune (\d+) name "([^"]*7600[^"]*)" .* ok\s*$' } | Select-Object -First 1
+if (-not $line) { 'RESULT error no RX 7600 tune line'; 'SUMMARY {"job":"pc1-7600-reset","status":"failed"}'; exit 2 }
+$ord = [int]([regex]::Match($line, '^tune (\d+)').Groups[1].Value)
+"RESULT before $line"
+$r = @(& $tool --card $ord --reset 2>&1 | ForEach-Object { "$_" }) | Where-Object { $_ -match '^tune ' } | Select-Object -First 1
+"RESULT reset $r"
+Start-Sleep -Seconds 2
+$after = @(& $tool --tune 2>&1 | ForEach-Object { "$_" }) | Where-Object { $_ -match ('^tune ' + $ord + ' ') } | Select-Object -First 1
+"RESULT after $after"
+$ok = ($after -match ' gmax 0 ') -and ($after -match ' plimit 0 ')
+"SUMMARY {""job"":""pc1-7600-reset"",""status"":""$(if ($ok) { 'done' } else { 'check' })"",""ordinal"":$ord}"
+exit 0
diff --git a/docs/analysis/class-v6/amd-intel-energy/run-ae-pc2-b580-energy-20261008.ps1 b/docs/analysis/class-v6/amd-intel-energy/run-ae-pc2-b580-energy-20261008.ps1
new file mode 100644
index 000000000..1e1c416e3
--- /dev/null
+++ b/docs/analysis/class-v6/amd-intel-energy/run-ae-pc2-b580-energy-20261008.ps1
@@ -0,0 +1,70 @@
+# AMD-and-Intel energy lane (Igneum 2.0 register: GPU-03's Intel cell, ECO-05's owed B580 watts), 8 October 2026.
+# The second Windows rig's Intel Arc B580, measure only, unelevated, no knob written (Intel has none in our tools): the class v6 kit's own
+# igneum-worker-opencl.exe is the load (--bench-pack; the app is never quit, paused or resumed, no api/cards), and
+# ze-power.exe (fetch job fetch-ae-ze-power-20261008, sha256 354abddd...; Level Zero sysman energy counters through the
+# driver's ze_loader.dll, read only) samples every 5 s beside it. First: --memprobe at 256 and 1024 MiB (the dependent
+# random-read ceiling; two sizes to see whether the rate falls with the footprint beyond the caches, a TLB-reach sign).
+# The sampler is a child process whose pid is written to this job's folder and stopped by that pid only.
+$ErrorActionPreference = 'Continue'
+$env:IGNEUM_V6_KIT_ID = 'fetch-class-v6-kit-20261008'
+$jobName = 'pc2-b580-energy'
+$started = Get-Date
+function Stamp { (Get-Date).ToUniversalTime().ToString('yyyy-MM-ddTHH:mm:ssZ') }
+function Summary([string] $status, [hashtable] $extra) {
+  $o = [ordered]@{ job = $jobName; status = $status; duration_s = [int]((Get-Date) - $started).TotalSeconds; finished_at = (Stamp) }
+  foreach ($k in $extra.Keys) { $o[$k] = $extra[$k] }
+  'SUMMARY ' + ($o | ConvertTo-Json -Compress -Depth 5)
+}
+"RESULT start $(Stamp) job=$jobName machine=$env:COMPUTERNAME app_version=$env:IGNEUM_APP_VERSION"
+$jobs = Split-Path $env:IGNEUM_JOB_DIR
+$ze = Join-Path (Join-Path $jobs 'fetch-ae-ze-power-20261008') 'ze-power.exe'
+if (-not (Test-Path $ze)) { $ze = @(Get-ChildItem -LiteralPath $jobs -Recurse -Filter 'ze-power.exe' -ErrorAction SilentlyContinue | Select-Object -First 1 | ForEach-Object { $_.FullName }) | Select-Object -First 1 }
+if (-not $ze -or -not (Test-Path $ze)) { "RESULT error ze-power.exe missing (fetch job fetch-ae-ze-power-20261008 first)"; Summary 'failed' @{ error = 'no sampler' }; exit 2 }
+"RESULT sampler $ze sha256 $((Get-FileHash -Algorithm SHA256 $ze).Hash.ToLower())"
+@(& $ze 2>&1 | ForEach-Object { "$_" }) | ForEach-Object { "RESULT ze_idle $_" }
+# the kit: IGNEUM_V6_KIT_ID, else the newest job folder holding the worker and packs\hl-v6-foldrw
+$kit = $null
+if ($env:IGNEUM_V6_KIT_ID) { $kit = Join-Path $jobs $env:IGNEUM_V6_KIT_ID }
+if (-not $kit -or -not (Test-Path $kit)) { $kit = @(Get-ChildItem -LiteralPath $jobs -Directory -ErrorAction SilentlyContinue | Where-Object { (Test-Path (Join-Path $_.FullName 'packs\hl-v6-foldrw\kernel_bound.cl')) -and (Test-Path (Join-Path $_.FullName 'bin\windows\igneum-worker-opencl.exe')) } | Sort-Object LastWriteTime -Descending | ForEach-Object { $_.FullName }) | Select-Object -First 1 }
+if (-not $kit) { "RESULT error no class v6 kit on this PC"; Summary 'failed' @{ error = 'no kit' }; exit 2 }
+$exe = Join-Path $kit 'bin\windows\igneum-worker-opencl.exe'; $packs = Join-Path $kit 'packs'
+"RESULT worker $exe sha256 $((Get-FileHash -Algorithm SHA256 $exe).Hash.ToLower()) kit=$(Split-Path $kit -Leaf)"
+$list = @(& $exe --list 2>&1 | ForEach-Object { "$_" }); $list | ForEach-Object { "RESULT list $_" }
+$dev = $null; foreach ($l in $list) { if ($l -match '^\s*\[(\d+)\].*(B580|Arc|Battlemage)' -and $l -notmatch 'dup') { $dev = [int]$Matches[1]; break } }
+if ($null -eq $dev) { "RESULT error no Arc B580 in --list"; Summary 'failed' @{ error = 'no B580' }; exit 2 }
+$w = @(Get-CimInstance Win32_Process -Filter "Name = 'igneum-worker-opencl.exe'" -ErrorAction SilentlyContinue | ForEach-Object { "$($_.ProcessId):[$($_.CommandLine -replace '\s+', ' ')]" })
+$loaded = ($w | Where-Object { $_ -match "--device\s+$dev(\s|$)" }).Count -gt 0
+"RESULT device $dev card_state=$(if ($loaded) { 'loaded' } else { 'quiet' }) workers=[$($w -join ' ')]"
+foreach ($mib in @(256, 1024)) { @(& $exe --memprobe --probe-mib $mib --device $dev 2>&1 | ForEach-Object { "$_" }) | ForEach-Object { "RESULT memprobe $mib $_" } }
+$pidFile = Join-Path $env:IGNEUM_JOB_DIR 'ze.pid'
+$rows = @()
+foreach ($pk in @('hl-v6-foldrw', 'hl-v6-all')) {
+  $d = Join-Path $packs $pk
+  if (-not (Test-Path (Join-Path $d 'kernel_bound.cl'))) { "RESULT error pack $pk missing"; continue }
+  $nb = if ($pk -eq 'hl-v6-all') { 30 } else { 60 }
+  $samp = Join-Path $env:IGNEUM_JOB_DIR ("ze-$pk.txt")
+  $sp = Start-Process -FilePath $ze -ArgumentList @('-l', '5') -RedirectStandardOutput $samp -WindowStyle Hidden -PassThru
+  Set-Content -LiteralPath $pidFile -Value $sp.Id
+  $t0 = Get-Date
+  $out = @(& $exe --bench-pack --pack $d --device $dev --batch-log2 24 --batches $nb 2>&1 | ForEach-Object { "$_" })
+  $t1 = Get-Date
+  $zpid = [int](Get-Content -LiteralPath $pidFile); Stop-Process -Id $zpid -Force -ErrorAction SilentlyContinue; Remove-Item -LiteralPath $pidFile -ErrorAction SilentlyContinue
+  foreach ($l in $out) { if ($l -match '^(pack |class |RESULT |FAIL|error|warm-up|kernel|exchange)') { "RESULT bench $pk out $l" } }
+  $res = $out | Where-Object { $_ -match '^RESULT ' } | Select-Object -Last 1
+  $fp = ''; $mhs = 0; $check = ''
+  if ($res -match 'fingerprint=([0-9a-f]{16})') { $fp = $Matches[1] }
+  if ($res -match 'mhs=([0-9.]+)') { $mhs = [double]$Matches[1] }
+  if ($res -match 'check=(\w+)') { $check = $Matches[1] }
+  # ze lines arrive every 5 s; the lines after the first 20 s of the bench (warm-up and dataset build excluded): the sampler
+  # started with the bench, so the first four sample rounds are dropped
+  $lines = @(); if (Test-Path -LiteralPath $samp) { $lines = @(Get-Content -LiteralPath $samp) }
+  $lines | Select-Object -First 6 | ForEach-Object { "RESULT ze $pk $_" }
+  $dom = @{}
+  $round = @{}
+  foreach ($l in $lines) { if ($l -match '^intel (\d+) dom (\d+) card (\d) watts ([\d.]+)') { $k = "$($Matches[1])/$($Matches[2])/card$($Matches[3])"; if (-not $round.ContainsKey($k)) { $round[$k] = 0 }; $round[$k]++; if ($round[$k] -gt 4) { if (-not $dom.ContainsKey($k)) { $dom[$k] = @() }; $dom[$k] += [double]$Matches[4] } } }
+  $dm = [ordered]@{}
+  foreach ($k in $dom.Keys) { $dm[$k] = [math]::Round((($dom[$k] | Measure-Object -Average).Average), 1); "RESULT ROW card=ArcB580 pack=$pk domain=$k mhs=$mhs watts=$($dm[$k]) uj_per_hash=$(if ($mhs -gt 0) { [math]::Round($dm[$k] / $mhs, 3) } else { 0 }) samples=$($dom[$k].Count) seconds=$([int]($t1 - $t0).TotalSeconds) fingerprint=$fp check=$check $(Stamp)" }
+  $rows += [ordered]@{ pack = $pk; mhs = $mhs; watts = $dm; fingerprint = $fp; check = $check }
+}
+Summary $(if ($rows.Count) { 'done' } else { 'failed' }) @{ rows = $rows; device = $dev }
+exit $(if ($rows.Count) { 0 } else { 1 })
diff --git a/docs/analysis/class-v6/amd-intel-energy/ze-power.c b/docs/analysis/class-v6/amd-intel-energy/ze-power.c
new file mode 100644
index 000000000..1a7f99c8b
--- /dev/null
+++ b/docs/analysis/class-v6/amd-intel-energy/ze-power.c
@@ -0,0 +1,98 @@
+/* ze-power: Intel GPU power from the Level Zero sysman energy counters (AMD-and-Intel energy lane, 8 October 2026).
+ * Windows: loads ze_loader.dll (installed by the Intel graphics driver) at run time; no SDK, no headers.
+ *   ze-power.exe            one reading: two energy-counter reads 1 s apart per power domain
+ *   ze-power.exe -l N       a line every N seconds until killed
+ * Line: intel <dev> dom <k> card <0|1> watts <W> energy_uj <E> ts_us <T> name "<first printable strings>"
+ * The counter is the device's own (package or card domain, as the driver exposes it), not the wall. Read only:
+ * no set call exists in this program. */
+#include <stdio.h>
+#include <stdint.h>
+#include <string.h>
+#include <stdlib.h>
+#include <windows.h>
+typedef int32_t zr;
+typedef void *H;
+typedef struct { uint64_t energy; uint64_t timestamp; } ecnt;
+typedef zr (__cdecl *f_init)(uint32_t);
+typedef zr (__cdecl *f_get)(uint32_t *, H *);
+typedef zr (__cdecl *f_get2)(H, uint32_t *, H *);
+typedef zr (__cdecl *f_card)(H, H *);
+typedef zr (__cdecl *f_cnt)(H, ecnt *);
+typedef zr (__cdecl *f_props)(H, void *);
+#define MAXD 8
+#define MAXP 8
+static void names(H dev, f_props gp, char *out, size_t cap) {
+    out[0] = 0;
+    if (!gp) return;
+    static unsigned char buf[8192];
+    memset(buf, 0, sizeof buf);
+    *(uint32_t *)(buf + 0) = 0x1;   /* ZES_STRUCTURE_TYPE_DEVICE_PROPERTIES */
+    *(uint32_t *)(buf + 16) = 0x3;  /* core: ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES */
+    if (gp(dev, buf) != 0) return;
+    size_t o = 0; int run = 0; size_t start = 0;
+    for (size_t i = 0; i < sizeof buf && o + 2 < cap; i++) {
+        unsigned char c = buf[i];
+        if (c >= 32 && c < 127) { if (!run) { start = i; run = 1; } }
+        else { if (run && i - start >= 4) { size_t n = i - start; if (o + n + 2 >= cap) break; if (o) out[o++] = '|'; memcpy(out + o, buf + start, n); o += n; } run = 0; }
+    }
+    out[o] = 0;
+}
+int main(int argc, char **argv) {
+    int loop = 0;
+    if (argc >= 3 && strcmp(argv[1], "-l") == 0) loop = atoi(argv[2]);
+    HMODULE m = LoadLibraryA("ze_loader.dll");
+    if (!m) { printf("error no ze_loader.dll (%lu)\n", GetLastError()); return 2; }
+    f_init zesInit = (f_init)GetProcAddress(m, "zesInit");
+    f_get zesDriverGet = (f_get)GetProcAddress(m, "zesDriverGet");
+    f_get2 zesDeviceGet = (f_get2)GetProcAddress(m, "zesDeviceGet");
+    f_get2 enumPwr = (f_get2)GetProcAddress(m, "zesDeviceEnumPowerDomains");
+    f_card cardPwr = (f_card)GetProcAddress(m, "zesDeviceGetCardPowerDomain");
+    f_cnt getE = (f_cnt)GetProcAddress(m, "zesPowerGetEnergyCounter");
+    f_props gp = (f_props)GetProcAddress(m, "zesDeviceGetProperties");
+    f_init zeInit = (f_init)GetProcAddress(m, "zeInit");
+    f_get zeDriverGet = (f_get)GetProcAddress(m, "zeDriverGet");
+    f_get2 zeDeviceGet = (f_get2)GetProcAddress(m, "zeDeviceGet");
+    if (!enumPwr || !getE) { printf("error the loader has no sysman power entry points\n"); return 2; }
+    H drv[MAXD], dev[MAXD * 4]; uint32_t nd = 0, ndev = 0; const char *path = "zesInit";
+    zr r = zesInit ? zesInit(0) : -1;
+    if (r == 0 && zesDriverGet && zesDeviceGet) {
+        uint32_t c = MAXD; if (zesDriverGet(&c, drv) == 0) nd = c;
+        for (uint32_t i = 0; i < nd; i++) { uint32_t k = MAXD; if (zesDeviceGet(drv[i], &k, dev + ndev) == 0) ndev += k; }
+    }
+    if (ndev == 0 && zeInit && zeDriverGet && zeDeviceGet) {
+        path = "zeInit+ZES_ENABLE_SYSMAN"; SetEnvironmentVariableA("ZES_ENABLE_SYSMAN", "1");
+        r = zeInit(0);
+        uint32_t c = MAXD; if (r == 0 && zeDriverGet(&c, drv) == 0) nd = c;
+        for (uint32_t i = 0; i < nd; i++) { uint32_t k = MAXD; if (zeDeviceGet(drv[i], &k, dev + ndev) == 0) ndev += k; }
+    }
+    printf("info path %s init %d drivers %u devices %u\n", path, (int)r, nd, ndev);
+    if (ndev == 0) { printf("error no sysman device\n"); return 3; }
+    H pw[MAXD * 4][MAXP + 1]; uint32_t np[MAXD * 4]; int iscard[MAXD * 4][MAXP + 1]; char nm[MAXD * 4][256];
+    for (uint32_t d = 0; d < ndev; d++) {
+        names(dev[d], gp, nm[d], sizeof nm[d]);
+        uint32_t k = MAXP; np[d] = 0;
+        if (enumPwr(dev[d], &k, pw[d]) == 0) np[d] = k;
+        for (uint32_t j = 0; j < np[d]; j++) iscard[d][j] = 0;
+        H c = 0;
+        if (cardPwr && cardPwr(dev[d], &c) == 0 && c) { pw[d][np[d]] = c; iscard[d][np[d]] = 1; np[d]++; }
+        printf("info dev %u domains %u name \"%s\"\n", d, np[d], nm[d]);
+    }
+    ecnt prev[MAXD * 4][MAXP + 1];
+    for (uint32_t d = 0; d < ndev; d++) for (uint32_t j = 0; j < np[d]; j++) { memset(&prev[d][j], 0, sizeof(ecnt)); getE(pw[d][j], &prev[d][j]); }
+    int every = loop > 0 ? loop : 1;
+    for (;;) {
+        Sleep(every * 1000);
+        for (uint32_t d = 0; d < ndev; d++) for (uint32_t j = 0; j < np[d]; j++) {
+            ecnt e; memset(&e, 0, sizeof e);
+            zr q = getE(pw[d][j], &e);
+            double dt = (double)(e.timestamp - prev[d][j].timestamp) * 1e-6, de = (double)(e.energy - prev[d][j].energy) * 1e-6;
+            double w = dt > 0 ? de / dt : 0;
+            printf("intel %u dom %u card %d watts %.2f energy_uj %llu ts_us %llu rc %d name \"%s\"\n", d, j, iscard[d][j], w,
+                   (unsigned long long)e.energy, (unsigned long long)e.timestamp, (int)q, nm[d]);
+            prev[d][j] = e;
+        }
+        fflush(stdout);
+        if (loop <= 0) break;
+    }
+    return 0;
+}