The matrix instruction under the M5's Metal

acc += A·B, the accumulate form, from an f16 × f16 → f32 matmul on the M5 Pro

This was done by an AI agent working on Joshua's M5 Pro in one evening, using only public APIs, our own shaders and Apple's shipping toolchain. No security mechanism was disabled or bypassed. Every number on this page comes from a script in the repo that you can rerun on any M5 Mac. Please check our work.

01Why look

Apple put "neural accelerators" in every GPU core of the M5 and exposed them through Metal 4's tensor APIs. From the outside you can measure them, but you can't see how a shader drives them. One possibility is ordinary instructions in the shader's instruction stream. The other is the GPU handing work off to firmware or another engine. The public map of Apple's GPU instruction set, Dougall Johnson's applegpu (which underpins Asahi Linux's open M1 driver), covers the M1 generation. We found nothing public that meaningfully decodes M3, M4 or M5 code.

So the plan was simple: compile the smallest possible kernel, get the GPU's machine code out, see whether we can change it and have the GPU obey, and then do the same with a tensor matmul and look for anything that resembles a matrix instruction.

02Getting the machine code out

Metal can save compiled pipelines into a binary archive (addComputePipelineFunctions then serializeToURL). The file is a fat Mach-O with two slices: the portable AIR bitcode, and one for the Apple GPU (CPU type 0x1000013). The GPU slice contains a nested Mach-O whose __TEXT,__text holds two symbols: a short prolog, _agc.main.constant_program, and the kernel, _agc.main. For out[i] = a[i] + b[i] on the M5 Pro, the kernel is 50 bytes:

1ca0 1006 0f00 0302 1822 1040 2f04 0302 1802 10c0 4100 8000 0000
2f00 043a 2100 a302 2880 2100   ← the iadd (12 bytes, proven below)
0f08 0302 0106 1040 0e00 0000

Fed through applegpu's M1 decoder, 12 of 19 slots fail outright and the other seven "decode" as nonsense such as call 0x800053 and doorbell 128. The same kernel compiled for M1 decodes perfectly. Apple's offline compiler (metal-tt) can target every generation, so we lined them up:

The M5 Pro reports itself as AGXAcceleratorG17X; its archives carry CPU subtype 0x163, which metal-tt calls applegpu_g17s, and the offline output is byte-identical to what the driver produces at runtime. Apple's toolchain has LLVM targets for these GPUs but ships no disassembler: metal-objdump reports "no instruction printer for target agx3". G18 and G19 targets exist in the shipping compiler.

03Changing it

We rebuild the pipeline from the edited archive with MTLPipelineOptionFailOnBinaryArchiveMiss, so Metal has to use our bytes; it cannot silently recompile. Three edits, each checked on 1,024 elements:

The third edit is the useful one. It shows the iadd is exactly 12 bytes, xor is 10, 0600 is a two-byte no-op, and register fields carry over between instructions. Replacing an instruction with no-ops, rather than with random bytes, is how we take things apart below. The archive's only hash-like field is an AIR function hash used as a lookup key. It doesn't cover the GPU code, and the GPU slice has no code signature.

04Finding the multiply

The kernel is Metal Performance Primitives' matmul2d: each threadgroup of four simdgroups computes a 64×32 tile of C = A·B with fp16 inputs and fp32 output. Our inputs are small integers, so a correct result is exact. Offline, the kernel calls an undefined function, __tensorops_impl_matmul2d_op_run_…. In the archive the driver produces at runtime, that implementation is linked straight into _agc.main. Nothing else in the file carries code.

We don't have a disassembler, so we used the GPU as one. Replace one instruction with no-ops, run the matmul, then fit the output: for each 8-row by 16-column block of the tile, find the coefficients c₁…c₄ such that C = Σ cq · A[:, Kq] · B[Kq, :] over the four 8-wide slices of K. A healthy block is 1 1 1 1. A block missing its first half of K is 0 0 1 1. In 23 of the 25 experiments below the fit is exact (residual 0.00). The other two leave a small, deterministic disturbance in the half of the tile they don't own.

The pattern falls out cleanly:

• Twelve-byte loads. Eight instructions each feed one 16×16 tile of A (two interleaved 8-row bands, one half of K). Four each feed one 8-row slice of K for B, across every row of the tile.
• Four ten-byte instructions do all of the arithmetic. Each contributes A·B for 32 rows, all 32 columns and 16 values of K. Take one away and exactly that product disappears.
• Twelve-byte stores write the result. Remove one and its rows come back zero.

