`llvm_intrinsic` cannot bind `air.simdgroup_matrix_8x8_load` on Apple GPU — `pop.call_llvm_intrinsic` legalization rejects the pointer operand

TL;DR

On Apple Silicon, max.gpu.compute.arch.mma_apple.apple_mma_load_8x8 is documented to return "a
SIMD vector of 2 elements for this thread's fragment," and — under every calling convention we
tried, including a maximally-favorable minimal probe — always lowers to a manual 2-element load
plus a shufflevector widening into the 64-lane MMA fragment value, never to a single native
air.simdgroup_matrix_8x8_load instruction. Apple's own Metal compiler front-end (confirmed via
llama.cpp's compiled Metal kernels) emits that native instruction routinely for the identical
operation. We tried to bind it ourselves via Mojo's own documented low-level escape hatch,
std.sys.llvm_intrinsic[name, ReturnType](*args) — the same generic mechanism
apple_mma_load_8x8's own sibling (_mma_apple_8x8, the multiply-accumulate op) already uses
internally. That fails to compile, with a specific, reproducible MLIR legalization error,
whenever the call includes a pointer-typed operand — which this intrinsic's real ABI requires. We
independently reproduced this failure in a brand-new, isolated toolchain install (see
"Verification" below), on mojo==1.1.0.dev2026080606 / max==26.6.0.dev2026080606.

This was found while investigating a real, measured ~1.73x-2.88x per-dispatch GPU cost gap (under
realistic dense batching) between our own Apple-8x8-MMA GEMM kernel and llama.cpp's hand-written
Metal GEMM kernel, at an identical, real LLM qkv_proj shape and identical multiply-accumulate
instruction count. We closed most of that gap with an orthogonal, unrelated fix (wider cooperative
memory loads), narrowing it from ~2.88x to ~1.73x — but every kernel variant we tried plateaus at
that ~1.73x floor, and manual fragment construction (this issue) is the most plausible remaining,
not-yet-closable contributor, because it is a confirmed toolchain ceiling, not a tuning choice we
can work around from user code.

---

1. Minimal, self-contained reproduction (fails to compile)

File: repro_native_intrinsic_fails.mojo (same directory as
this issue body). No imports outside standard std.*/max.gpu.* — no project-specific modules.

It defines a small GPU kernel that:

1. Stages an 8x8 tile into AddressSpace.SHARED memory (identical convention apple_mma_load_8x8
   itself uses).
2. Calls a new function, apple_mma_load_8x8_native, which attempts to bind
   air.simdgroup_matrix_8x8_load directly via llvm_intrinsic, using the exact 4-argument ABI
   read from Apple's own compiled Metal output (see below) — (pointer, dims, strides, offsets).
3. Feeds the result into a multiply-accumulate call (also via llvm_intrinsic, bound the same way
   _mma_apple_8x8 already binds it internally — this half is not the part that fails) so the
   whole pipeline is dispatched through DeviceContext.enqueue_function, which is what actually
   forces Mojo to lower the kernel and hit the failure below.

The call site, verbatim:

from std.sys import llvm_intrinsic

comptime _Dims = SIMD[DType.int64, 2]

def apple_mma_load_8x8_native[
    dtype: DType
](
    ptr: UnsafePointer[
        Scalar[dtype], address_space = AddressSpace.SHARED, origin=MutUntrackedOrigin
    ],
    stride: Int,
) -> SIMD[dtype, 64]:
    return llvm_intrinsic["llvm.air.simdgroup_matrix_8x8_load", SIMD[dtype, 64]](
        ptr,
        _Dims(Int64(stride), 8),
        _Dims(1, Int64(stride)),
        _Dims(0, 0),
    )

Why this ABI specifically: read directly from xcrun metal-objdump --disassemble output on
llama.cpp's own compiled kernel_mul_mm_f16_f32.metallib (Apple's own Metal front-end's emitted
IR, not guessed), which declares and calls:

declare <64 x float> @air.simdgroup_matrix_8x8_load.v64f32.p3f32(
    float addrspace(3)* nocapture readonly, <2 x i64>, <2 x i64>, <2 x i64>
) local_unnamed_addr
%419 = tail call fast <64 x float> @air.simdgroup_matrix_8x8_load.v64f32.p3f32(
    float addrspace(3)* nocapture readonly %196,
    <2 x i64> <i64 64, i64 8>, <2 x i64> <i64 1, i64 64>, <2 x i64> zeroinitializer
)

i.e. (pointer, dims={stride, 8}, strides={1, stride}, offsets={0, 0}). llama.cpp's flash-attention
kernel issues the same shape for the F16 element type
(@air.simdgroup_matrix_8x8_load.v64f16.p3f16), confirming the ABI generalizes across dtype.

