[techmore]

What

agx3_matrix_encode_mma8 only packed the 14-byte AGX3C (M5) form. This commit adds the 12-byte AGX3B f32matmad8x8 that Apple's compiler emits on G16 (applegpu_g16g, M4) for simdgroup_float8x8 / simdgroup_half8x8 multiply-accumulate. It uses the same anonymous fields as the AGX3C encoder:

Field	Meaning
2, 18, 35, 52	D, A, B, C registers (16-bit units, register-pair form, r0..r95)
26, 43, 60	A, B, C/D types (1 = fp16, 2 = fp32; C/D is fp32)
27/29, 44/46, 61/63	A, B, C cache / discard hints
11	D cache hint
59	2 = accumulate, 1 = simdgroup_multiply (no C)
69..74	wait mask, six slots (AGX3C has eight)

A and B may have different types, as Metal's half/float mixes do. This works through a new optional b_input_type in agx3_matrix_mma8_desc (0 = same as A). AGX3C still takes one input type and now refuses a different B explicitly.

Why

The AGX3B table already has this form, but nothing encoded it, so cooperative matrices existed only on M5. On M4, the same instruction emitted by another Mesa tree took MLX 4-bit matmul from about 0.8-0.9 to 1.3-1.6 TFLOP/s, and 512-token prefill of 4-bit models went up 1.1-2.1x. These numbers come from our receipts, measured on the GravityLinux Apple9 stack, not on this tree. PR 2 (m4/agx3b-cooperative-matrix) wires this encoding up, opt-in.

Evidence

• Witnesses. We built MSL kernels with metal-tt -arch applegpu_g16g. Each kernel loads every A, B and C fragment from a distinct offset and stores every D to one. So each MAC's operand registers are known from the loads and stores around it, not from the MAC itself.
  • All 177 MACs found that way reproduce byte for byte with this encoder: 123 fp32, 51 with fp16 A and B, and three simdgroup_multiply or mixed-type examples.
  • The test carries 54 of them, which cover every cache, discard, wait-slot and operand-type combination seen.
  • It also checks refusals: a seventh wait slot, bf16, D past r95, an odd register, C together with c_zero, and AGX3C with a mixed B type.
  • The hint and wait bits in the vectors are read back from Apple's bytes. The registers and types are the independent part.
• Cross-check against the generated table. We had first recovered the bit positions independently by solving those MACs: operand bits, the pair-swapped C field, hints at 32/47/48/62/71..73, wait at 12..17 and accumulate at 88. They are exactly where agx3b_tables.c puts these fields.
• Hardware, outside this tree. On the M4 Mac mini (J773g, T8132/G16G), with GravityLinux's Apple9 backend on a Gravity kernel:
  • Vulkan cooperative-matrix GEMMs using the fp32-operand form match the CPU bit for bit from 8x8x8 to 1024x1024x1024, row- and column-major.
  • omarchy-mlx greedy text is identical with the matrix route on and off on 13 models.
  • Our lowering widened fp16 to fp32 first, so the fp16-operand encodings never ran on hardware. They come from Metal output only.

Tests run

• 2026-11-01, macOS (Apple Silicon host, Apple clang 21), this branch: agx_tests 503 tests, 502 passed, 1 skipped (NativeFMA.ExactDyadicOracle, which also skips on the base). The base has 502 tests, 501 passed. NativeForms.AGX3BMma8MetalWitnesses passes. See the INDEX notes for how the macOS build was set up.
• Earlier round, M4 Mac mini, aarch64 Arch Linux, GCC, same base: agx_tests 496 passed and 1 skipped in a release build; 502 passed and 1 skipped with -Db_ndebug=false. The Vulkan driver and drm-shim built.
• Standalone: all 177 witnesses checked against the encoder, 177/177.

Remaining hardware test plan

• No hardware run is needed for this PR by itself, because it changes no shader. Hardware checks belong to PR 2.
• One open point: the witnesses are G16G. G17P shares the AGX3B table, but we did not check any G17P Metal output. If someone has applegpu_g17p output for simdgroup_float8x8, comparing it against these vectors would close this.

Risks and limitations

• The fp16-accumulator form (cf 00 56 ..., a different layout) is not mapped. bf16 is refused, because we have no AGX3B bf16 witness.
• The commit adds a trailing member to agx3_matrix_mma8_desc. Existing aggregate initialisers keep their meaning.

🤖 Generated with Claude Code

[greptile-apps]

Confidence Score: 5/5

[Low impact] The PR appears safe to merge; the previous finding is fixed and no new actionable issues were found.

Summary

Adds the 12-byte AGX3B matrix multiply-accumulate form with separate A and B types.

• The matrix encoder now packs AGX3B’s 12-byte 8x8 instructions.

_{Reviews (2) · Last reviewed commit: "agx: compare the AGX3C B type against A'..." · Reviewed by Greptile}

[techmore]

Aurora downstream draft: aurora-silicon#191