Counter ASIC 2.0 layer 7: integer matrix family design (vendor primitives cited: PTX dp4a and mma .u8/.s8, AMD v_dot4_i32_iu8 and WMMA iu8 on RDNA 3 and 4 via LLVM and GPUOpen, CDNA 3 MFMA i8, Metal 4 matmul2d char x char -> int found in MSL 4.1 table 7.3; dot4 and mm8 semantics, reserve entry R1, emulation rule); M5 Max dot4 probe numbers in the bench log; PC playbook dot4-probe.ps1 prepared, not published
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
This commit is contained in:
igneum-josh
parent
5d5ba15ed7
commit
f59708dff7
3 changed files with 241 additions and 0 deletions
159
docs/analysis/int8-matrix-family.md
Normal file
159
docs/analysis/int8-matrix-family.md
Normal file
|
|
@ -0,0 +1,159 @@
|
|||
# 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 or 2), NVIDIA OpenCL inline PTX, and CUDA `__dp4a` | owed: PC job prepared (`relay/playbooks/dot4-probe.ps1`, exe `dot4-probe-cl.exe` cross-compiled, sha256 in the bench-log entry); not published until the coordinator's "go PC" | | | | | | |
|
||||
| RX 9070 XT (PC 1), AMD OpenCL `__builtin_amdgcn_sudot4` | owed, same job (the job runs every listed device: the 5090 on NVIDIA's OpenCL, the 9070 XT, the gfx1036) | | | | | | |
|
||||
|
||||
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). 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 (inline PTX through NVIDIA OpenCL; `__dp4a` through CUDA if the PC has nvcc) | owed, PC job prepared, waiting for "go PC" |
|
||||
| `sudot4` on the 9070 XT through Adrenalin's OpenCL C, and whether `cl_khr_integer_dot_product` appears on the 3683.0 platform | owed, same job; the probe prints both |
|
||||
| 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 |
|
||||
|
|
@ -1524,3 +1524,18 @@ What is measured: one BLS12-381 aggregate signature over 16 summed G1 keys plus
|
|||
| on, split 90 s | v3 | 0 / 2 | none / 3 | 278 / 265 | apart | none | 3 on n0 | 2 (n0 reconnected 6 s after the heal, A's chain at about 58 DAA, inside the table) |
|
||||
|
||||
Reading (the NEW finding, ledger C4). With the module off GHOSTDAG alone converges on the heavier chain and the losing side's records re-determine (F24 works when the chain moves). With the module on the overlay holds during the split (A, with 30% of the frozen table, locks nothing; B locks 7 and 8) and then fails at the heal in the shipped node: B's certificates for blocks off n0's chain are "kept pending until the chain decides (no lock at this index)", n0's chain never decides because GHOSTDAG keeps its heavier tip and nothing turns the certificate into a fork-choice constraint, and once n0's last lock (index 7, DAA 209) is one window old (DAA 329) the frozen table stops applying on A's chain ("no frozen table (no lock on this chain inside the window)"), A's two keys are 100% of A's own window (B's post-cut blocks are red there) and n0 locks 10, 11, 12 alone; B's certificates for 10 and 11 then log CONFLICTING on n0 (n0 log, 17:27:04 to 17:29:54 BST). A finality fork from a 96-s honest partition, no attacker, table intact at the heal; the 150-s run and the v2 control end the same way. The spec's fork choice ("GHOSTDAG among tips through all certified checkpoints", 3.5) is therefore implemented only for certificates over blocks already on the node's chain. Fix named in the ledger entry: verify an off-chain certificate against the table at its own block and let it constrain fork choice (a certificate-driven reorg), then re-determine. Raw: `scratchpad fud-a/c4-results-*.md`, node logs `c4-on90-tmp/`, `c4-v2-control-tmp/`.
|
||||
|
||||
## 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.
|
||||
|
|
|
|||
67
relay/playbooks/dot4-probe.ps1
Normal file
67
relay/playbooks/dot4-probe.ps1
Normal 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
|
||||
Loading…
Reference in a new issue