igneum/docs/analysis/int8-matrix-family.md

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