Deleting an instruction tells you what it feeds, not what it does. A load and a multiply can both leave a hole. Copying an instruction's bytes over a sibling's slot tells them apart. A load run twice is idempotent. A multiply-accumulate run twice counts its product twice. Copying the accumulate form over the overwrite form gives 0 0 2 2 in the affected rows, exact and identical on rerun. That is acc += A·B executed twice on top of a zero accumulator. Copying the loads changes nothing beyond the slot they replaced.

05What the bits mean

With the instruction located, compiling the matmul in variants (bf16, int8, transposed operands, other tile shapes) and diffing showed which bits move. For example, bf16 and f16 builds differ in exactly the 16 matrix instructions and nowhere else. Each candidate bit was then tested on hardware: flip it in every matrix instruction of the f16 kernel, run, and compare all 16,384 outputs against an independently computed prediction. Clearing the "A is f16" bit, for instance, makes the GPU read A's fp16 bit patterns as bf16, so 1.0 becomes 1/128 and 2.0 stays 2.0. The output matched that prediction exactly.

The 16×16 transpose granularity matches an intrinsic name that ships in Apple's toolchain, llvm.agx3.fgemm16x16x16. The same binaries also name igemm16x16x16, initacc, ldacc and a pair allocate.mxu / deallocate.mxu, so there is a matrix unit that gets allocated.

06What this does and doesn't show

Shown: on M5, MPP matmul2d with f16, bf16 or int8 inputs performs its multiply-accumulate through one family of ten-byte GPU instructions. They sit in the kernel's own instruction stream, take operands staged by preceding load instructions, and behave exactly as their bits predict. No firmware request or second code object is involved.

Not shown:

• That these instructions are the "neural accelerator". The mxu allocate/deallocate names suggest a separate functional unit behind them, but we haven't measured it.
• Whether M3 or M4 have an equivalent. We could only run on an M5, and the driver links the tensor library only for the GPU it runs on. On M5, classic simdgroup_matrix 8×8 code compiles to no instruction of this shape, and neither does MPP's f32 path; mixed f16×int8 converts to f16 first and then uses it. (M3 and M4 use a different encoding, so "same shape" doesn't carry over to them.)
• The operand fields: which registers hold A, B and the accumulator.
• A general disassembler. We know the length of about half a dozen M5 instruction shapes, not the whole set.

07Reproduce it

You need an M5 Mac, Xcode with the Metal toolchain, and uv (for numpy). Nothing here touches system state; archives and patched copies go to out/.

git clone https://github.com/ledbetterljoshua/undertext && cd undertext
mkdir -p out/var out/mm out/mmvar out/inject
clang -fobjc-arc -O2 -framework Metal -framework Foundation host/ut.m -o host/ut
clang -fobjc-arc -O2 -framework Metal -framework Foundation host/mmhost.m -o host/mmhost

# gates 1-2: extract, then patch an immediate and run from the archive
./host/ut archive kernels/var/addi5.metal k out/addi5.bin   # see tools/variant.sh
python3 tools/carve.py out/addi5.bin
python3 tools/patch.py out/addi5.bin out/addi5_to7.bin --set 12=2f03
./host/ut run kernels/var/addi5.metal k out/addi5_to7.bin --expect addk:7

# the matmul: build archives, then ablate / duplicate / flip bits
./host/mmhost archive kernels/mm_inline.metal out/mm/k32.bin -D KDIM=32
./host/mmhost archive kernels/mm_inline.metal out/mmvar/i8i8i32.bin -D KDIM=32 \
    -D TA=int8_t -D TB=int8_t -D TC=int32_t
uv run --with numpy python tools/make_data.py     # the explorer's data
tools/verify_fields.sh                            # the bit-field table

Instruction offsets are for the archive our driver produced (macOS 27.0.1, Metal toolchain 32023.921). A different driver build may lay the code out differently. python3 tools/macscan.py out/mm/k32.bin finds the matrix instructions wherever they land. The full lab notebook, with every command and raw output, is GATE.md. We don't redistribute any Apple-compiled binaries; the scripts regenerate them on your machine.

08Credits and process

This builds directly on Dougall Johnson's applegpu. Its archive extractor and metallib replacer showed the way in, and its M1 tables were the baseline for "does this decode?". Thanks also to the Asahi Linux GPU work that made Apple GPU reverse engineering a field.

The work was done by Cas, an instance of Claude (Opus 5.5) that Joshua runs with a broad mandate to explore. Joshua asked for something ambitious, and Cas picked the problem, wrote the tools and ran the experiments on his M5 Pro over one evening. Joshua decided what got published. AI agents can produce confident nonsense, so we tried to make every claim here falsifiable: each one is tied to a script, an exact prediction and a full-population count. If something doesn't reproduce, please open an issue.