igneum-network/igneum
1
0
Fork 0
Code Issues Pull requests Releases Activity

Merge ca2-analysis (c9ebe7b) into release-0.3.11: docs and standalone probe sources the evidence rows cite (C38); nothing a built artefact reads

Browse source 
# Conflicts:
#	docs/bench-log.md
This commit is contained in:
igneum-labs 2026-10-05 23:03:45 +00:00
commit bfdc632a14
7 changed files with 1018 additions and 0 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

175
docs/analysis/int8-matrix-family.md Normal file
View file

@ -0,0 +1,175 @@
# Layer 7: the integer matrix family (INT8 x INT8 into INT32) as a reserved instruction family, design
5 October 2026 (night), Counter ASIC 2.0 (`docs/plans/counter-asic-2.md`, layer 7), branch `ca2-analysis`. Design
only: nothing here touches the generator, a vector or a node. Every figure is cited (vendor document, URL, section) or
measured (machine, date, command) or labelled approximate.
## 1. The primitive per vendor, from the vendor documents
| Vendor, hardware | Per-lane dot4 (4 bytes x 4 bytes into a 32-bit integer) | Warp or wave matrix (int8 tiles, int32 accumulate) | Source |
|---|---|---|---|
| NVIDIA, sm_61 and later (Pascal on) | PTX `dp4a.atype.btype d, a, b, c` with `.atype = .btype = {.u32, .s32}`: "Four-way byte dot product which is accumulated in 32-bit result"; semantics `d = c; for i in 0..3: d += Va[i] * Vb[i]` with the bytes sign- or zero-extended by type; introduced in PTX ISA 5.0, "Requires sm_61 or higher". CUDA: `__device__ int __dp4a(int srcA, int srcB, int c)` ("Four-way signed int8 dot product with int32 accumulate") and the unsigned form, plus `char4`/`uchar4` overloads | `mma.sync` with `.u8`/`.s8` A and B and `.s32` C and D: shape `.m8n8k16` "requires sm_75 or higher" (Turing on, PTX 6.5); shapes `.m16n8k16` and `.m16n8k32` require sm_80 (Ampere on, PTX 7.0); sparse `.m16n8k32` and `.m16n8k64` with `.u8`/`.s8` also exist | PTX ISA 9.4, section 9.7.1.24 (dp4a) and 9.7.16.5 (mma), https://docs.nvidia.com/cuda/parallel-thread-execution/index.html ; CUDA Math API, integer intrinsics, https://docs.nvidia.com/cuda/cuda-math-api/cuda_math_api/group__CUDA__MATH__INTRINSIC__INT.html ; read 5 October 2026 |
| AMD RDNA 3 (gfx11) | `v_dot4_i32_iu8` (VOP3P; each operand signed or unsigned by a per-operand bit, optional clamp) reached from clang/HIP/OpenCL C as `__builtin_amdgcn_sudot4(bool a_signed, int a, bool b_signed, int b, int acc, bool clamp)` (LLVM feature `dot8-insts`: "Has v_dot4_i32_iu8, v_dot8_i32_iu4 instructions"); `v_dot4_u32_u8` as `__builtin_amdgcn_udot4` (`dot7-insts`: "Has v_dot4_u32_u8, v_dot8_u32_u4"); `v_dot4_i32_i8` as `__builtin_amdgcn_sdot4` (`dot1-insts`: "Has v_dot4_i32_i8 and v_dot8_i32_i4"). gfx11's common feature set carries dot7, dot8, dot9, dot10 and dot12 | `V_WMMA_I32_16X16X16_IU8`: `__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32` and `_w64` (feature `wmma-256b-insts`), a 16x16x16 tile per wave | LLVM `clang/include/clang/Basic/BuiltinsAMDGPU.td` (main, read 5 October 2026), lines defining `__builtin_amdgcn_sdot4`, `udot4`, `sudot4`, `wmma_i32_16x16x16_iu8_w32`; AMD GPUOpen, "How to accelerate AI applications on RDNA 3 using WMMA", https://gpuopen.com/learn/wmma_on_rdna3/ ; the RDNA 3 ISA PDF itself did not download tonight (AMD's CDN refused curl and the fetcher timed out), so the instruction names are from the compiler and GPUOpen, not quoted from the ISA guide |
| AMD RDNA 4 (gfx12, the 9070 XT) | the same `sudot4` and `udot4` builtins: LLVM's `FeatureISAVersion12_Generic` carries `FeatureDot7Insts` and `FeatureDot8Insts` and not `FeatureDot1Insts`, so `__builtin_amdgcn_sdot4` is NOT exposed on gfx12 and `sudot4` with both operands signed is the signed form to use | `__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12` and `_w64_gfx12` (feature `wmma-128b-insts`): the int8 WMMA exists on RDNA 4 with a narrower per-lane operand (2 ints per lane for A and B against 4 on RDNA 3); AMD's RDNA 4 WMMA guide names the same builtin | LLVM `llvm/lib/Target/AMDGPU/AMDGPU.td` (`FeatureISAVersion12_Generic`) and `BuiltinsAMDGPU.td` (main, 5 October 2026); AMD GPUOpen, "WMMA guide for AMD RDNA 4 architecture GPUs, part 2", https://gpuopen.com/learn/wmma-guide-amd-rdna-4-gpus-part-2/ ; the RDNA 4 ISA guide (AMD document 70651, April 2025) was not readable tonight (docs.amd.com returned 401 to a direct fetch) |
| AMD CDNA 3 (MI300) | the same VOP3P dot instructions (approximate: not checked in the CDNA 3 guide tonight) | `V_MFMA_I32_16X16X32_I8` and `V_MFMA_I32_32X32X16_I8` (opcodes 87 and 86 in the VOP3P-MFMA table) | AMD Instinct MI300 CDNA 3 ISA Reference Guide, 5 August 2025, https://www.amd.com/content/dam/amd/en/documents/instinct-tech-docs/instruction-set-architectures/amd-instinct-mi300-cdna3-instruction-set-architecture.pdf (downloaded and grepped 5 October 2026) |
| AMD, OpenCL on Adrenalin (Windows) | `cl_khr_integer_dot_product` is NOT in the 24 extensions Adrenalin lists for gfx1201 (the list: fp64, the int32 and int64 atomics, 3d image writes, byte addressable store, fp16, gl sharing, amd device attribute query, amd media ops and media ops2, d3d10, d3d11 and dx9 sharing, image2d from buffer, subgroups, gl event, depth images, mipmap image and writes, amd copy buffer p2p); the platform is OpenCL 2.1 so the OpenCL C 3.0 feature macro `__opencl_c_integer_dot_product_input_4x8bit` is not expected. What IS reachable: the Adrenalin OpenCL compiler is clang (driver string `PAL,LC`), and `__builtin_amdgcn_sudot4` from OpenCL C has been shown to emit `V_DOT4_I32_IU8` on a Radeon 780M (gfx1103, RDNA 3, driver 32.0.31041, Windows 11) at 2.7x the scalar fallback (1.27 to 3.42 TMAC/s) | not from OpenCL C | Adrenalin 26.9.2 extension list, https://geeks3d.com/20260904/amd-radeon-adrenalin-26-9-x-graphics-driver/ ; the OpenCL C route: https://github.com/1640675651/CPPminer/pull/1 (third party, one machine; the 9070 XT run of this document's probe is the check) ; `cl_khr_integer_dot_product` itself: OpenCL C 3.0 specification section 6.2.2.16, `int dot(char4, char4)` and `int dot_acc_sat(char4, char4, int)`, https://registry.khronos.org/OpenCL/specs/3.0-unified/html/OpenCL_Ext.html |
| Apple, Metal (MSL 4.1, 4 June 2026) | none. MSL has no dp4a or packed byte dot product: the built-in `dot(T x, T y)` is a geometric function on floating-point vectors (section 6.9); the integer functions of section 6.4 have no dot form. A per-lane dot4 is scalar emulation (section 4 below measures it) | `simdgroup_matrix<T, 8, 8>` exists for T = half, bfloat (Metal 3.1 and later) and float only (section 2.4: "T is half, bfloat ... or float"); no integer SIMD-group matrix. BUT Metal 4's tensor operation `mpp::tensor_ops::matmul2d` (section 7.2.1, table 7.3, "MatMul2D data type supported") lists A `char` x B `char` into C `int` (Metal 4) and `uchar` x `uchar` into `int` (Metal 4 and OS 26.4), plus `char` x `int4b_format` into `int`. So Apple has an exact int8 x int8 into int32 matrix path, on tensors (device or threadgroup memory, or a `cooperative_tensor` per SIMD-group or threadgroup), not on registers, and only through Metal 4's tensor API. Which GPU families run it in hardware (the M5's neural accelerators) against emulation is in the Metal Feature Set Tables, which the spec defers to and which were not read tonight | Metal Shading Language Specification version 4.1, https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf , sections 2.4, 6.9, 7.2.1 table 7.3 (PDF downloaded and text-extracted 5 October 2026) |
The Apple finding, stated plainly: the brief's expectation ("Apple has no int8 matrix or dot path") is half right. There
is no per-lane dot4 and no integer `simdgroup_matrix`. There is an exact `char x char -> int` matmul2d in Metal 4
(table 7.3). Two things about it are unverified tonight and matter for conformance: whether the int accumulate wraps or
saturates (the spec text I extracted says nothing either way; a vector at the int32 edge on the M5 settles it), and the
feature-set table (which Apple GPUs run it natively). What is settled: a generator op that is a per-lane dot4 has no
Apple intrinsic and costs scalar emulation; a generator op that is a whole-unit 8x8x16 or 16x16x16 int8 tile has a
native path on all three vendors (mma.sync on sm_75+, WMMA on RDNA 3 and 4, matmul2d on Metal 4), with Apple's path
living in a different API shape (tensors, not register fragments).
## 2. The family's semantics as a generator op (integer only, bit-exact)
Two forms are proposed; the reserve can hold both as separate families or one.
### 2.1 `dot4`: per-lane
```
dot4 dst = dst + dot4_u8(src, src2)
where dot4_u8(a, b) = sum over i in 0..3 of byte_i(a) * byte_i(b), bytes zero-extended, sum modulo 2^32
```
- Bytes are UNSIGNED. Reason, measured below: on Apple the unsigned emulation costs 1.6 ALU-chain steps per dot4
and the signed one 4.7 (section 4), while NVIDIA (`dp4a.u32.u32`) and AMD (`V_DOT4_U32_U8`, `udot4`, `dot7-insts`)
carry the unsigned form natively as they carry the signed one. Signed bytes buy nothing for the hash (the input is a
pseudo-random register) and cost the vendor without the intrinsic 3x more.
- Accumulation wraps modulo 2^32 like every other op in section 1 (spec 1.14 item 5). The maximum dot of four unsigned
bytes is 4 x 255 x 255 = 260,100, so no single dot4 overflows; the wrap is in the running sum, which is why the
AMD `clamp` bit and the OpenCL `dot_acc_sat` form are NOT the primitive (saturation would change results).
- Operands: `dst`, `src`, `src2` with `src != dst` as for `mad`; `src2` may equal either.
- Verifier: one closed-form integer expression per lane; the register-major interpreter of 1.11 adds four byte
multiplies and adds per lane. The CPU reference in the probes (`dot4_ref` in `proto-opencl/dot4-probe.c`) is this
expression.
### 2.2 `mm8`: the 32-lane unit as one int8 tile
The 32 lanes of a unit (spec 1.9) hold, in `src`, a 4-byte row fragment of an 8 x 16 int8 matrix A and, in `src2`,
a 4-byte column fragment of a 16 x 8 int8 matrix B, in exactly the layout of PTX `mma.m8n8k16` with `.u8` operands
(PTX ISA 9.4 section 9.7.16.5, "Matrix Fragments for mma.m8n8k16", the integer-type layout):
```
lane l (0..31): A[row = l >> 2][k = 4 * (l & 3) .. 4 * (l & 3) + 3] = the 4 bytes of src (byte 0 = lowest k)
B[k = 4 * (l & 3) .. +3][col = l >> 2] = the 4 bytes of src2
result C[r][c] = sum over k in 0..15 of A[r][k] * B[k][c] (uint8 x uint8, 16 products, exact, at most 1,040,400)
mm8 dst = dst + C[l >> 2][2 * (l & 3) + bit] bit = an immediate 0 or 1 drawn by the generator
```
Every lane receives one of the two C elements its lane position owns in the PTX fragment (`c0` for bit 0, `c1` for
bit 1), added into `dst` modulo 2^32. The whole op is a function of the unit's `src` and `src2` across all 32 lanes,
like `shfl`, so it needs the unit to be exactly 32 logical lanes (the wave64 rule of 1.9 applies: a wave64 device
holds two units and the local-memory path is used).
How each vendor runs it:
| Vendor | Native form | Cost per `mm8` (approximate until measured) |
|---|---|---|
| NVIDIA sm_75+ | one `mma.sync.aligned.m8n8k16.row.col.s32.u8.u8.s32` per warp, A and B fragments straight from `src` and `src2`, C = 0 in, `c0`/`c1` out, one add | one tensor instruction plus one add |
| AMD RDNA 3 and 4 | one `V_WMMA_I32_16X16X16_IU8` per wave32 with the 8x16 and 16x8 tiles zero-padded into 16x16 (the WMMA fragment layout differs from PTX's: a fixed permutation of bytes between lanes, which is a few `ds_bpermute` or `v_perm` operations, bit-exact) | one WMMA plus the permutation and the pad |
| AMD CDNA | `V_MFMA_I32_16X16X32_I8` with padding | as above |
| Apple, Metal 4 | `matmul2d<descriptor(8, 8, 16)>` on `uchar` A and B into an `int` cooperative tensor (table 7.3 row "uchar, uchar, int", OS 26.4), the fragments written from registers into a threadgroup tensor first (32 lanes x 8 bytes = 256 bytes), the C element read back per lane | one tensor op plus two threadgroup round trips; on Apple GPUs without the neural accelerators the runtime's emulation, unmeasured |
| Any vendor, fallback | 16 scalar byte products per lane after gathering the 16 bytes of B's column from the 4 lanes that hold them (4 shuffles or one 64-byte threadgroup exchange) | 4 shuffles plus 4 `dot4` emulations: on Apple about 4 x 1.6 = 6.4 ALU steps plus the shuffles (approximate, from the probe) |
Verifier: the unit evaluates C as 8 x 8 x 16 = 1,024 unsigned byte products once per `mm8` instruction and hands
each lane its element. That is 1,024 multiply-adds per instruction per unit, against 64 x 8 = 512 instructions per
hash: at W_new = 4 (section 3) a program carries about 2.6 `mm8` per iteration, 21 per hash, 21,500 multiply-adds per
unit per hash, under 10 microseconds on one core (approximate), far inside the 0.63 ms the verifier already spends per
unit (spec 1.11). The simulation stays exact because every product and sum is an integer with a defined wrap.
### 2.3 Which form to reserve
`mm8` is the one that takes matrix hardware at GPU scale from a chip (the plan's layer 7 row): a chip without tensor
units pays 1,024 products per unit per instruction where a GPU pays one tensor instruction. `dot4` is a per-lane ALU op
that a chip matches with four 8-bit multipliers, which is cheap silicon; it adds little chip resistance and costs Apple
emulation. Recommendation: reserve `mm8`; keep `dot4` out, or in only as `dot4_u8` behind `mm8`.
## 3. The genesis reserve entry (spec text for 1.13.2)
Proposed wording, to go under 1.13.2 as the first named reserve family once the conformance runs of section 5 pass:
> Reserve family R1, `mm8` (integer matrix). Semantics: section 2.2 of `docs/analysis/int8-matrix-family.md`,
> uint8 operands from `src` and `src2` in the m8n8k16 fragment layout, one int32 element of C per lane selected by
> the immediate `bit`, added into `dst` modulo 2^32. Weight at unlock `W_new = 4` points, taken proportionally from the
> ten live non-load families (the load weight and count are untouched, 1.13.1). Edge vectors, each a hand-built unit
> run on every vendor: all bytes 0xFF in A and B (C = 16 x 65,025 = 1,040,400 everywhere); all bytes 0x80 (C = 16 x
> 16,384 = 262,144); A all zero (C = 0); `dst` = 0xFFFFFFFF with a nonzero C (the wrap); alternating 0x00 and 0xFF by
> lane (the fragment mapping: C[r][c] nonzero only where the row and column bytes meet); `bit` = 0 and 1 on the same
> fragments. Unlock: at the start of era n = 4 (two years after genesis, DAA 62,208,000), or earlier by the 90%
> signalling path of section 5.7; never by a release.
The era-4 choice is deliberate: two years is long enough for the three vendors' tensor paths (and Apple's Metal 4
feature-set coverage) to be in every miner's driver, and short enough to land before any chip built against the
launch instruction set has paid back (approximate; a chip programme is 12 to 24 months, approximate, from memory).
Reserve rule for a vendor that can only emulate. Spec 1.13.2 as written requires conformance on every vendor; it says
nothing about cost. Proposed addition:
> A family enters the reserve when it is bit-exact on every vendor of 1.15. A vendor that reaches the result only by
> emulation (no instruction or library path) does not block entry if the measured penalty of the emulation on that
> vendor, on the family's own probe (a dependent chain of the op, G ops/s against the same vendor's integer ALU chain),
> is at most 8x per op, AND the family's weight at unlock keeps the emulating vendor's hash-rate loss under 5% on the
> memory-hard hash (the hash is latency-bound, so a per-op penalty on 4% of the instructions is a small fraction of a
> hash whose time is 128 dependent DRAM reads; the 5% is checked on the vendor's card with the family live, not
> computed). A family whose emulation exceeds either bound stays out of the reserve until the vendor ships a path.
With tonight's numbers: on Apple the unsigned `dot4` emulation is 1.6x per op (inside the bound); the signed one 4.7x
(inside, but why pay it); `mm8` through Metal 4's matmul2d is a path, not an emulation, and its cost is owed.
## 4. dp4a-class throughput, measured so far
Probe: a dependent chain of one dot4 per step per lane (`acc = dot4(x, y, acc); x = x * K + acc; y = rotl(y, 7) ^
(acc + s)`), 1,048,576 lanes x 4,096 steps, best of 3, device time, bit-exact against a CPU reference on two lanes
per run, beside the ALU chain of the 9070 XT bench-log entry (`x = x * K + rotl(y, 7); y = (y ^ x) + s`, 5 ops per
step counted). Sources: `proto-metal/dot4-probe.swift` (Metal), `proto-opencl/dot4-probe.c` (OpenCL: scalar, the
`cl_khr_integer_dot_product` `dot`, AMD `__builtin_amdgcn_sudot4`, NVIDIA inline PTX `dp4a.s32.s32`),
`proto-cuda/dot4-probe.cu` (CUDA `__dp4a` and the scalar emulation, for a PC with nvcc). Each OpenCL variant is built on
its own and a variant the platform cannot compile prints a "build failed" row.
| Card, API | Date, command | ALU chain, G steps/s | dot4 signed emulation, G dot4/s | dot4 unsigned emulation, G dot4/s | dot4 intrinsic, G dot4/s | Penalty of the emulation per op (ALU steps per dot4) | ok (bit-exact) |
|---|---|---|---|---|---|---|---|
| Apple M5 Max, Metal | 5 October 2026 20:0x UTC, `with-lock.sh measure ./dot4-probe` (swiftc -O), GPU start-to-end time | 879.8 (4.882 ms) | 188.2 (22.82 ms) | 548.2 (7.834 ms) | none exists | signed 4.7x, unsigned 1.6x | yes, all three kernels |
| Apple M5 Max, Apple OpenCL 1.2 | same, `with-lock.sh measure ./dot4-probe-cl --device 0`, event time | 871.5 (4.928 ms) | 188.4 (22.80 ms) | not in this probe | `cl_khr_integer_dot_product` not listed; the kernel using `dot(char4, char4)` compiled anyway and ran at 846 G/s but MISMATCHED the CPU reference on every lane checked (Apple's `dot` on char4 is not an integer dot; the extension macro must gate it) | signed 4.6x | alu and dot4e yes; dot4_khr NO |
| RTX 5090 (PC 1, ae432dc7), NVIDIA OpenCL 3.0 CUDA, driver 617.14 | 5 October 2026 20:29 UTC, job `run-dot4-20261005` (`relay/playbooks/dot4-probe.ps1`, both mining cards switched off in the app first, restored after; `node tools/jobs.mjs run-dot4-20261005`), event time | 8,753.5 (0.491 ms) | 1,239.1 (3.466 ms) | not in the OpenCL probe | 7,453.6 (0.576 ms) via inline PTX `dp4a.s32.s32` | emulation 7.1x; the intrinsic 1.17x (the chain is one dp4a plus 3 ops against 5 ops), so the emulation costs 6.0x the instruction | yes, all three |
| RX 9070 XT (PC 1, gfx1201, eGPU), AMD OpenCL 2.0 AMD-APP 3683.0 (PAL,LC) | same job, same time | 701.4 (6.124 ms) | 480.8 (8.932 ms) | not in the OpenCL probe | 664.3 (6.465 ms) via `__builtin_amdgcn_sudot4(true, a, true, b, acc, false)`: the Adrenalin OpenCL C compiler accepts the clang builtin and emits `v_dot4_i32_iu8` | emulation 1.46x; the intrinsic 1.06x, so the emulation costs 1.38x the instruction | yes, all three; the older 3652.0 platform entry for the same card gave 696.2 / 501.7 / 683.6 |
| Ryzen 9800X3D gfx1036 (PC 1, integrated RDNA 2, 2 CUs) | same job | 40.6 (105.9 ms) | 15.8 (272.3 ms) | | `sudot4` does not build: "needs target feature dot8-insts" (RDNA 2 has `dot1-insts`' `v_dot4_i32_i8`, the `sdot4` builtin, which the probe did not try) | emulation 2.6x | alu and dot4e yes |
| Every PC device | `cl_khr_integer_dot_product` not listed on NVIDIA (OpenCL 3.0) or AMD (2.0); the pragma draws "unknown OpenCL extension" on both and the `dot(char4, char4)` kernel does not build | | | | | | |
Reading across the three cards. Per dot4 at the hardware rate: the 5090 does 7.45 T dot4/s (one `dp4a` per step, 0.85
of its ALU-chain step rate), the 9070 XT 0.66 T (0.95 of its ALU-chain rate), the M5 Max 0.55 T at best (the unsigned
emulation; no instruction). On the ALU chain the 5090 is 12.5x the 9070 XT and 10x the M5 Max; on hardware dot4 it is
11.2x the 9070 XT, so the family does not widen the AMD gap, and 13.6x the M5 Max, so Apple's emulation widens its gap
by 1.4x on this op (approximate: one probe shape, the ratios of best-of-3 numbers). The signed emulation is where the
vendors differ most: 7.1x the ALU step on NVIDIA, 4.7x on Apple, 1.46x on AMD (AMD's compiler and byte-permute
hardware make the four sign-extended products nearly free; the NVIDIA OpenCL compiler does not pattern-match the
emulation into `dp4a`, which the 6x gap between `dot4e` and `dot4_nv` shows). None of this is a hash-rate number: the
hash is bound by 128 dependent DRAM reads, and a family at W_new = 4 adds about 21 of these ops per hash per lane
against about 1.2 microseconds of memory latency per hash per lane (approximate), so the per-op penalties above turn
into hash-rate losses well under 5% on every card, to be measured with the family live.
Reading of the Mac numbers. The ALU chain's 880 G steps/s on the M5 Max is the integer baseline (5 ops per step
counted, so about 4.4 T int ops/s, approximate; the 5090's 8,754 G steps/s is about 43.8 T, against the whitepaper's
104.8 peak INT32 TOPS which counts a multiply-add as two). A signed dot4 emulated as `int4(as_type<char4>(a))` products costs
4.7 of those steps; the unsigned form 1.6 steps. The 3x gap between the two is the sign extension (Metal lowers the
unsigned byte extraction to masks that fold into the multiplies, approximate reading of the result, not of the
compiled code). Both are far under the 8x bound of section 3, and the hash spends its time on DRAM reads, so a per-lane
`dot4` family would cost Apple a few percent at W_new = 4 (to be measured with the family live, not computed). The
Apple OpenCL `dot(char4, char4)` mismatch is the kind of thing the edge vectors of section 3 exist to catch.
## 5. What is owed or unverified
| Item | State |
|---|---|
| dp4a throughput on the RTX 5090 through NVIDIA OpenCL inline PTX | measured (section 4); the CUDA `__dp4a` form (`proto-cuda/dot4-probe.cu`) is unrun (no nvcc job tonight) and is a cross-check, not a gap |
| `sudot4` on the 9070 XT through Adrenalin's OpenCL C; `cl_khr_integer_dot_product` on the 3683.0 platform | measured: the builtin works and emits the instruction; the extension is not listed and the `dot(char4, char4)` kernel does not build |
| `sdot4` (`dot1-insts`) on RDNA 2 (gfx1036) | not tried; the probe only carries `sudot4` |
| Metal 4 `matmul2d` uchar x uchar into int on the M5 Max: wrap or saturate at the int32 edge, native or emulated, throughput | owed (a second Metal probe; the API needs a tensor set-up the dot4 probe does not have) |
| Metal Feature Set Tables: which Apple GPU families run int8 matmul2d natively | not read tonight |
| RDNA 3 and RDNA 4 ISA guides: the instruction text itself (names taken from LLVM and GPUOpen) | AMD's CDN refused the downloads tonight |
| `mm8` on AMD: the exact byte permutation between the PTX m8n8k16 fragment layout and the RDNA WMMA 16x16x16 layout | design, to be written with the kernel |
| The hash-rate cost of the family live at W_new = 4 on each vendor (the 5% rule of section 3) | owed, needs the generator change (not tonight) |
| Edge vectors of section 3 as files | owed, with the generator change |

283
docs/analysis/sram-mirror.md Normal file
View file

@ -0,0 +1,283 @@
# Layer 6: the SRAM mirror of the cache against published SRAM density, year 0 to 10
5 October 2026 (night), Counter ASIC 2.0 (`docs/plans/counter-asic-2.md`, layer 6), branch `ca2-analysis`. Every figure
below is either cited (paper, vendor document, URL, date) or labelled approximate. Nothing here is a measurement of a
chip. Numbers in this file were computed with the arithmetic shown; the script is in section 10.
Revision 2 (same night): the first draft priced the mirror from bit-cell area times a 0.70 array factor. The
coordinator's chip-economics research (sources below) showed that shipped cache-only dies land at about half that
density once assist circuits, redundancy, TSVs, power and test are in. Every table now carries two columns: the
shipped-product density as the headline and the bit-cell figure as the lower bound. The conclusion did not move; the
cost per die rose 2 to 3x.
## 1. The question
The lottery hash derives every dataset item from a 256 MiB cache (spec 01 sections 1.5 and 1.8). A chip that holds the
cache in on-die SRAM can recompute items instead of reading the dataset (ledger M16, the recompute attacker). Layer 6
asks whether the cache size, as the specification schedules it, keeps that SRAM mirror unaffordable for ten years of
the genesis schedule, and if not what growth rule would.
Two things also sit in a chip's SRAM budget if it mirrors the full read-only working set: the layer 5 hot table (32,
64 or 96 MB, a class parameter on `readwidth` b970dda, coordinator's note of 5 October) beside the 256 MiB cache. The
per-warp scratch of layer 3 (32 or 128 KB per warp, written, not read-only) is not mirrorable and is left out of the
mirror; it is counted in the 6 GB working-set budget in section 7.
## 2. What the specification schedules for the cache
| Quantity | Rule | Where |
|---|---|---|
| Dataset | 2 GiB at genesis plus 0.5 GiB per year (`N_d` grows about 23 KiB per day) | spec 01 section 1.13.3, Designed |
| Cache | 256 MiB, "prototype value, to be fixed at gate 1"; the rule that fixes it: "the cache must exceed the largest on-chip cache of any card that mines, and 96 MiB of L2 on the 5090 is the figure to beat" | spec 01 sections 1.5 and 1.16 |
| Cache growth | None. No section of `docs/spec/` grows the cache (grep of `docs/spec` for cache growth, schedule, doubling: only the dataset rule of 1.13.3 and the README's "growth" word, which refers to it) | this analysis, 5 October 2026 |
So the plan's layer 6 row ("already in the design; confirm the schedule") is half right: dataset growth is in the
design, cache growth is not. The cache is flat at 256 MiB for every year of the schedule as the spec stands. M16's
closing line names the rule the cache should get ("exceeds what one die can hold, and grows") as a gate 1 decision
that has not been taken.
## 3. SRAM density, cited: bit cells per node and shipped cache dies
### 3.1 Bit cells
| Node (vendor) | HD 6T bit cell, um^2 | Raw density, Mbit/mm^2 (1/cell) | Year of volume (approximate) | Source |
|---|---|---|---|---|
| N7 (TSMC) | 0.027 | 37.0 | 2018 | WikiChip, "TSMC Details 5 nm" (ISSCC/IEDM disclosures), https://fuse.wikichip.org/news/3398/tsmc-details-5-nm/ |
| N5 (TSMC) | 0.021 | 47.6 | 2020 | same (two N5 cells: HD 0.021, HP 0.025) |
| N3B (TSMC) | 0.0199 | 50.3 | 2022 to 2023 | WikiChip, "IEDM 2022: Did We Just Witness The Death Of SRAM?", https://fuse.wikichip.org/news/7343/iedm-2022-did-we-just-witness-the-death-of-sram/ (TSMC's IEDM 2022 N3 paper) |
| N3E (TSMC) | 0.021 | 47.6 | 2023 | same; Tom's Hardware, "TSMC's 3nm Node: No SRAM Scaling", https://www.tomshardware.com/news/no-sram-scaling-implies-on-more-expensive-cpus-and-gpus |
| N2 (TSMC) | 0.0175 | 57.1 | 2025 to 2026 | TSMC at IEDM 2024, reported by Tom's Hardware, https://www.tomshardware.com/tech-industry/tsmc-shares-deep-dive-details-about-its-cutting-edge-2nm-process-node-at-iedm-2024-35-percent-less-power-or-15-percent-more-performance ; ISSCC 2025 paper "A 38.1Mb/mm2 SRAM in a 2nm-CMOS-Nanosheet Technology", https://research.tsmc.com/page/memory/4.html |
| Intel 18A | 0.021 | 47.6 | 2025 to 2026 | ISSCC 2025 paper 29.2, "A 0.021 um^2 High-Density SRAM in Intel 18A RibbonFET Technology with PowerVia", https://www.researchgate.net/publication/389644177 ; IEEE Spectrum 26 Feb 2025, https://spectrum.ieee.org/sram-intel-tsmc |
| Samsung SF3 / SF2 | not disclosed as a bit cell area in anything found tonight (Samsung's ISSCC papers give assist circuits and macro figures, not the HD cell) | | | search of ISSCC 2021 to 2025 coverage, 5 October 2026; left out of the tables |
The stall. N3B's cell is 5% smaller than N5's and N3E's is the same size as N5's (0.021 um^2 both): zero SRAM
scaling from N5 to N3E (WikiChip IEDM 2022 article above; Tom's Hardware above; SemiAnalysis "TSMC's 3nm Conundrum",
https://newsletter.semianalysis.com/p/tsmcs-3nm-conundrum-does-it-even). N2's nanosheet cell recovers 17% (0.021 to
0.0175 um^2). So across 2020 to 2026 the HD bit cell shrank once, by 17%.
Macro density from the bit cell. WikiChip's and SemiAnalysis's convention is bit-cell density times about 0.70 for
the assist and periphery overhead (SemiAnalysis, December 2022: TSMC N5 HD SRAM macro 31.8 Mib/mm^2 after about 30%
assist overhead; WikiChip's 31.8 Mib/mm^2 for the 0.021 um^2 cell is the same arithmetic). The two ISSCC 2025 macros
bracket it: TSMC N2 38.1 Mb/mm^2 at a 0.0175 um^2 cell is 67%; Intel 18A 38.1 Mb/mm^2 array density and 34.3 Mb/mm^2
for the volume macro at a 0.021 um^2 cell are 80% and 72%. That is a macro on a test chip. It is the LOWER BOUND on
die area, not the die.
### 3.2 Shipped cache dies (what a whole die of SRAM really holds)
| Product | SRAM | Die | Node | MB per mm^2 | Source |
|---|---|---|---|---|---|
| AMD 3D V-Cache (Zen 3 SRAM chiplet) | 64 MB | 41 mm^2 | TSMC 7 nm | 1.56 | AMD at Hot Chips 33, reported by Tom's Hardware, August 2021, https://www.tomshardware.com/news/amd-unveils-more-ryzen-3d-packaging-and-v-cache-details-at-hot-chips ("the 3D V-Cache SRAM measures 41 mm^2", "64 MB of 7 nm SRAM"); the densest cache-only die that has shipped |
| Graphcore GC200 (with compute) | 900 MB | 823 mm^2 | 7 nm | 1.09 | coordinator's chip-economics research, 5 October 2026 (vendor figures) |
| Groq TSP | 220 MB | 725 mm^2 | 14 nm | 0.30 | same |
The V-Cache die is a pure SRAM die with its TSVs, redundancy, test and power: 1.56 MB/mm^2 at N7 against the bit-cell
figure 37.0 Mbit/mm^2 = 4.6 MB/mm^2 and the 0.70-macro figure 3.2 MB/mm^2. The shipped die is 0.48 of the macro
figure. The headline column below scales the V-Cache density to other nodes by the bit-cell ratio (0.027 / cell), an
approximation that assumes the periphery and TSV overheads scale with the cell, which they do not fully (so the
headline column is itself slightly optimistic for the attacker at N5 and below).
### 3.3 GPU on-die SRAM, the reticle, wafer prices
GPU on-die SRAM for scale: the RTX 5090 carries 96 MB of L2 (98,304 KB) on a 750 mm^2 TSMC 4N die with 92.2 billion
transistors; the full GB202 has 128 MB; the RTX 4090 had 72 MB and the RTX 3090 6 MB (NVIDIA, "RTX Blackwell GPU
Architecture" whitepaper v1.1, appendix table "L2 Cache Size", https://images.nvidia.com/aem-dam/Solutions/geforce/blackwell/nvidia-rtx-blackwell-gpu-architecture.pdf).
At the V-Cache density scaled to N5 (2.0 MB/mm^2) that L2 is about 48 mm^2 of the 750 (6%), approximate. The
RX 9070 XT carries 64 MB of Infinity Cache plus 8 MB of L2 (vendor figures, approximate, bench-log "the 9070 XT on the
eGPU").
Reticle: the EUV field is 26 x 33 mm = 858 mm^2, about 830 mm^2 usable after scribe lanes (SemiAnalysis, "Die Size
And Reticle Conundrum", https://newsletter.semianalysis.com/p/die-size-and-reticle-conundrum-cost ; WikiChip "Mask",
https://en.wikichip.org/wiki/mask). The 5090's 750 mm^2 is 90% of it.
Wafer prices (approximate; TSMC publishes none, every figure is supply-chain reporting): N7 about $9,500, N5 and N3
about $20,000 (Silicon Analysts, "Wafer Pricing by Node", September 2026, https://siliconanalysts.com/data/wafer-pricing);
N2 about $30,000 (Tom's Hardware, https://www.tomshardware.com/tech-industry/semiconductors/tsmc-could-charge-up-to-usd45-000-for-1-6nm-wafers-rumors-allege-a-50-percent-increase-in-pricing-over-prior-gen-wafers).
## 4. Die area to mirror the cache, per node, two columns
Headline = V-Cache density (41 mm^2 per 64 MiB at N7) scaled by the bit-cell ratio. Lower bound = bits / (raw
density x 0.70). Columns: the 256 MiB cache alone, the cache plus the 96 MB hot table of layer 5 (as MiB), and the
larger caches of the options in section 7. Area in mm^2; a figure over 830 is split into the dies shown.
| Node | 256 MiB, headline | 256 MiB, lower bound | 256 + 96, headline | 256 + 96, lower bound | 512 MiB, headline / lower | 1 GiB, headline / lower | 4 GiB, headline / lower |
|---|---|---|---|---|---|---|---|
| N7 | 164 | 83 | 226 | 114 | 328 / 166 | 656 / 331 | 2,624 (4 dies) / 1,325 (2 dies) |
| N5 | 128 | 64 | 175 | 89 | 255 / 129 | 510 / 258 | 2,041 (3 dies) / 1,031 (2 dies) |
| N3B | 121 | 61 | 166 | 84 | 242 / 122 | 483 / 244 | 1,934 (3 dies) / 977 (2 dies) |
| N3E, Intel 18A | 128 | 64 | 175 | 89 | 255 / 129 | 510 / 258 | 2,041 (3 dies) / 1,031 (2 dies) |
| N2 | 106 | 54 | 146 | 74 | 213 / 107 | 425 / 215 | 1,701 (3 dies) / 859 (2 dies) |
One reticle (830 mm^2) holds, at the headline density, 1.3 GiB of SRAM at N7, 1.6 GiB at N5, N3E and 18A, 1.9 GiB at
N2 (lower-bound column: 2.5, 3.2, 3.9 GiB).
Against the figures the ledger carries: M16's "100 to 300 mm^2" (low end from a 0.02 um^2 cell with overhead, high
end from wafer-scale parts at about 1 MB per mm^2) brackets the headline 106 to 164 mm^2 well; the plan's "about
45 mm^2 at a leading node" is below even the lower bound and should be read as the bit-cell area with no overhead.
The right figures for the ledger are 106 to 164 mm^2 (shipped density) with 54 to 83 mm^2 as the floor.
## 5. Cost per good die, two columns
Dies per 300 mm wafer by the usual approximation pi x 150^2 / A minus the edge term pi x 300 / sqrt(2A); yield by
Poisson exp(-A x D0) with D0 = 0.1 defects per cm^2 (an assumption, approximate; SRAM arrays carry redundancy so
real yield is higher, which lowers these costs). Cost per good die = wafer price / (dies x yield). Packaging, test,
the logic beside the SRAM and the design (masks at N5 and below run into the tens of millions of dollars,
approximate) are not in these numbers; they are per-die silicon only. Headline / lower bound in each cell.
| Node, wafer price | 256 MiB | 256 + 96 MiB | 1 GiB | 4 GiB |
|---|---|---|---|---|
| N7, $9,500 | 164 mm^2, 379 dies, yield 0.85: $30 / $13 | $44 / $19 | $224 / $75 | $896 (4 dies) / $456 (2 dies) |
| N5, $20,000 | 128 mm^2, 495 dies, 0.88: $46 / $21 | $68 / $30 | $306 / $111 | $1,512 (3 dies) / $621 (2 dies) |
| N3B, $20,000 | 121 mm^2, 524 dies, 0.89: $43 / $20 | $63 / $28 | $280 / $103 | $1,371 (3 dies) / $569 (2 dies) |
| N3E, 18A, $20,000 | $46 / $21 | $68 / $30 | $306 / $111 | $1,512 / $621 |
| N2, $30,000 | 106 mm^2, 600 dies, 0.90: $56 / $26 | $81 / $37 | $343 / $131 | $1,641 (3 dies) / $696 (2 dies) |
Reading. The silicon for a 256 MiB mirror is $30 to $56 per die at shipped density (2 to 3x the first draft's
figure), under $90 with the hot table. A funded chip programme pays that without noticing: it was never the SRAM
that priced the recompute attacker out, and the plan's premise for layer 6 ("the SRAM mirror stays unaffordable")
does not hold for the cache as a mirror and did not hold at genesis either. A 1 GiB cache is a 425 to 656 mm^2 die
($224 to $343), affordable too; 4 GiB is a 3 to 4 die part at about $900 to $1,600 of silicon, which is a different
product but not an impossible one (the attacker's problem at that size is the 1,024 dependent cross-die reads per
hash, section 6).
## 6. What the mirror buys the attacker, year by year
From M16 (`docs/analysis/m16-recompute-attacker-2026-10-05.md`): with the cache on die the attacker recomputes 128
items per hash at about 1,170 integer operations and 8 dependent 64-byte cache reads each, about 150,000 operations
and 1,024 dependent SRAM reads per hash. At a 5090-class integer budget (about 50 T op/s, approximate) that is
0.33 Ghash/s against the honest 141 Mhash/s projected for version 2 programs: 2.4x at equal silicon before any
fixed-function factor, 3x to 6x with one (approximate). The SRAM is 106 to 164 mm^2 of that chip at the headline
density (14 to 22% of a 750 mm^2 die; the m16 model's 13 to 40% band holds), so the mirror is cheap and the recompute
route is bound by integer throughput, not by SRAM.
The layer 5 hot table changes nothing in that arithmetic: the hot table is read-only and derived from the day key
like the cache, so a chip mirrors it in the same SRAM (another 32 to 96 MB, 24 to 48 mm^2 at N5 headline) and reads
it at SRAM latency, which is exactly what a GPU's L2 does with it. Layer 5 taxes the DRAM-only chip (the one without
SRAM); it does not tax the SRAM chip.
Dataset growth does not touch the recompute attacker: the attacker never holds the dataset. It taxes the
partial-store attacker (O-1.6, the time-memory curve, not drawn) and the honest card.
Year by year under the schedule as it stands (flat 256 MiB), the mirror's area at the best node available that
year, headline density. Node years are approximate; the density trend from 2018 to 2025 is 37.0 to 57.1 Mbit/mm^2
raw, 1.54x in 7 years, about 6% per year, and it came in one step (N2); the extrapolation past 2026 assumes that
average holds (approximate, and optimistic for the attacker: A16 and A14 have no disclosed SRAM cell yet).
| Year | Calendar (approximate) | Dataset, GiB | Cache (spec) | Best node | Mirror of the cache, headline (lower bound), mm^2 | With a 96 MiB hot table, headline, mm^2 | Mirror as a share of a 750 mm^2 die |
|---|---|---|---|---|---|---|---|
| 0 | 2027 | 2.0 | 256 MiB | N2 (cited) | 106 (54) | 146 | 14% |
| 1 | 2028 | 2.5 | 256 MiB | N2 or A16 | 103 (52) | 142 | 14% |
| 2 | 2029 | 3.0 | 256 MiB | trend | 95 (48) | 130 | 13% |
| 3 | 2030 | 3.5 | 256 MiB | trend | 89 (45) | 123 | 12% |
| 4 | 2031 | 4.0 | 256 MiB | trend | 84 (43) | 116 | 11% |
| 5 | 2032 | 4.5 | 256 MiB | trend | 79 (40) | 109 | 11% |
| 6 | 2033 | 5.0 | 256 MiB | trend | 75 (38) | 103 | 10% |
| 7 | 2034 | 5.5 | 256 MiB | trend | 71 (36) | 97 | 9% |
| 8 | 2035 | 6.0 | 256 MiB | trend | 67 (34) | 92 | 9% |
| 9 | 2036 | 6.5 | 256 MiB | trend | 63 (32) | 87 | 8% |
| 10 | 2037 | 7.0 | 256 MiB | trend | 59 (30) | 82 | 8% |
Reading. A flat cache's mirror shrinks from 14% to 8% of a large die over the decade, and a 5090-class consumer GPU
already carries 96 MB of L2 on one die with the full GB202 at 128 MB; at the 2020 to 2025 pace of GPU L2 growth
(6 MB, 72 MB, 96 MB on the three NVIDIA flagships in the whitepaper table) a consumer GPU could hold 256 MiB on die
within the decade. The spec's own rule for the cache ("must exceed the largest on-chip cache of any card that
mines") would then be broken by a flat cache. That is the real reason to grow it: not to price a chip out (section
5 shows the SRAM cannot do that) but to keep the cache out of every GPU's own cache, so the honest hash stays
DRAM-latency-bound and the recompute route stays a route only a custom chip can take.
## 7. Answer to the layer 6 question, and the options
Does the flat 256 MiB cache keep the SRAM mirror unaffordable through year 10? No. It is affordable at year 0 ($30 to
$56 of silicon per die at shipped density, section 5) and gets cheaper. What keeps the recompute attacker near 1x is
M16's integer arithmetic and the mixer-cost lever (4x the mixer cost puts the equal-silicon gain at 0.36x, bounded
by the CPU verify gate), not the cache size. The cache size does one other job, keeping the cache larger than any
GPU's L2, and that job needs growth.
Options for the cache rule, with the honest costs each implies. Verifier fill time is 0.2 s per 256 MiB on one core
(spec 1.12: "a 0.2 s CPU cache fill", from the measured 175 to 190 ms of section 1.8.3), scaled linearly; the
verifier holds the whole cache (section 1.11), so its memory is the cache size plus the program and the interpreter.
GPU fill: 0.67 ms per 256 MiB on the 5090 (section 1.8.3), linear. The GPU dataset build (13.4 ms per 1 GiB on the
5090, section 1.8.3) depends on the dataset size, not the cache size; a larger cache spreads the build's 8 dependent
reads per item over more memory, which on a GPU means more of them miss L2 and the build slows by some factor
between 1x and the L2-to-DRAM latency ratio, which is a measurement to take (approximate; owed). Mirror area is at N2
headline density (lower bound in brackets), the node of the first years; at the trend's year-10 density divide by
about 1.8.
| Option | Rule | Cache at year 0 / 4 / 10 | Mirror at N2, headline (lower bound), year 0 / 4 / 10, mm^2 | Dies at year 10 (830 mm^2 reticle), headline | Verifier fill, one core, year 0 / 10 | Verifier memory, year 10 | GPU cache fill (5090), year 10 | Keeps the cache above a 96 MB L2 at year 10 | Keeps it above a 256 MB L2 |
|---|---|---|---|---|---|---|---|---|---|
| A, as specified | flat 256 MiB | 256 / 256 / 256 MiB | 106 (54) / 106 / 106 | 1 | 0.2 / 0.2 s | 256 MiB | 0.7 ms | yes, 2.7x | no |
| B | cache = dataset / 8 (today's ratio) | 256 / 512 / 896 MiB | 106 (54) / 213 (107) / 372 (188) | 1 | 0.2 / 0.7 s | 896 MiB | 2.3 ms | yes, 9.3x | yes, 3.5x |
| C | cache doubles when the dataset doubles (the dataset's own clock: year 4, then year 12) | 256 / 512 / 512 MiB | 106 (54) / 213 (107) / 213 (107) | 1 | 0.2 / 0.4 s | 512 MiB | 1.3 ms | yes, 5.3x | yes, 2x |
| D | cache = dataset / 4 | 512 / 1,024 / 1,792 MiB | 213 (107) / 425 (215) / 744 (376) | 1 | 0.4 / 1.4 s | 1.75 GiB | 4.7 ms | yes | yes, 7x |
| E, one reticle | cache sized so the mirror exceeds one reticle at the node of the day: 2 GiB at N2 headline density (section 4; 4 GiB on the lower bound), growing with density | 2 GiB / about 2.3 / about 3.5 GiB | 850 / 850 / 850 (by construction) | 2 | 1.6 / 2.8 s | 3.5 GiB | 5.4 / 9.4 ms | yes | yes |
Where the working set enters (coordinator's budget: 1 GiB table + hot table + scratch for every resident warp +
buffers under 6 GB on an 8 GB card): the cache is not in the miner's working set at hash time (the dataset is built
from it once a day and the cache can be dropped or kept), so options A to D do not move that budget; the dataset's own
growth does (2 GiB at genesis, 4 GiB at year 4, 7 GiB at year 10, which is past an 8 GB card at about year 8 on its
own). Option E's 2 GiB cache would have to be built on the card and dropped, which is fine for a 16 GB card and tight
on an 8 GB one at build time (2 GiB cache + 2 GiB dataset + hot table). The per-warp scratch at 170 SMs x 64 warps
(approximate, readwidth) is 340 MB at 32 KB and 1.36 GB at 128 KB per warp; with the 1 GiB table, a 96 MB hot table
and buffers that is 1.5 to 2.5 GB at the prototype dataset size, 2.5 to 3.5 GB at the 2 GiB genesis size, inside
6 GB either way.
Recommendation. Option C (the cache doubles when the dataset doubles) is the one that keeps the spec's own rule true
with the smallest verifier cost: it ties the cache to a clock the spec already has, keeps `AND MASK` (a power of two
every step, which is the 1.13.3 option (b) argument again), costs the verifier 0.4 s and 512 MiB at year 4 and nothing
more until year 12, and keeps the cache 2x above a 256 MB GPU L2 if one appears. It does not price a chip out; nothing
about cache size does (section 5). The lever that does is the mixer cost multiplier of M16, which is the gate 1
decision to take beside this one. Option B is the same idea in a smooth form and costs the verifier 0.7 s at year 10.
Option E is the only one that makes the mirror a multi-die part and it costs every verifier 1.6 s and 2 GiB at
genesis (at the headline density; the lower-bound density would ask for 4 GiB and 3.2 s), which fails the spirit of
the 10 ms verify gate (the fill is once a day, but a light node joining pays it on every day it syncs across).
Decision for the project lead, at gate 1: A, B, C, D or E above, together with M16's mixer multiplier. Nothing here changes a
vector today: the cache size is a prototype value of spec 1.16 and the growth rule would be a new sentence in 1.13.3.
## 8. Why the latency bound is the property to lean on (citations behind the plan's rule)
The plan's "what stays true" paragraph says DRAM latency is the same physics for everyone and bandwidth per watt is
what a custom memory chip buys. The sources behind that:
| Claim | Figure | Source |
|---|---|---|
| Random-access DRAM latency is the same across memory types | Row cycle time 40 to 48 ns across DDR4, GDDR5 and HBM2 | Li, Reddy and Jacob, "A Performance and Power Comparison of Contemporary DRAM Architectures", MEMSYS 2018 (coordinator's chip-economics research, 5 October 2026) |
| Latency does not scale, bandwidth does | DRAM latency improved about 1.3x in two decades while bandwidth improved about 20x | K. Chang, "Understanding and Improving the Latency of DRAM-Based Memory Systems", PhD thesis, CMU, 2017 (same research) |
| No mining chip has bought latency with exotic memory | No shipped mining chip has used HBM or stacked memory; the Ethash chips used DDR3, GDDR6 and undisclosed types | same research; the Ethash chip gain of about 3x in the plan came from bandwidth per watt, not latency |
| The honest hash is latency-bound on every card measured | The hash runs within a few percent of 1/128 of each card's dependent random-read ceiling (5090, 9070 XT, M5 Max) | `docs/bench-log.md`, "the 9070 XT on the eGPU", 5 October 2026 (measured) |
Reading for layer 6: an SRAM mirror beats DRAM latency by about 10x per read (a 64 MiB buffer inside the 9070 XT's
Infinity Cache chased at 9.2 G loads/s against 2.5 in GDDR6, the same bench-log entry; the 5090's L2 at 5.8x the
hash rate of its 1 GiB dataset, M16), which is why the recompute attacker is bound by the 1,024 dependent SRAM reads
and the 150,000 integer operations per hash and not by the SRAM's size or price. The cache size decides whether the
mirror is one die or several (section 4); it does not decide whether the mirror exists.
## 9. What is cited, what is approximate, what is owed
| Item | Status |
|---|---|
| Bit cells for N7, N5, N3B, N3E, N2, Intel 18A | cited (section 3.1) |
| Shipped cache-die density (AMD V-Cache 64 MB on 41 mm^2 at 7 nm; Graphcore GC200; Groq TSP) | cited (section 3.2; V-Cache checked against Tom's Hardware's Hot Chips 33 report, 5 October 2026; the Graphcore and Groq rows are from the coordinator's research and were not re-checked tonight) |
| Scaling the V-Cache density to other nodes by the bit-cell ratio | approximate, stated |
| Samsung SF2 or SF3 bit cell | not found; left out |
| Array efficiency 0.70 | WikiChip's and SemiAnalysis's convention, bracketed by two ISSCC 2025 macros (67 to 80%); a macro figure, used only as the lower bound |
| Wafer prices | approximate, supply-chain reporting, cited |
| D0 = 0.1 per cm^2, Poisson yield | assumption, stated |
| Node years and the 6% per year density trend past 2026 | approximate, extrapolated from cited 2018 to 2025 points |
| GPU L2 sizes | cited (NVIDIA whitepaper); AMD Infinity Cache approximate |
| Latency citations (MEMSYS 2018, Chang 2017, mining-chip memory types) | from the coordinator's research, not re-read tonight |
| Recompute attacker arithmetic | M16, which is itself arithmetic on measured rates, not a chip measurement |
| Dataset-build slowdown at a larger cache on a GPU | owed, a measurement (5090 at a 512 MiB and 1 GiB cache) |
| The on-die emulation of M16 (inline kernel with a 64 MiB cache inside the 5090's L2) | still a PC job (M16) |
## 10. The arithmetic
```
MiB = 2^20; bits = cache_MiB * MiB * 8
headline_mm2 = cache_MiB * (41 / 64) * (cell_um2 / 0.027) (V-Cache: 41 mm2 per 64 MiB at N7, scaled by cell)
raw_Mbit_per_mm2 = 1 / cell_um2 (1e6 cells per mm2 per um2 of cell)
lower_bound_mm2 = bits / (raw * 0.70 * 1e6)
dies_per_wafer = pi * 150^2 / area - pi * 300 / sqrt(2 * area)
yield = exp(-area_mm2 * 0.001) (D0 = 0.1 per cm2)
cost_per_good_die = wafer_price / (dies * yield); over 830 mm2: k = ceil(area / 830) dies of area / k, cost x k
reticle_GiB = 830 / (mm2 per MiB) / 1024
```
Run on 5 October 2026 with Python 3 on the M5 Max; the printed tables are the ones above, rounded.

26
docs/bench-log.md
View file

@ -1883,3 +1883,29 @@ Reading, with the miner-on pairs above (empty shard 15,670 MiB, full prototype s
|---|---|
| Unit tests | `cargo test --release -p kaspa-consensus-core -p igneum-exec --lib -- proving config::params::tests::override_params_carry_the_proving_v1 config::params::tests::consensus_digest` on this Mac (target `vendor/igneum-node/target-pv1`, 19:09Z): consensus core 13 passed (the segment record round trip, signature and the three nested sections; the credit split; the params switch and the digest that moves only once the switch is set), exec 8 passed (the segment grid and the split; the record checks: alignment, block, chain length, the veto naming the field, the deadline, the window; the chain rule both ways; the unproven restart; the shard side at 90%; the pool offering the segment section). The six full node suites go to PC 2 as a build job when the fleet is back |
| The fast-time 3-node harness (`tools/proving-v1/net.mjs`, 29950+, suffix 956, every node in trust mode, three vmine voters, v0 at DAA 60, v1 at DAA 120, 4 blocks a segment, unproven after 60 DAA, a tenth to the aggregator; fork b177718e built on this Mac) | run 2, 19:13:01Z to 19:16:19Z, under the run lock: PASSED, 21 checks in 197.3 s (`tools/proving-v1/report-2026-10-05.json`). v1 start = chain block 119 on all three nodes; the native statement identical on all three. Known-finished: segment 119..122's fresh-chain record submitted to n1 at t=131.1 s, relayed, verified (trust) and PAID on n0 1.0 s later at chain block 129, 253,611,648,000,000,000 wei = a tenth of the four credits, the same on every node, the payout address holding it. Chain rule: segment 123..126's fresh-chain record refused ("does not chain to segment 119..122 ... proven (record paid at chain block 129)"), the continuing one (chain_len 8) accepted and paid. Known-failed: segment 127..130 left without a record: a fresh-chain record for 131..134 refused while 127..130 was pending ("pending until DAA 191"); at DAA 192 the status read unproven, a late record for 127..130 refused ("unproven: carried after the deadline"), the fresh-chain record for 131..134 accepted and paid with chain_len 4; `segmentsInWindow` proven 3, unproven 1. The shard side: a v1 shard's `shardWei` = 90% of its block's credit. Run 1 (19:10Z) failed in its own tooling (the signer's argument order), fixed. Run 3 on the FINAL fork tree (ece42979 on the 0.3.10 commit 21d4c73c, protocol 15, N = 8 both in the params default and `--segment 8`, the fast-time file's four fields), 20:52:41Z to 20:56:45Z: PASSED, 21 checks in 244.4 s (segments of 8: 119..126 paid in 1.0 s after submission, 127..134 refused fresh and paid continuing with chain_len 16, 135..142 left unproven and skipped, 143..150 restarted the chain) |
## 5 October 2026 (night), dp4a-class throughput on the M5 Max: the dot4 emulation against the ALU chain (Counter ASIC 2.0 layer 7)
Apple M5 Max, macOS 26, branch `ca2-analysis` (base `readwidth` 4badcee). The probes are standalone (no pack, no lottery kernel): `proto-metal/dot4-probe.swift` (built `swiftc -O -o dot4-probe dot4-probe.swift -framework Metal` under `with-lock.sh build`), `proto-opencl/dot4-probe.c` (built `cc -std=c99 -O2 -o dot4-probe-cl dot4-probe.c -framework OpenCL`), both run under `with-lock.sh measure` (exclusive; nothing else built or measured on the Mac during the runs). Shape: a dependent chain of one dot4 per step per lane, `acc = dot4(x, y, acc); x = x * 0x9E3779B1 + acc; y = rotl(y, 7) ^ (acc + s)`, 1,048,576 lanes x 4,096 steps, work-group 256, best of 3 with a fresh seed per repetition, device time (Metal: command buffer GPU start to end; OpenCL: event profiling). Beside it the ALU chain of the 9070 XT entry (`x = x * K + rotl(y, 7); y = (y ^ x) + s`, 5 ops per step counted). Every kernel is checked bit for bit against a CPU reference on lanes 0 and 1,048,575 in every repetition ("ok"). Design context: `docs/analysis/int8-matrix-family.md`.
| API, kernel | What one step is | best ms | G steps/s | ns per dependent step | ok |
|---|---|---|---|---|---|
| Metal, `probe_alu` | mul, add, rotate, xor, add | 4.882 | 879.8 (about 4.4 T int ops/s at 5 per step, approximate) | 1,192 | yes |
| Metal, `probe_dot4s` | signed dot4 emulated: `int4(as_type<char4>(a))` x same for b, 4 products summed into a wrapping int, plus the 3-op chain | 22.820 | 188.2 G dot4/s | 5,571 | yes |
| Metal, `probe_dot4u` | unsigned dot4 emulated: `uint4(as_type<uchar4>(a))`, same chain | 7.834 | 548.2 G dot4/s | 1,913 | yes |
| Apple OpenCL 1.2, `alu` | as Metal | 4.928 | 871.5 | 1,203 | yes |
| Apple OpenCL 1.2, `dot4e` | signed dot4 emulated with `convert_int4(as_char4(a))` | 22.797 | 188.4 G dot4/s | 5,566 | yes |
| Apple OpenCL 1.2, `dot4_khr` | `acc + dot(as_char4(x), as_char4(y))` under `#pragma OPENCL EXTENSION cl_khr_integer_dot_product : enable` | 5.076 | 846.2 | 1,239 | NO: mismatched the CPU reference on every lane checked in all 3 repetitions |
Reading: on this GPU a signed-byte dot4 costs 4.7 ALU-chain steps and an unsigned-byte one 1.6; Metal has no dp4a and no integer simdgroup matrix (MSL 4.1 sections 2.4 and 6.9), so these are the honest Apple costs of a per-lane dot4 family, and an unsigned definition is 3x cheaper for Apple at no cost to NVIDIA or AMD (both carry the unsigned form, PTX `dp4a.u32.u32`, AMD `v_dot4_u32_u8`). Apple's OpenCL does not list `cl_khr_integer_dot_product`; its `dot` on `char4` compiled anyway and returned something other than the integer dot (the mismatch), which is why a family's conformance vectors must gate every vendor path on the feature macro, not on "it compiled". Not run here: NVIDIA and AMD. The PC job is prepared and not published (coordinator's rule): `relay/playbooks/dot4-probe.ps1` with `dot4-probe-cl.exe` (proto-opencl/dot4-probe.c cross-compiled with mingw as `x86_64-w64-mingw32-gcc -std=c99 -O2 -static -DIGNEUM_CL_DYNAMIC -DCL_TARGET_OPENCL_VERSION=120 -I proto-cuda/nvrtc/redist/include`, sha256 `5adaeb1aceb03dc41135baabe0b53f1ed5fac891a5b3c3849645b03efe4416f4`, 161,863 bytes); it runs the scalar, KHR, AMD `__builtin_amdgcn_sudot4` and NVIDIA inline-PTX `dp4a` variants on every OpenCL GPU of the machine with the mining cards switched off through `/api/cards` and restored after. The CUDA form (`proto-cuda/dot4-probe.cu`, `__dp4a`) needs nvcc on the PC and is the cross-check.
**PC 1, 5 October 2026 20:29 UTC, the same probe on the RTX 5090 and the RX 9070 XT** (machine ae432dc7, Windows 11; fetch job `fetch-dot4-20261005` placed `dot4-probe-cl.exe` sha256 `5adaeb1a…6416f4`, run job `run-dot4-20261005` ran `relay/playbooks/dot4-probe.ps1`: the app's `nvidia:0` and `amd:1:gfx1201` cards switched off through `POST api/cards`, the probe run on every OpenCL device, the cards restored with their settings (identities 8 and 2, power cap 80% and none); `node tools/jobs.mjs run-dot4-20261005`; 101 s wall, every kernel under 10 ms; device event time, best of 3, same lanes and steps as the Mac rows):
| Device, platform | alu, G steps/s (ms) | dot4e signed emulation, G dot4/s (ms) | dot4 instruction, G dot4/s (ms) | `cl_khr_integer_dot_product` | ok |
|---|---|---|---|---|---|
| RTX 5090, NVIDIA OpenCL 3.0 CUDA, driver 617.14 | 8,753.5 (0.491) | 1,239.1 (3.466), 7.1x the ALU step | 7,453.6 (0.576) via inline PTX `dp4a.s32.s32`, 1.17x the ALU step | not listed; the `dot(char4,char4)` kernel does not build | yes |
| RX 9070 XT (gfx1201), AMD-APP 3683.0 (PAL,LC), OpenCL 2.0 | 701.4 (6.124) | 480.8 (8.932), 1.46x | 664.3 (6.465) via `__builtin_amdgcn_sudot4`, 1.06x | not listed; same | yes |
| RX 9070 XT, the older 3652.0 platform entry (dup) | 696.2 (6.169) | 501.7 (8.561) | 683.6 (6.283) | not listed | yes |
| gfx1036 (integrated RDNA 2, 2 CUs), 3683.0 | 40.6 (105.9) | 15.8 (272.3), 2.6x | `sudot4` does not build: "needs target feature dot8-insts" | not listed | alu and dot4e yes |
Reading: one `dp4a` on the 5090 costs about one ALU-chain step (7.45 T dot4/s, 0.85 of the chain's 8.75 T steps/s); one `v_dot4_i32_iu8` on the 9070 XT the same (0.66 T, 0.95 of its chain). Emulating the signed dot4 costs 6.0x the instruction on NVIDIA (the OpenCL compiler does not fold the four sign-extended products into `dp4a`) and 1.38x on AMD. Vendor ratios: the 5090 is 12.5x the 9070 XT on the ALU chain and 11.2x on hardware dot4; against the M5 Max's best (unsigned emulation, 0.55 T) it is 10x on the chain and 13.6x on dot4. The hash itself is bound by DRAM reads, so these per-op numbers bound a family's cost and are not hash rates (`docs/analysis/int8-matrix-family.md` section 4). Adrenalin's OpenCL C accepts the clang builtin and emits the instruction on RDNA 4 (the third-party RDNA 3 report of the same route is now confirmed on this card); no PC platform lists the Khronos integer-dot extension. The 5090 SM clock read 2,505 MHz before and after (nvidia-smi; 2,850 MHz while mining in the telemetry entry), so the card was idle for the probe.

115
proto-cuda/dot4-probe.cu Normal file
View file

@ -0,0 +1,115 @@
// dot4-probe (CUDA): dp4a-class throughput on NVIDIA, standalone (no pack, no lottery kernel).
// Counter ASIC 2.0 layer 7 (docs/analysis/int8-matrix-family.md), 5 October 2026. PC job; not run on the Mac.
//
// Three dependent chains, same shape as the OpenCL --memprobe ALU chain (proto-opencl/host.c, probe_alu) and the
// Metal probe (proto-metal/dot4-probe.swift): 1,048,576 lanes x 4,096 steps, best of 3, event time.
// alu x = x * K + rotl(y, 7); y = (y ^ x) + s the card's integer baseline, 5 ops per step counted
// dot4i acc = __dp4a(x, y, acc) (PTX dp4a.s32.s32, sm_61+) one hardware dot4 per step per lane
// dot4e the scalar emulation of the same (4 sign-extended byte products summed, wrapping int32)
// The emulation and the intrinsic must agree bit for bit with the CPU reference (checked on two lanes per run).
//
// Build (Windows, CUDA Toolkit): nvcc -O2 -arch=sm_120 -o dot4-probe-cuda.exe dot4-probe.cu
// Build (Linux): nvcc -O2 -arch=sm_120 -o dot4-probe-cuda dot4-probe.cu
// Run: dot4-probe-cuda [--lanes N] [--steps N] [--reps N] [--device N]
#include <cuda_runtime.h>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <cstdint>
__host__ __device__ inline uint32_t pm_mix(uint32_t x) { x ^= x >> 16; x *= 0x7feb352du; x ^= x >> 15; x *= 0x846ca68bu; x ^= x >> 16; return x; }
__host__ __device__ inline uint32_t rotl32(uint32_t v, uint32_t n) { return (v << n) | (v >> (32u - n)); }
__host__ __device__ inline int32_t dot4_emul(uint32_t a, uint32_t b, int32_t acc) {
int32_t r = acc;
for (int i = 0; i < 4; ++i) {
int32_t ba = (int32_t)(int8_t)((a >> (8 * i)) & 0xffu);
int32_t bb = (int32_t)(int8_t)((b >> (8 * i)) & 0xffu);
r = (int32_t)((uint32_t)r + (uint32_t)(ba * bb));
}
return r;
}
__global__ void probe_alu(uint32_t steps, uint32_t seed, uint32_t* out) {
uint32_t g = blockIdx.x * blockDim.x + threadIdx.x;
uint32_t x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u;
for (uint32_t s = 0; s < steps; ++s) { x = x * 0x9E3779B1u + rotl32(y, 7u); y = (y ^ x) + s; }
out[g] = x ^ y;
}
__global__ void probe_dot4i(uint32_t steps, uint32_t seed, uint32_t* out) {
uint32_t g = blockIdx.x * blockDim.x + threadIdx.x;
uint32_t x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u;
int32_t acc = (int32_t)pm_mix(x);
for (uint32_t s = 0; s < steps; ++s) {
acc = __dp4a((int)x, (int)y, acc);
x = x * 0x9E3779B1u + (uint32_t)acc;
y = rotl32(y, 7u) ^ ((uint32_t)acc + s);
}
out[g] = (uint32_t)acc ^ x ^ y;
}
__global__ void probe_dot4e(uint32_t steps, uint32_t seed, uint32_t* out) {
uint32_t g = blockIdx.x * blockDim.x + threadIdx.x;
uint32_t x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u;
int32_t acc = (int32_t)pm_mix(x);
for (uint32_t s = 0; s < steps; ++s) {
acc = dot4_emul(x, y, acc);
x = x * 0x9E3779B1u + (uint32_t)acc;
y = rotl32(y, 7u) ^ ((uint32_t)acc + s);
}
out[g] = (uint32_t)acc ^ x ^ y;
}
static uint32_t lane_ref(const char* name, uint32_t g, uint32_t seed, uint32_t steps) {
uint32_t x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u;
if (strcmp(name, "alu") == 0) {
for (uint32_t s = 0; s < steps; ++s) { x = x * 0x9E3779B1u + rotl32(y, 7u); y = (y ^ x) + s; }
return x ^ y;
}
int32_t acc = (int32_t)pm_mix(x);
for (uint32_t s = 0; s < steps; ++s) {
acc = dot4_emul(x, y, acc);
x = x * 0x9E3779B1u + (uint32_t)acc;
y = rotl32(y, 7u) ^ ((uint32_t)acc + s);
}
return (uint32_t)acc ^ x ^ y;
}
#define CK(x) do { cudaError_t e = (x); if (e != cudaSuccess) { printf("CUDA error %s at %s:%d\n", cudaGetErrorString(e), __FILE__, __LINE__); return 1; } } while (0)
int main(int argc, char** argv) {
uint32_t lanes = 1u << 20, steps = 4096u; int reps = 3, device = 0;
for (int i = 1; i < argc; ++i) {
if (!strcmp(argv[i], "--lanes") && i + 1 < argc) lanes = (uint32_t)strtoul(argv[++i], 0, 10);
else if (!strcmp(argv[i], "--steps") && i + 1 < argc) steps = (uint32_t)strtoul(argv[++i], 0, 10);
else if (!strcmp(argv[i], "--reps") && i + 1 < argc) reps = atoi(argv[++i]);
else if (!strcmp(argv[i], "--device") && i + 1 < argc) device = atoi(argv[++i]);
else { printf("unknown argument %s\n", argv[i]); return 2; }
}
CK(cudaSetDevice(device));
cudaDeviceProp p; CK(cudaGetDeviceProperties(&p, device));
printf("dot4-probe (CUDA) on %s, sm_%d%d, %d SMs, %d MHz, lanes %u, steps %u, best of %d, event time\n", p.name, p.major, p.minor, p.multiProcessorCount, p.clockRate / 1000, lanes, steps, reps);
uint32_t* d_out; CK(cudaMalloc(&d_out, (size_t)lanes * 4));
uint32_t* h_out = (uint32_t*)malloc((size_t)lanes * 4);
cudaEvent_t e0, e1; CK(cudaEventCreate(&e0)); CK(cudaEventCreate(&e1));
printf("| kernel | lanes | steps | best ms | G steps/s (= G dot4/s for dot4 rows) | ns per dependent step | lanes 0 and last ok |\n|---|---|---|---|---|---|---|\n");
const char* names[3] = { "alu", "dot4i", "dot4e" };
for (int k = 0; k < 3; ++k) {
float best = 1e30f; int ok = 1;
for (int r = 0; r < reps; ++r) {
uint32_t seed = 0x2468aceu + (uint32_t)r * 0x9E3779B9u;
CK(cudaEventRecord(e0));
if (k == 0) probe_alu<<<lanes / 256, 256>>>(steps, seed, d_out);
else if (k == 1) probe_dot4i<<<lanes / 256, 256>>>(steps, seed, d_out);
else probe_dot4e<<<lanes / 256, 256>>>(steps, seed, d_out);
CK(cudaEventRecord(e1)); CK(cudaEventSynchronize(e1)); CK(cudaGetLastError());
float ms = 0; CK(cudaEventElapsedTime(&ms, e0, e1)); if (ms < best) best = ms;
CK(cudaMemcpy(h_out, d_out, (size_t)lanes * 4, cudaMemcpyDeviceToHost));
uint32_t gs[2] = { 0u, lanes - 1u };
for (int j = 0; j < 2; ++j) { uint32_t want = lane_ref(names[k], gs[j], seed, steps); if (h_out[gs[j]] != want) { ok = 0; printf("MISMATCH %s lane %u: gpu %08x cpu %08x\n", names[k], gs[j], h_out[gs[j]], want); } }
}
double sps = (double)lanes * (double)steps / (best / 1000.0);
printf("| %s | %u | %u | %.3f | %.2f | %.3f | %s |\n", names[k], lanes, steps, best, sps / 1e9, best * 1e6 / (double)steps, ok ? "yes" : "NO");
printf("RESULT DOT4 vendor=nvidia device=\"%s\" kernel=%s lanes=%u steps=%u best_ms=%.3f gsteps_per_s=%.2f ok=%d\n", p.name, names[k], lanes, steps, best, sps / 1e9, ok);
}
printf("dot4-probe: done\n");
return 0;
}

165
proto-metal/dot4-probe.swift Normal file
View file

@ -0,0 +1,165 @@
// dot4-probe: dp4a-class throughput on Apple silicon, standalone (no pack, no lottery kernel).
// Counter ASIC 2.0 layer 7 (docs/analysis/int8-matrix-family.md), 5 October 2026.
//
// Metal has no dp4a intrinsic and no integer simdgroup_matrix (MSL 4.1 section 2.4 lists half, bfloat and float only;
// the Metal 4 tensor op matmul2d does carry char x char -> int, MSL 4.1 table 7.3, measured separately when it is).
// So the per-lane dot4 here is the scalar emulation a conforming Apple miner would run: four sign-extended bytes of
// each operand multiplied and summed into a wrapping int32 accumulator, exactly the PTX dp4a semantics
// (PTX ISA 9.4 section 9.7.1.24: d = c; d += Va[i] * Vb[i] for i in 0..3, bytes sign- or zero-extended).
//
// Two kernels, same shape as the OpenCL --memprobe ALU chain (proto-opencl/host.c, probe_alu: 1,048,576 lanes x 4,096
// steps, best of 3):
// alu x = x * K + rotate(y, 7); y = (y ^ x) + s the card's integer baseline, 5 ops per step counted
// dot4 acc = dot4(x, y, acc); x = x * K + acc; y = rotate(y, 7) ^ acc one dependent dot4 per step per lane
// Rates: G steps/s per lane-step, so G dot4/s for the second kernel. Timing is the command buffer's GPU start to end.
//
// Build: swiftc -O -o dot4-probe dot4-probe.swift -framework Metal
// Run: ./dot4-probe [--lanes N] [--steps N] [--reps N] [--signed|--unsigned]
import Foundation
import Metal
let source = """
#include <metal_stdlib>
using namespace metal;
inline uint pm_mix(uint x) { x ^= x >> 16; x *= 0x7feb352du; x ^= x >> 15; x *= 0x846ca68bu; x ^= x >> 16; return x; }
// dp4a, signed bytes, wrapping int32 accumulate: the exact PTX dp4a.s32.s32 semantics.
inline int dot4_s(uint a, uint b, int acc) {
int4 va = int4(as_type<char4>(a));
int4 vb = int4(as_type<char4>(b));
return acc + va.x * vb.x + va.y * vb.y + va.z * vb.z + va.w * vb.w;
}
// dp4a, unsigned bytes, wrapping uint32 accumulate: dp4a.u32.u32.
inline uint dot4_u(uint a, uint b, uint acc) {
uint4 va = uint4(as_type<uchar4>(a));
uint4 vb = uint4(as_type<uchar4>(b));
return acc + va.x * vb.x + va.y * vb.y + va.z * vb.z + va.w * vb.w;
}
kernel void probe_alu(constant uint& steps [[buffer(0)]], constant uint& seed [[buffer(1)]],
device uint* out [[buffer(2)]], uint g [[thread_position_in_grid]]) {
uint x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u;
for (uint s = 0u; s < steps; ++s) { x = x * 0x9E3779B1u + rotate(y, 7u); y = (y ^ x) + s; }
out[g] = x ^ y;
}
kernel void probe_dot4s(constant uint& steps [[buffer(0)]], constant uint& seed [[buffer(1)]],
device uint* out [[buffer(2)]], uint g [[thread_position_in_grid]]) {
uint x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u;
int acc = int(pm_mix(x));
for (uint s = 0u; s < steps; ++s) {
acc = dot4_s(x, y, acc);
x = x * 0x9E3779B1u + uint(acc);
y = rotate(y, 7u) ^ (uint(acc) + s);
}
out[g] = uint(acc) ^ x ^ y;
}
kernel void probe_dot4u(constant uint& steps [[buffer(0)]], constant uint& seed [[buffer(1)]],
device uint* out [[buffer(2)]], uint g [[thread_position_in_grid]]) {
uint x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u;
uint acc = pm_mix(x);
for (uint s = 0u; s < steps; ++s) {
acc = dot4_u(x, y, acc);
x = x * 0x9E3779B1u + acc;
y = rotate(y, 7u) ^ (acc + s);
}
out[g] = acc ^ x ^ y;
}
"""
// CPU reference of the dot4 chain for one lane, to check the kernel is the arithmetic it claims (bit-exact).
func pmMix(_ v: UInt32) -> UInt32 {
var x = v
x ^= x >> 16; x = x &* 0x7feb352d; x ^= x >> 15; x = x &* 0x846ca68b; x ^= x >> 16
return x
}
func dot4sRef(_ a: UInt32, _ b: UInt32, _ acc: Int32) -> Int32 {
var r = acc
for i in 0..<4 {
let ba = Int32(Int8(truncatingIfNeeded: a >> (8 * UInt32(i))))
let bb = Int32(Int8(truncatingIfNeeded: b >> (8 * UInt32(i))))
r = r &+ ba &* bb
}
return r
}
func dot4uRef(_ a: UInt32, _ b: UInt32, _ acc: UInt32) -> UInt32 {
var r = acc
for i in 0..<4 {
let ba = UInt32(UInt8(truncatingIfNeeded: a >> (8 * UInt32(i))))
let bb = UInt32(UInt8(truncatingIfNeeded: b >> (8 * UInt32(i))))
r = r &+ ba &* bb
}
return r
}
func rotl(_ v: UInt32, _ n: UInt32) -> UInt32 { (v << n) | (v >> (32 - n)) }
func laneRef(kernel: String, g: UInt32, seed: UInt32, steps: UInt32) -> UInt32 {
var x = pmMix(g ^ seed), y = x ^ 0x5bd1e995
switch kernel {
case "probe_alu":
for s in 0..<steps { x = x &* 0x9E3779B1 &+ rotl(y, 7); y = (y ^ x) &+ s }
return x ^ y
case "probe_dot4s":
var acc = Int32(bitPattern: pmMix(x))
for s in 0..<steps { acc = dot4sRef(x, y, acc); x = x &* 0x9E3779B1 &+ UInt32(bitPattern: acc); y = rotl(y, 7) ^ (UInt32(bitPattern: acc) &+ s) }
return UInt32(bitPattern: acc) ^ x ^ y
default:
var acc = pmMix(x)
for s in 0..<steps { acc = dot4uRef(x, y, acc); x = x &* 0x9E3779B1 &+ acc; y = rotl(y, 7) ^ (acc &+ s) }
return acc ^ x ^ y
}
}
var lanes = 1 << 20, steps: UInt32 = 4096, reps = 3
var args = Array(CommandLine.arguments.dropFirst())
while !args.isEmpty {
let a = args.removeFirst()
switch a {
case "--lanes": lanes = Int(args.removeFirst())!
case "--steps": steps = UInt32(args.removeFirst())!
case "--reps": reps = Int(args.removeFirst())!
default: print("unknown argument \(a)"); exit(2)
}
}
guard let dev = MTLCreateSystemDefaultDevice() else { print("no Metal device"); exit(1) }
let lib: MTLLibrary
do { lib = try dev.makeLibrary(source: source, options: nil) } catch { print("compile failed: \(error)"); exit(1) }
let queue = dev.makeCommandQueue()!
let outBuf = dev.makeBuffer(length: lanes * 4, options: .storageModeShared)!
print("dot4-probe on \(dev.name), lanes \(lanes), steps \(steps), best of \(reps), GPU start-to-end time")
print("| kernel | lanes | steps | best ms | G steps/s (= G dot4/s for dot4 rows) | ns per dependent step | lane 0 ok |")
print("|---|---|---|---|---|---|---|")
for name in ["probe_alu", "probe_dot4s", "probe_dot4u"] {
let fn = lib.makeFunction(name: name)!
let pso = try! dev.makeComputePipelineState(function: fn)
let tg = min(256, pso.maxTotalThreadsPerThreadgroup)
var best = Double.infinity
var okAll = true
for r in 0..<reps {
var st = steps
var seed = UInt32(0x2468ace) &+ UInt32(r) &* 0x9E3779B9
let cb = queue.makeCommandBuffer()!
let enc = cb.makeComputeCommandEncoder()!
enc.setComputePipelineState(pso)
enc.setBytes(&st, length: 4, index: 0)
enc.setBytes(&seed, length: 4, index: 1)
enc.setBuffer(outBuf, offset: 0, index: 2)
enc.dispatchThreads(MTLSize(width: lanes, height: 1, depth: 1), threadsPerThreadgroup: MTLSize(width: tg, height: 1, depth: 1))
enc.endEncoding()
cb.commit()
cb.waitUntilCompleted()
let ms = (cb.gpuEndTime - cb.gpuStartTime) * 1000.0
if ms < best { best = ms }
// bit-exactness of the kernel against the CPU reference on two lanes
let p = outBuf.contents().bindMemory(to: UInt32.self, capacity: lanes)
for g in [UInt32(0), UInt32(lanes - 1)] {
let want = laneRef(kernel: name, g: g, seed: seed, steps: steps)
if p[Int(g)] != want { okAll = false; print("MISMATCH \(name) lane \(g): gpu \(String(p[Int(g)], radix: 16)) cpu \(String(want, radix: 16))") }
}
}
let stepsPerS = Double(lanes) * Double(steps) / (best / 1000.0)
print(String(format: "| %@ | %d | %u | %.3f | %.2f | %.3f | %@ |", name, lanes, steps, best, stepsPerS / 1e9, best * 1e6 / Double(steps), okAll ? "yes" : "NO"))
}
print("dot4-probe: done")

187
proto-opencl/dot4-probe.c Normal file
View file

@ -0,0 +1,187 @@
/* dot4-probe (OpenCL): dp4a-class throughput on AMD and NVIDIA through OpenCL, standalone (no pack, no lottery kernel).
* Counter ASIC 2.0 layer 7 (docs/analysis/int8-matrix-family.md), 5 October 2026. PC job; the Mac can run it on Apple
* OpenCL for the scalar rows only.
*
* Dependent chains, same shape as --memprobe's ALU chain (host.c, probe_alu) and the Metal and CUDA probes:
* 1,048,576 lanes x 4,096 steps, best of 3, device event time.
* alu x = x * K + rotate(y, 7); y = (y ^ x) + s the card's integer baseline
* dot4e scalar emulation: 4 sign-extended byte products summed into a wrapping int accumulator
* dot4_amd __builtin_amdgcn_sudot4(a, 1, b, 1, acc, 0): V_DOT4_I32_IU8 with both operands signed, no clamp (RDNA 3;
* the same builtin on RDNA 4 is the thing to check: LLVM gates it on the dot7-insts / dot8-insts features)
* dot4_khr cl_khr_integer_dot_product: acc + dot(as_char4(a), as_char4(b)) (OpenCL C 3.0 extension, section
* 6.2.2.16 of the OpenCL C specification; listed by a platform or not)
* dot4_nv inline PTX "dp4a.s32.s32" (NVIDIA's OpenCL compiler accepts inline PTX asm; PTX ISA 9.7.1.24)
* Each variant is built on its own: a variant the platform cannot compile prints one "build failed" row with the first
* line of the build log and the run goes on. Every variant that runs is checked bit for bit against the CPU reference
* on two lanes, so an intrinsic with different semantics (a saturating accumulate, say) shows as NO in the ok column.
*
* Build (Mac, Apple OpenCL): cc -std=c99 -O2 -o dot4-probe-cl dot4-probe.c -framework OpenCL
* Build (Windows, mingw, no SDK): x86_64-w64-mingw32-gcc -std=c99 -O2 -static -DIGNEUM_CL_DYNAMIC -DCL_TARGET_OPENCL_VERSION=120 \
* -I <redist>/include -o dot4-probe-cl.exe dot4-probe.c (OpenCL.dll loaded at run time)
* Run: dot4-probe-cl [--list] [--device N] [--lanes N] [--steps N] [--reps N]
*/
#define CL_TARGET_OPENCL_VERSION 120
#ifdef __APPLE__
#include <OpenCL/cl.h>
#else
#include <CL/cl.h>
#endif
#ifdef IGNEUM_CL_DYNAMIC
#include "cl_dynamic.h"
#endif
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <stdint.h>
static const char* COMMON =
"static inline uint pm_mix(uint x) { x ^= x >> 16; x *= 0x7feb352du; x ^= x >> 15; x *= 0x846ca68bu; x ^= x >> 16; return x; }\n"
"#define CHAIN_HEAD uint g = (uint)get_global_id(0); uint x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u; int acc = (int)pm_mix(x);\n"
"#define CHAIN_TAIL x = x * 0x9E3779B1u + (uint)acc; y = rotate(y, 7u) ^ ((uint)acc + s);\n"
"#define CHAIN_OUT out[g] = (uint)acc ^ x ^ y;\n";
static const char* K_ALU =
"__kernel void probe(uint steps, uint seed, __global uint* out) {\n"
" uint g = (uint)get_global_id(0); uint x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u;\n"
" for (uint s = 0u; s < steps; ++s) { x = x * 0x9E3779B1u + rotate(y, 7u); y = (y ^ x) + s; }\n"
" out[g] = x ^ y;\n"
"}\n";
static const char* K_EMUL =
"static inline int dot4e(uint a, uint b, int acc) {\n"
" int4 va = convert_int4(as_char4(a)); int4 vb = convert_int4(as_char4(b));\n"
" return acc + va.x * vb.x + va.y * vb.y + va.z * vb.z + va.w * vb.w;\n"
"}\n"
"__kernel void probe(uint steps, uint seed, __global uint* out) {\n"
" CHAIN_HEAD\n"
" for (uint s = 0u; s < steps; ++s) { acc = dot4e(x, y, acc); CHAIN_TAIL }\n"
" CHAIN_OUT\n"
"}\n";
static const char* K_AMD =
"__kernel void probe(uint steps, uint seed, __global uint* out) {\n"
" CHAIN_HEAD\n"
" for (uint s = 0u; s < steps; ++s) { acc = __builtin_amdgcn_sudot4(true, (int)x, true, (int)y, acc, false); CHAIN_TAIL }\n"
" CHAIN_OUT\n"
"}\n";
static const char* K_KHR =
"#pragma OPENCL EXTENSION cl_khr_integer_dot_product : enable\n"
"__kernel void probe(uint steps, uint seed, __global uint* out) {\n"
" CHAIN_HEAD\n"
" for (uint s = 0u; s < steps; ++s) { acc = acc + dot(as_char4(x), as_char4(y)); CHAIN_TAIL }\n"
" CHAIN_OUT\n"
"}\n";
static const char* K_NV =
"static inline int dot4nv(uint a, uint b, int acc) { int d; asm(\"dp4a.s32.s32 %0, %1, %2, %3;\" : \"=r\"(d) : \"r\"(a), \"r\"(b), \"r\"(acc)); return d; }\n"
"__kernel void probe(uint steps, uint seed, __global uint* out) {\n"
" CHAIN_HEAD\n"
" for (uint s = 0u; s < steps; ++s) { acc = dot4nv(x, y, acc); CHAIN_TAIL }\n"
" CHAIN_OUT\n"
"}\n";
static uint32_t pm_mix(uint32_t x) { x ^= x >> 16; x *= 0x7feb352du; x ^= x >> 15; x *= 0x846ca68bu; x ^= x >> 16; return x; }
static uint32_t rotl32(uint32_t v, uint32_t n) { return (v << n) | (v >> (32u - n)); }
static int32_t dot4_ref(uint32_t a, uint32_t b, int32_t acc) {
int32_t r = acc; int i;
for (i = 0; i < 4; ++i) { int32_t ba = (int8_t)((a >> (8 * i)) & 0xffu), bb = (int8_t)((b >> (8 * i)) & 0xffu); r = (int32_t)((uint32_t)r + (uint32_t)(ba * bb)); }
return r;
}
static uint32_t lane_ref(int alu, uint32_t g, uint32_t seed, uint32_t steps) {
uint32_t x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u, s; int32_t acc;
if (alu) { for (s = 0; s < steps; ++s) { x = x * 0x9E3779B1u + rotl32(y, 7u); y = (y ^ x) + s; } return x ^ y; }
acc = (int32_t)pm_mix(x);
for (s = 0; s < steps; ++s) { acc = dot4_ref(x, y, acc); x = x * 0x9E3779B1u + (uint32_t)acc; y = rotl32(y, 7u) ^ ((uint32_t)acc + s); }
return (uint32_t)acc ^ x ^ y;
}
typedef struct { cl_platform_id p; cl_device_id d; char pname[128], dname[128], driver[64], ver[64]; } Dev;
static Dev devs[32]; static int ndevs = 0;
static void enumerate(void) {
cl_platform_id ps[8]; cl_uint np = 0, i;
if (clGetPlatformIDs(8, ps, &np) != CL_SUCCESS) return;
for (i = 0; i < np; ++i) {
cl_device_id ds[8]; cl_uint nd = 0, j;
if (clGetDeviceIDs(ps[i], CL_DEVICE_TYPE_GPU, 8, ds, &nd) != CL_SUCCESS) continue;
for (j = 0; j < nd && ndevs < 32; ++j) {
Dev* v = &devs[ndevs++]; v->p = ps[i]; v->d = ds[j];
clGetPlatformInfo(ps[i], CL_PLATFORM_NAME, sizeof v->pname, v->pname, NULL);
clGetDeviceInfo(ds[j], CL_DEVICE_NAME, sizeof v->dname, v->dname, NULL);
clGetDeviceInfo(ds[j], CL_DRIVER_VERSION, sizeof v->driver, v->driver, NULL);
clGetDeviceInfo(ds[j], CL_DEVICE_VERSION, sizeof v->ver, v->ver, NULL);
}
}
}
static int run_variant(cl_context ctx, cl_command_queue q, cl_device_id dev, const char* name, const char* body, int alu, cl_uint lanes, cl_uint steps, int reps, cl_mem out, uint32_t* host, const char* vendor, const char* dname) {
const char* srcs[2] = { COMMON, body }; cl_int err; cl_program prog; cl_kernel k; int r, ok = 1; double best = -1;
prog = clCreateProgramWithSource(ctx, 2, srcs, NULL, &err);
if (err != CL_SUCCESS) { printf("| %s | build failed | clCreateProgramWithSource %d | | | | |\n", name, (int)err); return 0; }
err = clBuildProgram(prog, 1, &dev, "-cl-std=CL1.2", NULL, NULL);
if (err != CL_SUCCESS) {
size_t n = 0; char* log; char* nl;
clGetProgramBuildInfo(prog, dev, CL_PROGRAM_BUILD_LOG, 0, NULL, &n); log = (char*)calloc(n + 1, 1);
if (n) clGetProgramBuildInfo(prog, dev, CL_PROGRAM_BUILD_LOG, n, log, NULL);
while (*log == '\n' || *log == '\r') ++log;
nl = strpbrk(log, "\r\n"); if (nl) *nl = 0;
printf("| %s | build failed | %.160s | | | | |\n", name, log);
printf("RESULT DOT4 vendor=%s device=\"%s\" kernel=%s build=failed\n", vendor, dname, name);
clReleaseProgram(prog); return 0;
}
k = clCreateKernel(prog, "probe", &err);
if (err != CL_SUCCESS) { printf("| %s | build failed | clCreateKernel %d | | | | |\n", name, (int)err); clReleaseProgram(prog); return 0; }
for (r = 0; r < reps; ++r) {
cl_uint seed = 0x2468aceu + (cl_uint)r * 0x9E3779B9u; size_t global = lanes, local = 256; cl_event ev; cl_ulong t0, t1; double ms; uint32_t gs[2]; int j;
clSetKernelArg(k, 0, sizeof(cl_uint), &steps); clSetKernelArg(k, 1, sizeof(cl_uint), &seed); clSetKernelArg(k, 2, sizeof(cl_mem), &out);
err = clEnqueueNDRangeKernel(q, k, 1, NULL, &global, &local, 0, NULL, &ev);
if (err != CL_SUCCESS) { printf("| %s | launch failed | %d | | | | |\n", name, (int)err); clReleaseKernel(k); clReleaseProgram(prog); return 0; }
clWaitForEvents(1, &ev);
clGetEventProfilingInfo(ev, CL_PROFILING_COMMAND_START, sizeof t0, &t0, NULL); clGetEventProfilingInfo(ev, CL_PROFILING_COMMAND_END, sizeof t1, &t1, NULL);
ms = (double)(t1 - t0) / 1e6; clReleaseEvent(ev);
if (best < 0 || ms < best) best = ms;
clEnqueueReadBuffer(q, out, CL_TRUE, 0, (size_t)lanes * 4, host, 0, NULL, NULL);
gs[0] = 0; gs[1] = lanes - 1;
for (j = 0; j < 2; ++j) { uint32_t want = lane_ref(alu, gs[j], seed, steps); if (host[gs[j]] != want) { ok = 0; printf("MISMATCH %s lane %u: gpu %08x cpu %08x\n", name, gs[j], host[gs[j]], want); } }
}
{
double sps = (double)lanes * (double)steps / (best / 1000.0);
printf("| %s | %u | %u | %.3f | %.2f | %.3f | %s |\n", name, lanes, steps, best, sps / 1e9, best * 1e6 / (double)steps, ok ? "yes" : "NO");
printf("RESULT DOT4 vendor=%s device=\"%s\" kernel=%s lanes=%u steps=%u best_ms=%.3f gsteps_per_s=%.2f ok=%d\n", vendor, dname, name, lanes, steps, best, sps / 1e9, ok);
}
clReleaseKernel(k); clReleaseProgram(prog); return 1;
}
int main(int argc, char** argv) {
cl_uint lanes = 1u << 20, steps = 4096u; int reps = 3, device = 0, list = 0, i; Dev* v; cl_int err; cl_context ctx; cl_command_queue q; cl_mem out; uint32_t* host; const char* vendor;
for (i = 1; i < argc; ++i) {
if (!strcmp(argv[i], "--list")) list = 1;
else if (!strcmp(argv[i], "--device") && i + 1 < argc) device = atoi(argv[++i]);
else if (!strcmp(argv[i], "--lanes") && i + 1 < argc) lanes = (cl_uint)strtoul(argv[++i], 0, 10);
else if (!strcmp(argv[i], "--steps") && i + 1 < argc) steps = (cl_uint)strtoul(argv[++i], 0, 10);
else if (!strcmp(argv[i], "--reps") && i + 1 < argc) reps = atoi(argv[++i]);
else { printf("unknown argument %s\n", argv[i]); return 2; }
}
#ifdef IGNEUM_CL_DYNAMIC
if (!ig_cl_load()) { printf("%s\n", ig_cl_error); return 1; }
#endif
enumerate();
if (list || ndevs == 0) { for (i = 0; i < ndevs; ++i) printf("[%d] %s | %s | driver %s | %s\n", i, devs[i].dname, devs[i].pname, devs[i].driver, devs[i].ver); if (ndevs == 0) printf("no OpenCL GPU devices\n"); return ndevs ? 0 : 1; }
if (device < 0 || device >= ndevs) { printf("no device %d (have %d)\n", device, ndevs); return 2; }
v = &devs[device];
vendor = strstr(v->pname, "NVIDIA") ? "nvidia" : (strstr(v->pname, "AMD") ? "amd" : (strstr(v->pname, "Apple") ? "apple" : "other"));
ctx = clCreateContext(NULL, 1, &v->d, NULL, NULL, &err); if (err != CL_SUCCESS) { printf("clCreateContext %d\n", (int)err); return 1; }
q = clCreateCommandQueue(ctx, v->d, CL_QUEUE_PROFILING_ENABLE, &err); if (err != CL_SUCCESS) { printf("clCreateCommandQueue %d\n", (int)err); return 1; }
out = clCreateBuffer(ctx, CL_MEM_READ_WRITE, (size_t)lanes * 4, NULL, &err); if (err != CL_SUCCESS) { printf("clCreateBuffer %d\n", (int)err); return 1; }
host = (uint32_t*)malloc((size_t)lanes * 4);
printf("dot4-probe (OpenCL) on [%d] %s | %s | driver %s | %s, lanes %u, steps %u, best of %d, device event time\n", device, v->dname, v->pname, v->driver, v->ver, lanes, steps, reps);
{
char ext[8192]; ext[0] = 0; clGetDeviceInfo(v->d, CL_DEVICE_EXTENSIONS, sizeof ext, ext, NULL);
printf("cl_khr_integer_dot_product listed: %s\n", strstr(ext, "cl_khr_integer_dot_product") ? "yes" : "no");
}
printf("| kernel | lanes | steps | best ms | G steps/s (= G dot4/s for dot4 rows) | ns per dependent step | lanes 0 and last ok |\n|---|---|---|---|---|---|---|\n");
run_variant(ctx, q, v->d, "alu", K_ALU, 1, lanes, steps, reps, out, host, vendor, v->dname);
run_variant(ctx, q, v->d, "dot4e", K_EMUL, 0, lanes, steps, reps, out, host, vendor, v->dname);
run_variant(ctx, q, v->d, "dot4_khr", K_KHR, 0, lanes, steps, reps, out, host, vendor, v->dname);
if (!strcmp(vendor, "amd")) run_variant(ctx, q, v->d, "dot4_amd", K_AMD, 0, lanes, steps, reps, out, host, vendor, v->dname);
if (!strcmp(vendor, "nvidia")) run_variant(ctx, q, v->d, "dot4_nv", K_NV, 0, lanes, steps, reps, out, host, vendor, v->dname);
clReleaseMemObject(out); clReleaseCommandQueue(q); clReleaseContext(ctx); free(host);
printf("dot4-probe: done\n");
return 0;
}

67
relay/playbooks/dot4-probe.ps1 Normal file
View file

@ -0,0 +1,67 @@
# Igneum run job: dp4a-class throughput (Counter ASIC 2.0 layer 7, docs/analysis/int8-matrix-family.md) on every OpenCL
# GPU of the machine: PC 1 (ae432dc7: RTX 5090 on NVIDIA's OpenCL, RX 9070 XT on the eGPU on AMD's, the gfx1036) or
# PC 2 (1ccfe586: RTX 5090). 5 October 2026. NOT published until the coordinator says "go PC <id>".
# Published as a plain `run` job (NOT --stop-miners), after a `fetch` job with --id fetch-dot4-20261005 that places
# dot4-probe-cl.exe (proto-opencl/dot4-probe.c cross-compiled with mingw, OpenCL.dll loaded at run time) in the jobs
# folder. The script switches off every NVIDIA and gfx1201 card in the app through POST <app.url>api/cards, waits for
# their workers to stop, runs the probe on every device the exe lists (the whole run is under a minute: three to five
# kernels x 3 repetitions x about 10 ms each), and switches the cards back on with the settings they had. Every result
# line starts with RESULT so `node tools/jobs.mjs <job id>` shows them; the probe's own "RESULT DOT4 ..." lines carry
# vendor, device, kernel, best ms and G steps/s, and ok=1 means bit-exact against the CPU reference on two lanes.
$ErrorActionPreference = 'Continue'
function Say([string] $m) { Write-Host ("[" + (Get-Date -Format 'HH:mm:ss') + "] " + $m) }
$jobs = Split-Path $env:IGNEUM_JOB_DIR
$fetched = Join-Path $jobs 'fetch-dot4-20261005'
$exe = Join-Path $fetched 'dot4-probe-cl.exe'
if (-not (Test-Path $exe)) { Write-Output "RESULT error probe missing at $exe (the fetch job runs first)"; exit 2 }
Write-Output "RESULT probe $exe sha256 $((Get-FileHash -Algorithm SHA256 $exe).Hash.ToLower())"
$list = & $exe --list 2>&1
$list | ForEach-Object { "RESULT list $_" }
$devs = @()
foreach ($l in $list) { if ($l -match '^\[(\d+)\]') { $devs += [int]$Matches[1] } }
if ($devs.Count -eq 0) { Write-Output 'RESULT error no OpenCL GPU device listed'; exit 2 }
# the app: switch off the NVIDIA card(s) and the 9070 XT, remember their settings
$appDir = $env:IGNEUM_APP_DIR
if (-not $appDir) { $appDir = Join-Path $env:LOCALAPPDATA 'igneum\app' }
$urlFile = Join-Path $appDir 'app.url'
$url = $null
if (Test-Path $urlFile) { $url = (Get-Content -LiteralPath $urlFile -Raw).Trim() }
$cards = @()
if ($url) {
try {
$st = Invoke-RestMethod -Uri ($url + 'api/state') -Method GET -TimeoutSec 10
$all = $st.mining.cards; if (-not $all) { $all = $st.cards }
$cards = @($all | Where-Object { $_.vendor -eq 'nvidia' -or ($_.vendor -eq 'amd' -and $_.key -match 'gfx1201') })
} catch { Say ("api/state: " + $_.Exception.Message) }
}
if ($cards.Count -gt 0) {
foreach ($c in $cards) { Write-Output ("RESULT card " + $c.key + " enabled=" + $c.enabled + " identities=" + $c.identities + " power_pct=" + $c.power_pct + " state=" + $c.state) }
$body = @{ cards = @($cards | ForEach-Object { @{ key = $_.key; enabled = $false; identities = [int]$_.identities; power_pct = [int]$_.power_pct } }) } | ConvertTo-Json -Depth 5
try { Invoke-RestMethod -Uri ($url + 'api/cards') -Method POST -Body $body -ContentType 'application/json' -TimeoutSec 10 | Out-Null; Say "cards off requested" } catch { Say ("api/cards off: " + $_.Exception.Message) }
$t = 0
while ($t -lt 90) {
Start-Sleep -Seconds 5; $t += 5
try {
$st = Invoke-RestMethod -Uri ($url + 'api/state') -Method GET -TimeoutSec 10
$all = $st.mining.cards; if (-not $all) { $all = $st.cards }
$still = @($all | Where-Object { ($cards.key -contains $_.key) -and -not ($_.state -eq 'off' -and $_.pid -eq 0) })
if ($still.Count -eq 0) { break }
} catch { }
}
Write-Output ("RESULT cards-off after " + $t + " s")
Start-Sleep -Seconds 5
} else { Write-Output 'RESULT card none-found (the app is not running or has no NVIDIA or gfx1201 card); measuring with whatever else runs on the GPUs' }
if (Get-Command nvidia-smi -ErrorAction SilentlyContinue) { & nvidia-smi --query-gpu=name,driver_version,clocks.sm,clocks.mem,temperature.gpu --format=csv,noheader 2>&1 | ForEach-Object { "RESULT gpu-before $_" } }
foreach ($d in $devs) {
Write-Output "RESULT probe device $d start $(Get-Date -Format HH:mm:ss)"
& $exe --device $d 2>&1 | ForEach-Object { if ($_ -match '^RESULT ') { $_ } else { "RESULT $_" } }
}
if (Get-Command nvidia-smi -ErrorAction SilentlyContinue) { & nvidia-smi --query-gpu=clocks.sm,clocks.mem,temperature.gpu --format=csv,noheader 2>&1 | ForEach-Object { "RESULT gpu-after $_" } }
if ($cards.Count -gt 0) {
$body = @{ cards = @($cards | ForEach-Object { @{ key = $_.key; enabled = [bool]$_.enabled; identities = [int]$_.identities; power_pct = [int]$_.power_pct } }) } | ConvertTo-Json -Depth 5
try { Invoke-RestMethod -Uri ($url + 'api/cards') -Method POST -Body $body -ContentType 'application/json' -TimeoutSec 10 | Out-Null; Write-Output "RESULT cards restored" } catch { Write-Output ("RESULT error card restore: " + $_.Exception.Message) }
}
exit 0