llvm_intrinsic's own generic signature, llvm_intrinsic[intrinsic_name: StringLiteral, ReturnType](*args), was discovered from the compiler's own suggested-fix diagnostic when calling
it with the wrong shape — this is Mojo's own documented mechanism, and it is the SAME mechanism
max.gpu.compute.arch.mma_apple's own _mma_apple_8x8 already uses to bind the sibling
multiply-accumulate intrinsic (confirmed by decompiling that module's own .mojoc bytecode: the
literal comptime string "llvm.air.simdgroup_matrix_8x8_multiply_accumulate" appears in its own
string-literal table). That multiply-accumulate half already compiles and runs correctly in
production today — only the load intrinsic, when given a pointer operand, is rejected.

Run:

mojo run repro_native_intrinsic_fails.mojo

2. Compile error output

Independently re-captured for this report (not just re-quoted from our earlier investigation) in a
brand-new, isolated toolchain install — see "Verification" below for exactly how. Exact terminal
output (warnings about unrelated deprecated APIs elided; nothing else edited):

/private/tmp/upstream_mma_repro_test/repro_native_intrinsic_fails.mojo:1:1: error: failed to lower module to LLVM IR for archive compilation, run LowerToLLVMPipeline failed
oss/modular/mojo/stdlib/std/sys/intrinsics.mojo:77:10: error: expected expanded operands: A, B, C
/private/tmp/upstream_mma_repro_test/repro_native_intrinsic_fails.mojo:77:81: note: called from
    return llvm_intrinsic["llvm.air.simdgroup_matrix_8x8_load", SIMD[dtype, 64]](
                                                                                ^
oss/modular/mojo/stdlib/std/sys/intrinsics.mojo:77:10: note: see current operation: %69 = "pop.call_llvm_intrinsic"(%67) <{fastmathFlags = #pop<fmf none>, hasSideEffects = true, intrin = "llvm.air.simdgroup_matrix_8x8_load" : !kgen.string}> : (!kgen.struct<(pointer<none, 3>, simd<2, si64>, simd<2, si64>, simd<2, si64>)>) -> !kgen.simd<64, f32>
oss/modular/mojo/stdlib/std/sys/intrinsics.mojo:77:10: error: failed to legalize operation 'pop.call_llvm_intrinsic' that was explicitly marked illegal: %68 = "pop.call_llvm_intrinsic"(%67) <{fastmathFlags = #pop<fmf none>, hasSideEffects = true, intrin = "llvm.air.simdgroup_matrix_8x8_load" : !kgen.string}> : (!kgen.struct<(pointer<none, 3>, simd<2, si64>, simd<2, si64>, simd<2, si64>)>) -> !kgen.simd<64, f32>
/private/tmp/upstream_mma_repro_test/repro_native_intrinsic_fails.mojo:77:81: note: called from
    return llvm_intrinsic["llvm.air.simdgroup_matrix_8x8_load", SIMD[dtype, 64]](
                                                                                ^
oss/modular/mojo/stdlib/std/sys/intrinsics.mojo:77:10: note: see current operation: %68 = "pop.call_llvm_intrinsic"(%67) <{fastmathFlags = #pop<fmf none>, hasSideEffects = true, intrin = "llvm.air.simdgroup_matrix_8x8_load" : !kgen.string}> : (!kgen.struct<(pointer<none, 3>, simd<2, si64>, simd<2, si64>, simd<2, si64>)>) -> !kgen.simd<64, f32>
/private/tmp/upstream_mma_repro_test/.pixi/envs/default/bin/mojo: error: failed to run the pass manager

The load-bearing lines: expected expanded operands: A, B, C and failed to legalize operation 'pop.call_llvm_intrinsic' that was explicitly marked illegal, both pointing into the toolchain's
own bundled stdlib at oss/modular/mojo/stdlib/std/sys/intrinsics.mojo:77:10 — i.e. inside
llvm_intrinsic's own generic implementation, not at our call site. The "note: called from"
lines point at our file only to show the call chain; the actual rejection is raised by the
compiler's own lowering pass before it ever reaches GPU-specific codegen (this compiles as far as
producing an MLIR module and then fails "failed to lower module to LLVM IR for archive compilation"
— it never gets far enough to attempt Metal/AIR codegen at all).

Additional diligence performed (in our source investigation) to rule out a call-site mistake,
summarized because it is directly relevant to whether this is fixable from user code:

• The identical failure occurs for both DType.float32 and DType.float16 return types.
• The identical failure occurs for both AddressSpace.SHARED (addrspace(3), used above) and
  device/global (addrspace(1)) pointers.
• A fake intrinsic name with 4 scalar (non-pointer) arguments compiles cleanly through this
  same legalization stage, failing only later with "intrinsic not found" (the expected failure
  mode for a made-up name) — i.e. the legalization pass itself does not reject arbitrary/unknown
  intrinsics outright.
• The real intrinsic name (llvm.air.simdgroup_matrix_8x8_load) called with an all-vector,
  no-pointer argument list (a deliberately wrong ABI) also compiles cleanly with no validation at
  this stage.
• Any pointer-plus-multiple-vector-arguments shape — the real ABI this intrinsic actually
  needs — either hits this exact clean rejection (real name) or crashes the Mojo compiler
  outright with a SIGSEGV (a fake name given the same pointer-plus-vectors shape).

That last point means the compiler has some specific, named awareness of
air.simdgroup_matrix_8x8_load (a fake name with the identical argument shape crashes rather than
producing this clean diagnostic) — this is a real, specific ceiling in the
pop.call_llvm_intrinsic legalization path for this intrinsic shape, not a Mojo-level ABI mistake
on the calling code's part, and not something achievable by varying argument order, count, or type
from user code.

3. Contrast: apple_mma_load_8x8's existing, working manual load+shuffle lowering

File: repro_manual_load.mojo (same directory). This one compiles
and runs correctly — it is the currently-only-available, public code path, included so a
Modular engineer can directly diff "works" against "fails to compile" (§1-2) side by side. It
calls the real, existing, documented max.gpu.compute.arch.mma_apple.apple_mma_load_8x8 under the
single most idiomatic/favorable calling convention constructible (shared-memory buffer that IS the
8x8 tile, row-major, stride equal to the tile's own width, one simdgroup, zero pointer-offset
arithmetic) — deliberately removing every confound a larger, real kernel's own tile-of-a-larger-
buffer indexing would introduce.

We disassembled our own production kernel that calls this same wrapper
(_qkv_proj_gemm_mma8x8, from the LLM-inference codebase this investigation was run in) via
xcrun metal-objdump --disassemble, and captured this pattern directly in its compiled LLVM IR.
Exact excerpt (trimmed for readability — the enormous literal poison/undef fill lists in the
shufflevector masks are abbreviated with [...]; nothing else is altered), with its precise
source citation for traceability:

Source: /tmp/redline_static_analysis/algenta_qkv_proj_gemm_mma8x8.func.ll, lines 228-230 and
243-244 (an ephemeral disassembly artifact from our own investigation's tooling pipeline, not
tracked in git — quoted here verbatim from the file while it was still present on disk):

; line 228-229: manual 2-element load from the shared-memory tile, then widen to 64 lanes
; (this thread's real 2 elements go into lanes [0,1]; the other 62 lanes are `undef` here)
%190 = load <2 x float>, <2 x float> addrspace(3)* %189, align 4
%191 = shufflevector <2 x float> %190, <2 x float> poison,
       <64 x i32> <i32 0, i32 1, [i32 undef x 62]>

; line 230: second shufflevector merges those 2 real lanes into a zero-filled 64-wide vector
; (lanes [2..63] become the literal float 0.0, matching the accumulator's own zero-init)
%192 = shufflevector <64 x float> %191, <64 x float> <float poison, float poison, [float 0.0 x 62]>,
       <64 x i32> <i32 0, i32 1, [i32 66..127]>

; ... (lines 231-242: the SAME two-shufflevector pattern repeated for the B operand and the
;      accumulator operand -- three total 2-element loads, six total shufflevectors, to build
;      the three 64-wide operands one multiply-accumulate call needs)

; line 243: the ACTUAL hardware MMA op -- takes the three manually-widened 64-wide operands
%203 = call <64 x float> @air.simdgroup_matrix_8x8_multiply_accumulate.f.f.v64f32.v64f32.v64f32.v64f32(
           <64 x float> %192, <64 x float> %197, <64 x float> %202) #2

; line 244: narrow the 64-wide result back down to this thread's real 2 elements before storing
%204 = shufflevector <64 x float> %203, <64 x float> undef, <2 x i32> <i32 0, i32 1>

This exact pattern (one 2-element load, two shufflevectors to widen each of the 3 MMA operands,
one more shufflevector to narrow the result back down) repeats identically for every
apple_mma_load_8x8 call site in that production kernel, for both DType.float32 and
DType.float16, and reproduces identically when disassembling repro_manual_load.mojo itself
(the minimal contrast file in this directory) — it is not specific to that one kernel's own
indexing.

For comparison, llama.cpp's compiled kernel_mul_mm_f16_f32.metallib — computing the
mathematically identical operation — loads the full 64-element fragment in a single instruction, no
manual construction:

declare <64 x half> @air.simdgroup_matrix_8x8_load.v64f16.p3f16(
    half addrspace(3)* nocapture readonly, <2 x i64>, <2 x i64>, <2 x i64>
) local_unnamed_addr
%262 = tail call fast <64 x half> @air.simdgroup_matrix_8x8_load.v64f16.p3f16(
    half addrspace(3)* nocapture readonly %261,
    <2 x i64> <i64 8, i64 8>, <2 x i64> <i64 1, i64 8>, <2 x i64> zeroinitializer
)

Confirmed not an Algenta-specific workaround: decompiling max.gpu.compute.arch.mma_apple's
own .mojoc bytecode shows its apple_mma_load_8x8 docstring states plainly, "Returns: SIMD
vector of 2 elements for this thread's fragment" — this is the wrapper's documented, intentional
behavior. There are zero occurrences of simdgroup_matrix_8x8_load anywhere in that module or in 7
other decompiled MAX/stdlib packages we checked. Modular's own production Apple-8x8 reference GEMM
kernel, linalg.gpu.apple.matmul_kernel, routes through this same wrapper — every Mojo/MAX
caller on this toolchain, including Modular's own shipping reference kernel, is limited to this
manual construction; there is currently no public code path to the native instruction at all.

4. The GEMM this was discovered in, and the correctness convention used

Discovered while benchmarking a real production kernel, _qkv_proj_gemm_mma8x8 (an Apple-8x8-MMA
GEMM implemented in Mojo, dispatched via max.gpu.host.DeviceContext), against the identical,
real qkv_proj projection shape from a Qwen3-0.6B-Base forward pass: M=256, K=1024, N=4096
(FLOPs = 2×M×K×N = 2,147,483,648, i.e. ~2.147 GFLOP). This is a compute-bound shape at this
kernel's realized arithmetic intensity (roughly 150-230 FLOP/byte against the host's measured ~411
GB/s streaming bandwidth ceiling and a ~16-18 TFLOP/s public-spec FP32 compute figure for this GPU
— the roofline ridge point works out to ~40 FLOP/byte, 4-6x below either side's realized intensity),
so multiply-accumulate throughput, not memory bandwidth, is the dominant cost lever at this shape.

Every correctness check in this investigation (including the two files in this issue) used the
same convention: an independent CPU-side Float64 reference implementation (a plain triple-loop
matmul, computed at full Float64 precision, never sharing code with the GPU kernel under test),
checked against the GPU kernel's real output read back from device memory, at both a small/ragged
shape and the real M=256/K=1024/N=4096 shape, reporting the worst absolute per-element delta.
Representative figures from kernels central to this investigation (F32-accumulate throughout;
tolerance is looser for F16-narrowed variants because of the format's own reduced mantissa, not a
weaker check):

• The wide-load kernel that closed most of the gap (_qkv_proj_gemm_mma8x8_vecload4, §5 below):
  worst delta 5.2e-8 (ragged shape) / 4.3e-7 (real qkv_proj shape) — both far under any
  reasonable F32 tolerance, confirming its memory-access rewrite is bit-for-bit equivalent math to
  the original kernel; only the fetch instructions differ.
• The minimal fragment-load contrast probe (§3, repro_manual_load.mojo, and its direct ancestor
  in our own repo, mma_forge/probes/mma_fragment_load_probe.mojo): worst delta 1.29e-7 (F32)
  / 5.0e-4 (F16) computing D = A @ A + 0 for a single 8x8 tile against the same
  independent-Float64-oracle convention — both PASS at that file's own thresholds (<= 0.01 F32,
  <= 0.05 F16). We re-ran this exact file in a fresh, isolated install while preparing this
  report (see "Verification" below) and got the identical two numbers, confirming both the
  correctness convention and the toolchain behavior are stable and reproducible.

The fragment-loading gap documented in this issue is not proven to be the sole or dominant
remaining cause of the residual GEMM performance gap described next — it is the most plausible,
concretely-identified, structurally-confirmed remaining contributor after every other lever we
could try from user code was tested and either refuted or already captured.

5. Real, measured batched-GPU timing comparison

All figures below are batched GPU-only marginal cost per dispatch — i.e. (GPUEndTime - GPUStartTime) captured around one shared MTLCommandBuffer encoding N identical GEMM dispatches
back-to-back (same command-buffer-batching methodology on both sides), fit via OLS across N in
{1,2,4,8,16,28,56,112}, at the real M=256/K=1024/N=4096 qkv_proj shape. This is deliberately NOT an
isolated single-dispatch number: an isolated-dispatch comparison at this same shape originally,
misleadingly, showed only a ~20-26% gap, because CPU-side dispatch/sync overhead dominates at this
scale for single dispatches — the batched-GPU-only methodology was built specifically to remove
that confound, and it revealed a much larger real gap.

llama.cpp's own kernel_mul_mm_f16_f32 batched marginal cost, measured the same way: ~0.205 -
0.209 ms/dispatch (converges from N=8 through N=112 in every clean run; GPU-only-per-dispatch at
that plateau is ~0.204ms — about half the isolated single-dispatch GPU-only median of 0.410ms).

Kernel	Batched marginal (mean)	Gap to llama.cpp (~0.207ms)
Original _qkv_proj_gemm_mma8x8 (all-F32 MMA operands, scalar cooperative loads)	0.5957 ms	2.88x
_qkv_proj_gemm_mma8x8 + BM64 tile-size change alone	0.5326 ms	2.57x
_qkv_proj_gemm_mma8x8_f16 (F16xF16→F32 MMA operands, matching llama.cpp's operand type, load path unchanged)	0.55-0.80 ms (unstable across N, sometimes worse)	2.7-3.9x
_qkv_proj_gemm_mma8x8_vecload4 (4-wide vectorized cooperative loads, MMA operand type and tile shape unchanged)	0.3548 ms	1.73x
BM64 + vecload4 stacked	0.3596 ms	1.74x (cannibalized — no compounding)
vecload4 + F16 stacked	0.3581 ms	1.73x (cannibalized — no compounding)

And the isolated-vs-batched ratio, showing the mechanism most directly relevant to this issue —
every Algenta-side variant we measured gets essentially zero cheaper per dispatch as batch
depth increases, while llama.cpp's kernel gets roughly 2x cheaper:

	Isolated N=1 GPU-only (mean)	Batched marginal (mean)	Marginal / Isolated ratio
Original _qkv_proj_gemm_mma8x8	0.5647 ms	0.5957 ms	1.055 (flat)
_qkv_proj_gemm_mma8x8_vecload4	0.3455 ms	0.3548 ms	1.027 (flat)
llama.cpp kernel_mul_mm_f16_f32	0.410 ms	~0.205-0.209 ms	~0.50 (real ~2x amortization)

What this rules in and out, precisely: five independent kernel-architecture variants and three
stacked combinations were measured this way. Wider cooperative memory loads (vecload4) is the
single dominant, load-bearing lever found — it alone narrows the gap from 2.88x to 1.73x. Neither
MMA operand narrowing (F32→F16, matching llama.cpp's own operand type) nor a tile-shape change
(BM64) compounds with it, tested twice independently (both land within noise of vecload4 alone).
A live GPU-clock (DVFS) capture additionally shows every Algenta-side variant reaches "Maximum"
clock state from its very first dispatch, while llama.cpp's kernel never once reaches "Maximum" in
any capture (caps at "Medium") — ruling out "insufficient clock boost" as an explanation for
Algenta's slower per-dispatch cost, since Algenta's kernel is already boosted higher. Native
fragment loading (this issue) is the one remaining, structurally-distinct, previously-untested axis
identified before the investigation concluded — every combination of load width, tile shape, and
MMA operand type available from Mojo/MAX user code has now been tried at this shape; this specific
mechanism (single native fragment-load instruction vs. manual per-thread reconstruction) could not
be tested at all, for the compiler reason in §1-2, and remains the most concrete, structurally-
motivated candidate for narrowing the residual ~1.73x further.

5b. Measured compiler impact: a native-Metal-fragment-load control kernel (a real number, not a projection)

§5 identified native fragment loading as "the one remaining, structurally-distinct, previously-
untested axis" and §6 below offers two fix directions — but neither could be measured from
Mojo/MAX itself, because the toolchain ceiling in §1-2 makes that impossible from user code. This
section closes that gap with a real, independently-built, independently-timed control experiment
written directly in Metal Shading Language (which, unlike Mojo, CAN express the native primitive):
a hand-written kernel, qkv_proj_vecload4_native_fragment.metal (+ dispatch harness,
native_fragment_oracle_harness.m, both at
benchmarks/llm_smoke/metal_native_fragment_oracle/ in the source repository), that is a
line-by-line port of _qkv_proj_gemm_mma8x8_vecload4 (§5's own gap-closing kernel — same
BM=32/BN=32/BK=16 tile, same 8-simdgroup/256-thread grid, same 4-wide cooperative device-memory
loads, same Int8+scale dequant, same accumulator ownership) with exactly one variable changed:
the fragment load/multiply-accumulate/store, from Mojo's only-available manual
2-element-load-plus-two-shufflevector construction (§3) to MSL's native simdgroup_load /
simdgroup_multiply_accumulate / simdgroup_store built-ins — the same primitives llama.cpp's own
kernel_mul_mm uses, and the exact instruction §1-2 show Mojo/MAX cannot bind from user code today.

Correctness (same independent CPU-Float64-oracle convention as §4): worst delta 4.32e-7 at
the real qkv_proj shape (M=256/K=1024/N=4096) — matches vecload4's own documented 4.3e-7 at this
exact shape almost to the digit, strong evidence the ported math is equivalent.

Build/dispatch confirmed real: xcrun metal-objdump --disassemble on the compiled AIR shows
air.simdgroup_matrix_8x8_load.v64f32.p3f32 (both operands, threadgroup address space),
air.simdgroup_matrix_8x8_multiply_accumulate.v64f32..., and
air.simdgroup_matrix_8x8_store.v64f32.p1f32 (device address space) — 15 total
simdgroup_matrix_8x8_* instructions, and exactly one shufflevector in the entire compiled
kernel (an unrelated scalar-broadcast splat for the wscale dequant, present in vecload4's own Mojo
kernel too) — the decisive structural contrast with the load-plus-two-shufflevector pattern §3
documents for every Mojo kernel.

Timing: measured with the exact batched-GPU-only methodology used throughout this
investigation (MTLCommandBuffer.GPUStartTime/GPUEndTime around one shared command buffer
encoding N identical dispatches, N in {1,2,4,8,16,28,56,112}, 50 warm reps/N, real qkv_proj shape),
5 independent process launches, machine GPU-load checked (ps aux/top -l1) before every run (no
competing GPU-heavy process at any point; pmset -g therm reported no thermal warning throughout):

Run	GPU-only marginal (OLS, N=1..112)	N>=8 plateau mean
1	0.4203 ms	0.4204 ms
2	0.4200 ms	0.4204 ms
3	0.4206 ms	0.4204 ms
4	0.4203 ms	0.4204 ms
5	0.4201 ms	0.4205 ms

Extraordinarily tight across all 5 runs — grand N>=8 plateau mean 0.4204 ms, standard
deviation 0.00015 ms across all 25 plateau points (N=8/16/28/56/112 x 5 runs) — a clean,
decisive, reproducible number, not a noisy one. (Isolated N=1 GPU-only shows the same bimodal
variance §5's own isolated-N=1 figures show for every kernel in this tournament — 0.344-0.552 ms
across the 5 runs — consistent with the already-documented real GPU idle<->active power-state
effect at short bursts, not a defect in this measurement; the batched N>=8 plateau, not the
isolated figure, is the number that matters here, per this investigation's own established
convention. The rotating-distinct-buffer-set regime used elsewhere in this investigation for a
subset of kernels was not additionally run for this control — the machine had only ~4-5GB free at
measurement time, and the hot-regime result below is already decisive without it.)

Comparison against §5's own table, and the investigation's pre-registered decision thresholds
(~0.21-0.25ms = missing fragment primitive explains most of the remaining gap; ~0.26-0.30ms =
important compiler loss, other bookkeeping/layout costs remain; ~0.31-0.34ms = useful but not
dominant; ~0.35ms+ = the API limitation is real but not the performance root cause, statistically
no better than vecload4's own Mojo result):

Kernel	Batched marginal	Gap to llama.cpp (~0.207ms)
_qkv_proj_gemm_mma8x8_vecload4 (Mojo, manual fragment — §5's own figure)	0.3548 ms	1.73x
qkv_proj_vecload4_native_fragment (MSL, native fragment — this section)	0.4204 ms	2.03x
Original _qkv_proj_gemm_mma8x8 (Mojo, manual fragment, no vecload4)	0.5957 ms	2.88x
llama.cpp kernel_mul_mm_f16_f32	~0.207 ms	1.00x

Verdict: 0.4204ms lands squarely in, and above, the pre-registered ">=0.35ms" bucket. This is
not merely "statistically no better than vecload4's own Mojo result" (that bucket's own wording) —
it is measurably, reproducibly worse: +18.5% slower than vecload4's manual-shufflevector
construction (0.4204 vs 0.3548ms), widening rather than narrowing the gap to llama.cpp (2.03x vs
vecload4's 1.73x). Native fragment loading, at least in this direct, otherwise-identical port, is
not the remaining bottleneck for this kernel/shape — swapping it in made this specific kernel
slower, not faster.

What this does and does not settle: this directly answers §5's own open question — how much of
vecload4's remaining ~1.73x gap to llama.cpp closes with native fragment loading — with a real
number: none of it; if anything the gap widens. It does not explain why (no Instruments
occupancy/register/scheduling capture was run for this specific kernel; that remains open, the same
way the deeper "why does llama.cpp amortize batching and this codebase's kernels don't" question in
§5 also remains open — a plausible candidate is that simdgroup_load/_store carry inherent
simdgroup-wide synchronization or scheduling cost that Mojo's manual per-thread 2-element load
avoids, but this is a hypothesis, not something this measurement isolates). It also does not
retroactively make the compiler limitation in §1-2 unimportant on its own terms — parity with
hand-written Metal, and unblocking future kernels/shapes where manual-construction cost might
dominate differently, are still worth having — but it does mean §5's own framing of native fragment
loading as the "most plausible remaining... contributor" to the residual gap is now falsified by
direct measurement for this specific kernel and shape, and should not be cited as a performance
argument for prioritizing this fix. Full raw run output is in the source investigation's own
telemetry ledger (benchmarks/llm_smoke/REALRACE_RESULTS.md / PROJECT_REDLINE_TELEMETRY.md,
"Project Redline reopening" entries).

6. Proposed fix direction (suggestions, not demands)

We do not have visibility into Mojo/MAX's compiler internals or design constraints, so both of
these are offered as directions for Modular's own compiler team to evaluate, not prescriptions:

A. A first-class, opaque matrix-fragment value type (larger-scope, higher-level). Rather than
exposing apple_mma_load_8x8 as a function returning a plain SIMD[dtype, 2] (the caller's literal
per-thread scalars) that must then be manually widened before every _mma_apple_8x8 call, a
dedicated fragment type — e.g. MatrixFragment8x8[dtype] — with only load_fragment(ptr, stride),
mma(a, b, c) -> MatrixFragment8x8, and store_fragment(ptr, stride, frag) operations, and
no per-thread-scalar indexing/construction exposed to user code at all, would let the backend
choose the native single-instruction load whenever the target ISA supports it (Apple AIR does),
falling back to the current manual construction only where the hardware truly requires it. This
mirrors how e.g. NVIDIA-targeting MMA APIs in other ecosystems typically model fragments as opaque,
backend-owned values rather than caller-visible scalar tuples — deferring lane-distribution
entirely to the backend, rather than baking a specific reconstruction strategy into a public,
2-element-return API contract that then can't be changed without a breaking API change.

A generalized sketch (this is a suggestion for Modular's own compiler team to evaluate, not a
concrete design): the same idea generalizes beyond MatrixFragment8x8[dtype] to a parameterized
type — e.g. CooperativeMatrixFragment[element_type, rows, cols, layout, scope] — where
rows/cols capture the fragment's logical tile shape (8x8 today, not hard-coded to it), layout
captures row-major vs. column-major / A-operand vs. B-operand orientation, and scope captures the
cooperative unit the fragment is distributed across (a single simdgroup on Apple GPUs today; the
same shape of type could describe e.g. an NVIDIA warp-level fragment without a new user-facing type
per backend). The only legal operations on such a value would be load_fragment, mma, and
store_fragment — no indexing, no construction from a raw SIMD vector, no way for user code to
inspect or assemble the per-lane layout at all; lane-distribution stays entirely backend-owned,
which is also what makes the backend free to choose the native instruction when one exists.

At the IR level, the natural complement of this user-facing type is a dedicated op family —
e.g. gpu.cooperative_matrix_load / gpu.cooperative_matrix_mma / gpu.cooperative_matrix_store
— that survives as a first-class, opaque operation through the optimizer pipeline all the way to
target lowering, instead of being desugared into generic load/shufflevector/intrinsic-call ops
early (which is what currently forces the manual-reconstruction pattern in §3, and is plausibly
also why pop.call_llvm_intrinsic's legalization in §1-2 has no path for a pointer operand on this
specific intrinsic — it is being pushed through a generic lowering path with no
cooperative-matrix-specific awareness). On Apple Silicon, this op would lower to a target-specific
gpu.apple.simdgroup_matrix_load/_mma/_store (a thin wrapper around the
air.simdgroup_matrix_8x8_* intrinsics this issue is about); on any backend without a native
cooperative-matrix instruction, it would lower to exactly the manual load-plus-shufflevector
sequence Mojo already generates today (§3) — i.e. the portable fallback is not new work, it is the
status quo, just moved one layer down, so a target that does have the native instruction is not
stuck paying the fallback's cost. Worth flagging honestly: §5b's own A/B measurement found that a
naive substitution of the native instruction alone did not win for the one kernel/shape we measured
it on, so the case for this direction should rest on portability/parity/future-shape value, not on
an assumed performance win.

B. A smaller, more targeted fix: independent of (A), simply making pop.call_llvm_intrinsic's
legalization pass accept a pointer-typed operand for llvm.air.simdgroup_matrix_8x8_load
specifically would at least unblock the existing low-level escape hatch (llvm_intrinsic) for
users who want to bind this instruction themselves today, without requiring any new public API
surface. Given that the compiler already has some specific, named awareness of this intrinsic
(§2's SIGSEGV-vs-clean-diagnostic distinction between a fake name and the real one, under the
identical pointer-plus-vectors argument shape), this may be a narrower, more contained fix than
(A) — though we have no visibility into how contained it actually is inside the legalization pass
itself.

Stated more generally, since this is the underlying ask rather than a fix scoped to one intrinsic
name: pop.call_llvm_intrinsic's operand-legalization should accept pointer-typed operands for
pointer-bearing intrinsics broadly, not only llvm.air.simdgroup_matrix_8x8_load. The
SIGSEGV-vs-clean-diagnostic distinction in §2 suggests today's legalization pass has per-intrinsic-
name awareness rather than a general "pointer operand is legal" rule, so any other pointer-bearing
AIR (or other target) intrinsic a user might want to bind through this same llvm_intrinsic escape
hatch would likely hit the identical rejection — this fix, if made general rather than
single-intrinsic-specific, would unblock that whole class at once.

7. Environment

Captured during the original investigation and independently reconfirmed while preparing this
report:

OS	macOS 26.5.2 (Build 25F84)
Xcode	26.6 (Build 17F113), at /Applications/Xcode-26.6.0.app
GPU	Apple M4 Max, integrated GPU, Metal backend (40 GPU cores, confirmed via system_profiler SPDisplaysDataType)
Host CPU	Apple M4 Max, 16 physical cores (12 performance + 4 efficiency)
Mojo	1.1.0.dev2026080606
MAX	26.6.0.dev2026080606
Package channel	https://conda.modular.com/max-nightly/ (+ conda-forge), via pixi
Build/run invocation	DEVELOPER_DIR=/Applications/Xcode-26.6.0.app/Contents/Developer pixi run mojo run <file>.mojo

This is a pinned nightly toolchain, chosen (in the original investigation) specifically because it
was empirically verified to run the Apple 8x8 simdgroup_matrix MMA primitive correctly on this
machine without a separate compiler crash an older, more widely-used pinned toolchain
(mojo==0.26.3, used elsewhere in the same codebase) hits for this primitive. We have not tested
whether the specific failure in this issue reproduces on other Mojo/MAX versions.

---

Verification / how this package was tested

Both .mojo files in this directory were re-run, independently of the original investigation, in
a brand-new, isolated pixi workspace created specifically for this report (not the source
repository's own mma_forge/ pixi environment, and with no files present in the directory besides
the two .mojo files and a fresh pixi.toml pinning the exact same mojo/max versions listed
in §7):

• repro_manual_load.mojo (§3, the "works" contrast case): compiled and ran successfully, real GPU
  dispatch, and reproduced the exact same correctness figures already on record from the
  source investigation: F32 fragment probe worst delta: 1.2914339464487057e-07, F16 fragment probe worst delta: 0.0005017651451950655, MMA FRAGMENT LOAD PROBE: PASS.
• repro_native_intrinsic_fails.mojo (§1-2, the bug): failed to compile with the exact error
  text quoted in §2 above, freshly captured from this isolated run (not copy-pasted from the
  original investigation's own transcript).

Honest limitation of this verification: this machine has no plain, non-pixi, system-wide
Mojo/MAX installation to test against — pixi/conda is, as far as we are aware, the standard
distribution mechanism for Mojo/MAX nightlies at this version, so a fresh pixi workspace pinned
to the same version, with zero other project files present, is the closest available
approximation of "a Modular engineer with a fresh MAX install" reachable from this environment. It
is a genuinely independent re-solve and re-install (its own .pixi/ environment, no shared state
with the source repository's own environment beyond the local package cache), not merely a re-run
inside the original project. We were not able to test on a second machine or a non-macOS/non-Apple-
Silicon host — this issue is specific to the Apple GPU backend (air.*-namespaced AIR intrinsics),
so cross-platform testing would not be meaningful for it regardless.

Both .mojo files in this directory are, to the best of our verification, genuinely self-contained
and dependency-free from the source repository's own modules (mma_forge, mma_native_forge,
attention_forge) — confirmed not just by inspection but by successfully building and running them
in a directory containing no other source files at all.

Start by running repro_native_intrinsic_fails.mojo and compare it with repro_manual_load.mojo. Read std/sys/intrinsics.mojo at line 77 and the reported pop.call_llvm_intrinsic legalization diagnostics. Done means the pointer-plus-vector ABI for air.simdgroup_matrix_8x8_load passes legalization and the reproduction compiles and runs.