Metal Backend Optimizations

Goal: match llama.cpp (Apple M4 Pro) performance. Current gap (2026-10-07 same-model, same-parameter A/B against llama.cpp c479922ac): decode is ~0.88-0.99× llama (7B Q4_K_M ≈ parity — 48.33 vs 48.59 t/s; 0.5B Q4_K_M 0.88×; Qwen3-0.6B Q8_0 0.89×) and long prefill is ~0.86-0.93× (0.5B Q4_K_M 0.86×; 7B Q4_K_M 0.93×; Qwen3-0.6B Q8_0 0.86×). The full protocol and every raw number are in §0.2; measured at minfer 97cc0d7 on macbook (macOS 27.0.1, Apple M4 Pro). The 2026-08-17 record (llama.cpp 88b47a755) is superseded and kept as history in §1.1. The re-measured minfer-only rows (throughput on the cached models + Metal-vs-CPU parity) live in §0.1, taken 2026-10-06 at 6b95763. See §1 Current state.

⚠️ The §0 progress table is the single source of truth for tracking; §1-§6 are the detailed explanations behind it. Update §0 first before changing code.

⚠️ Commit hashes in this document predate a repository history rewrite and do not resolve. Measured on 2026-10-10: 44 of the 47 hash-like tokens here fail git cat-file -t, and the three that pass are not this repository's commits anyway (llama.cpp's own c479922ac and 88b47a755, and an AppleClang version). They are kept as history, not as citations. Subject-matched equivalents from the current history — each verified with git cat-file -t — are listed in METAL-BACKEND-DESIGN.md §5.2; cite the equivalent and that section together, and treat a bare hash in this file as a pre-rewrite label rather than a locator.

📌 Compute-graph era note (2026-08-21+, Phase 6): inference now runs through the declarative compute graph (src/graph/, MetalBackend), which dispatches these kernels per op instead of the old whole-layer layer_gpu. Every kernel/optimization below remains valid — the gap is in wiring, not kernels: see §0.1 Graph-path (MetalBackend) integration status. The old-path measurements in §1 are pre-graph references; graph-path numbers live in §0.1.

📖 Backend design of record: METAL-BACKEND-DESIGN.md (MetalBackend lifecycle, dispatch, command buffers, model wiring, safety). This document stays the optimization ledger; commit hashes printed below predate a history rewrite and no longer resolve — see the design doc's §5.2 for subject-matched equivalents.

§0 Progress Overview (single tracking source)

Legend: [x] done (with commit) · [ ] to-do · [—] decided not to change. Commits are from git log; a few early items are marked "TODO-trace" (long history, to be filled in later).

§0.1 Graph-path (MetalBackend) integration status (2026-08-21)

The graph's MetalBackend (src/graph/metal_backend.rs) dispatches one op per call into a split-scoped MpsCommandBuffer. Which of the optimizations below already transfer:

Graph-path measured numbers (M4 Pro, --temp 0 greedy; old-path figures from §1/§1.6 — same models; the graph-path column was re-measured 2026-10-06 at 6b95763 on macbook (macOS 27.0.1, Apple M4 Pro), minfer bench -p <P> -n 128 -r 3 <model> with the Metal default except the CPU row (MINFER_DISABLE_MPS=1 -p 0). The pre-round figures it replaces were taken before the Metal round; the Δ column is against the old path, not against the pre-round figure):

Parity (re-measured 2026-10-06). graph_metal_matches_llama_reference (Qwen3-0.6B) reproduces the pinned llama-Metal greedy prefix [12095, 13, 576, 6722, 315, 9625, 374, 1083, 279] — the one model-level oracle that compares Metal against an external reference. The per-op Metal-vs-CPU gates (metal_*_matches_cpu in src/graph/metal_backend/tests.rs) and the four kernel isolation suites are green in the full macOS run (532 / 0 / 43 unit, 21 / 0 / 6 integration; the +1 is #329's dispatch-refusal gate). The Metal-vs-CPU logits row cannot be taken from graph_metal_matches_cpu_logits: in the graph era both ModelDef::forward and forward_graph route through Qwen2Graph::forward (src/models/qwen2/mod.rs:56/:96), so its max |Δ| = 0 compares the Metal graph with itself — the test's "CPU reference" comment is stale. Filed as #324; no parity regression is implied. The f16/bf16 device-gate max |Δlogit| numbers are in docs/SUPPORT-MATRIX.md (their model files are not cached for a fresh run here).

The old ~5.9 t/s CPU row is not reproduced at 6b95763: the same 0.5B Q4_0 file decodes at 149.69 ± 1.79 t/s on the CPU with -p 0. A 25× gap is not measurement noise, so the old figure describes a different configuration (it is quoted from the pre-graph §1 table and may predate the CPU kernel work); it is left visible in the Old path column but must not be read as the current CPU number.

Reading: G1 (attention dispatch: flash/split/parallel) closed the decode KV-growth regression (−44 %/−32 % → parity; 0.5B 2.1×, 7B ~1.5×); G3 (n_out tail-row reduction) cut the full-nt last-layer FFN + lm_head work so 0.5B prefill now exceeds the old path. Remaining graph-path gap: 7B prefill is now well ahead of the old path too; the residual llama gap (attention is not the bottleneck there — GEMM dominates and transfers fully) is the 2026-08-17 figure above and was not re-measured. Greedy outputs are byte-identical to the pre-G1 graph path (flash/split/parallel and tail-reduced paths all verified).

Windowed prefill (#315 / #359 / #369). The explicit prefill — the batched multi-sequence path — first ran the correctness kernel families kernel_gqa_attn_window_f32/_f16 (one-range attn_span) and kernel_gqa_attn_map_f32/_f16 (set-valued kv_map) where a single-sequence prefill keeps the tuned attn_flash_prefill (src/metal/kernels/attn_window.metal). Measured on this Mac (2026-10-07, macbook (macOS 27.0.1, Apple M4 Pro), 512 tokens, bar 0.8x named before the run) the one-range path reached only 0.539x the causal prefill's tokens/s for a two-sequence batch on the 0.5B Q4_0 (f32, hd 64) and 0.159x on Qwen3-0.6B Q8_0 (f16, hd 128); at the same shape 0.376x (~2.7x slower) and 0.096x (~10.4x slower), and the map path (a shared prefix read in place) 0.229x / 0.046x. [#359] added the fast one-range family kernel_flash_attn_window_blk_* (src/metal/kernels/fa_window.metal, a copy of the causal attn_flash_prefill tile with an explicit [lo, hi) mask, selected for nt > 1 at hd ∈ {64,128}) and re-measured 0.968x / 0.918x primary and 0.996x / 0.995x shape-matched; [#369] added the fast map family kernel_flash_attn_window_map_* (the same tile with a run-membership mask, fwin_map_has) and measured the map arm at 0.943x / 0.914x — all above the 0.8x bar, with MINFER_NO_WINDOW_FLASH=1 restoring the before numbers on both layouts as the A/B control. A map whose runs are far apart pays the gap between them (the kernel tiles the global union); the first-fit shared prefix the server emits is adjacent and does not. Protocol, node/kernel counts and the bar are in METAL-BACKEND-DESIGN.md §4.4.1.

Graph-path TODO (wire into MetalBackend): ① ✅ G1 attention dispatch; ② ✅ G2 rms_norm_256; ③ ✅ G3 n_out tail-row GetRows; ④ ✅ G4 fused decode QKV + bias+rope+store; ⑤ ✅ fused decode FFN gate+up (Op::FusedFFN); ⑥ ✅ G6 Qwen3 fused decode QKV with per-head Q/K RMSNorm (Op::FusedQkvNorm, 2026-08-27) — decode 3 matmul + 2 qk_norm + 2 rope + 2 store → 1 concat matmul

• 2 per-head norm + 1 no-bias rope+store. G1–G6 are wiring (kernels already isolated-tested); G3 additionally fixed two allocator liveness bugs (liveness used topo_order while the scheduler executes in build order; input buffers were freed and reused before host fills finished). The FFN fusion is gated on nf ≤ 16384 — the Q4_K concat matmul (od ≈ 37888) is slower than two separate matmuls on the 7B decode scalar kernel, so only the 0.5B class fuses FFN.

§0.2 Refreshed llama.cpp A/B (2026-10-07)

The §1.1 (2026-08-14/08-17) A/B against llama.cpp 88b47a755 was re-measured on the Mac against a HEAD llama.cpp build, c479922ac (build 11458, AppleClang 21.0.0.21000334, CMake Release -O3 -DNDEBUG; CMAKE_C_FLAGS_RELEASE additionally -ffp-contract=fast), at minfer 97cc0d7. Build: cmake --build ~/git/reading/llama.cpp/build-fpc --target llama-bench llama-cli -j. This is outcome 1 of #331: the pair is refreshed, not merely re-worded.

Protocol (gate-contract rules 3 and 5). Same GGUF file for both engines, from /Volumes/WD_BLACK/models/hf/Qwen/; same model, same prompt length and context for both. minfer:

./target/release/minfer bench -p P -n 128 -r 3 <model>                       # Metal default
MINFER_CACHE_TYPE=f16 ./target/release/minfer bench -p 430 -n 128 -r 3 <Q8_0 model>

llama.cpp (the 2026-08-17 pair used minfer --greedy vs llama-bench -b 512 -t 8; kept here — minfer bench is the current harness: pp times the prefill alone, and tg runs P prefill tokens and then times only the n_gen decode steps (the prefill happens before the timer in src/bench.rs's decode_once), which is llama-bench's tg caliber — the direct analogue of the old Generated: pure-decode number, not a prefill-inclusive average; both harnesses do a warmup pass before the measured reps and report a mean over -r 3):

~/git/reading/llama.cpp/build-fpc/bin/llama-bench -m <model> -p P -n 128 -r 3 -b 512 -t 8

Ambient state. macbook (macOS 27.0.1, Apple M4 Pro), hostname macbookpro-ysw, AC power (battery 80 %, not charging), no recorded thermal or performance warning; the box was otherwise loaded (Microsoft Edge + opencode/agent active, 15-min load average ≈ 7). Both engines ran back-to-back in one session, so the ratio is the robust quantity; each engine's absolute t/s carries the shared ambient load. Each cell is the harness mean ± stddev over the 3 measured reps.

Qwen2.5-0.5B-Instruct Q4_K_M (qwen2.5-0.5b-instruct-q4_k_m.gguf):

Qwen2.5-7B-Instruct Q4_K_M (split GGUF, entry part …-00001-of-00002.gguf):

Qwen3-0.6B Q8_0 (f16 KV cache both sides; minfer MINFER_CACHE_TYPE=f16, llama-bench default -ctk f16 -ctv f16):

For the Q4_K_M rows minfer's auto KV rule keeps the 0.5B class on f32 and the 7B class on f16 while llama-bench is f16 for both; this is stated, not corrected (the 2026-08-17 record had the same asymmetry).

Reading. The 2026-08-17 "long prefill ~2.6-2.7×" no longer holds: long prefill is now 0.86-0.93× (1.07-1.16× slower) and decode is 0.88-0.99× (7B at parity). Short-prompt pp30 stays dispatch/overhead-bound on both engines and is not a clean attention lever (0.72-0.85×).

§0.3 Small-batch prefill re-measurement (#40, 2026-10-09)

#40 closed with "the nt∈[2,8] small-od gap vs llama's ext kernel (0.5B pp4 141 vs 583) remains open — low value (2-8 token prompts), deferred": a 0.24× A/B taken 2026-08-21 against a then-HEAD llama.cpp. That figure is stale, and a reader of §0 could not tell the deferral from an oversight. Re-measured 2026-10-09 at minfer 270b6b3 on macbook (macOS 27.0.1, Apple M4 Pro), hostname macbookpro-ysw, AC power (~79 %), the box otherwise loaded (Edge + agent; 15-min load ≈ 7-8), against the §0.2 llama.cpp reference c479922ac (build 11458, AppleClang 21.0.0.21000334, -O3 -DNDEBUG -ffp-contract=fast).

Protocol. Same GGUF for both engines, from /Volumes/WD_BLACK/models/hf/Qwen/Qwen2.5-0.5B-Instruct-GGUF/: qwen2.5-0.5b-instruct-q4_k_m.gguf (the §0.2 model) and …q4_0.gguf. §0 #40 recorded the pp point ("pp4") but not the tg, so the §0.2 harness is used whole. Each engine ran its own contiguous sweep (minfer first, then llama) from a settled box, -r 5:

./target/release/minfer bench -p P -n 128 -r 5 <model>                      # Metal default
~/git/reading/llama.cpp/build-fpc/bin/llama-bench -m <model> -p P -n 128 -r 5 -b 512 -t 8

Reading. The recorded 0.24× is stale: minfer's pp4 roughly doubled since #40 and the ratio at pp2-8 is now ~0.3-0.65×. Repeats spread that band — the tiny-pp timing is a few ms of per-call setup, so the session's ambient dominates (llama's own pp30 moved 2845 → 2137 between §0.2 and this run; each engine's pp4 ratio across four sessions fell in 0.33-0.52 for k_m and 0.52-0.65 for q4_0). At pp30 the engines are within ~±15 %. The gap is real and lives exactly where #40 left it: the nt∈[2,8] small-od matmuls still take the serialized _multi kernel where llama runs kernel_mul_mv_ext (the batched matvec), because the GEMM gate is nt >= 2 && (od >= 2048 || nt >= 9) (src/metal/ops.rs:40), so a 0.5B's attn_qkv (od 1152) stays on the multi path through nt 8; the GEMM does kick in at nt >= 9, which is why the ratio improves through pp12-16.

Decision — kept deferred (2026-10-09). The fix is a mul_mv_ext-style batched-matvec kernel family (comparable surface to the q4_K decode port of #27) that speeds up only 2-8 token prompts, a window real chat/system prompts never occupy: the engines are already at 0.86-0.93× at pp30-495 (§0.2). The cost buys a few milliseconds of one-time prompt processing for single-word-style inputs and does not justify the kernel surface. Revisit trigger: a workload that actually runs 2-8 token prompts at volume (single-token completions, an embedding/similarity path, or speculative-draft prefill), or a batched-matvec kernel landing anyway for a larger reason. The work is filed as #431 so the finding is schedulable and searchable; the trigger above still decides when it is worth starting. Until then this is a decision, not an oversight.

✅ Done

Completed optimizations in detail (every done / decided-not item): §3.1 Correctness fixes · §3.2 CPU sampler · §3.3 Decode optimizations · §3.4 Prefill optimizations · §3.5 KV / long context · §3.6 Prefill GEMM gap investigation (resolved)

🔜 To-do (required path to match llama.cpp)

Principle (2026-08-12): we do NOT accept the current state — whatever llama.cpp can achieve, minfer must too. The former "accept the architecture floor" verdict is revoked; §4 is the only action path.

Note: §3.3's per-kernel table is superseded by the clean isolation A/B (§3.3): the trace mixed fused-vs-separate and different od per kernel name. The reliable decode gap = q6_K ffn_down kernel saturation (now fixed).

❌ Decided not to change (has measured or llama-source evidence)

1. Current state

1.1 Historical same-model, same-parameter A/B (2026-08-14 record, llama.cpp 88b47a755)

Historical record — superseded. These numbers are the 2026-08-14/08-17 A/B against llama.cpp 88b47a755; they were re-measured 2026-10-07 against c479922ac in §0.2. Kept as history.

minfer --greedy (pure decode, llama "Generation" caliber); llama.cpp llama-bench -b 512 -t 8 (pure eval). Model Qwen2.5-0.5B-Instruct. Prefill numbers updated after the prefill flash port (5974eb1): minfer pp430-435 uses the kernel_flash_attn_blk_f32 path (hd==64).

Q4_K_M (qwen2.5-0.5b-instruct-q4_k_m.gguf):

Q4_0 (qwen2.5-0.5b-instruct-q4_0.gguf):

Reading:

• Decode is now 72-88 % of llama (pure GPU 1.1-1.4×, default sampling 1.25×) — driven by rms_norm_256, the chunk-cap/sync fixes, and the per-kernel non-matmul profile (was 1.47× before 2026-08-10).
• Prefill (long) improved from 2.8-3.6× to ~2.6-2.7× after the prefill flash port (§3.4): GPU ~164 → ~144 ms at pp435 (~12 %; pp294 ~16 %). The residual gap is the non-attention 89 → ~44 ms structural difference (GEMMs + small kernels under-occupied, §3.6), NOT attention (now ~3-4 ms, llama-like).
• Short prefill (pp30): dominated by per-dispatch fixed overhead (~950 t/s at pp31 after #39 — the mmap first-access TLB cost is now absorbed at load; was ~430 t/s), so pp30 is no longer a meaningful attention lever; llama's pp30 is similarly launch-bound (0.5B 2686, 7B 328 t/s — minfer now 2.4×/1.3×).

Post-#28-32 (2026-08-21): the §1.1 0.5B numbers remain the reference — the GEMM-partial-tile barrier fix (#28) was a correctness fix (no regression), the unroll/barrier changes (#29/#30) improved 0.5B prefill ~7-9 % on top, and the lm_head output-rows-only fix (#32, §3.7) is a prefill/lm_head change that does not alter the 0.5B decode reference here. No re-measurement has invalidated any §1.1 figure. The user-facing 7B state is the current focus — see §1.6.

1.2 Per-token GPU decomposition (decode, nt==1, Q4_K_M 0.5B)

SUPERSEDED 2026-08-13 by the per-kernel trace (§3.3): the llama-side numbers here were inferred (no per-op timing existed). The trace shows the real picture differs: matmuls are NOT "zero gap" (minfer 1.6-3.9× slower per call), and small elementwise is NOT "4× slower" (parity). The Total is right; the category split was not.

1.3 Whole-pipeline comparison (decode token, nt==1)

1.4 Per-layer kernel sequence (Q4_K_M 0.5B, nt==1)

minfer — 20/layer (fused QKV OFF: Q5_0/Q5_0/Q8_0 mixed types, cannot concat):

RMSNorm → Wq/Wk/Wv 3×matmul → 3×add_bias → 2×RoPE → 2×KV store → attention split(partial+combine) → Wo matmul → residual → RMSNorm → fused gate+up matmul → SwiGLU → Ffn_down matmul → residual

×24 = 480 + output_norm 3 = 483. Q4_0 model (all Q4_0): fused QKV + BSR active → 12/layer ×24 + output 3 + GPU embed 1 = 292.

llama.cpp — 17/layer (flash_attn on):

RMSNorm → Wq/Wk/Wv 3×matmul → 2×RoPE → 2×KV store(f32→f16) → flash attention(1 dispatch) → Wo → residual → RMSNorm → gate+up 2×matmul → SwiGLU → Ffn_down → residual

×24 = 408 + output 3 + embed 1 ≈ 412 base; graph 822 nodes → ~490-530 dispatches (f16 cast/cont/reshape are non-no-op nodes).

1.5 Early performance milestones (Qwen2-0.5B, historical)

(Early numbers used the blended caliber; since 2026-08-06 Generated: is pure decode.)

1.6 7B state (2026-08-21 record, user-facing model, Q4_K_M)

Historical record. The llama-comparison gaps below are the 2026-08-21 standing; the refreshed 2026-10-07 A/B is in §0.2.

The 7B (qwen2.5-7b-instruct-q4_k_m, split GGUF) is the user-facing model; this is the current standing after the decode q4_K port (#27), the correctness/unroll/ barrier work (#28-30), the Phase-0 prefill decomposition (#31, §3.6), the lm_head output-rows-only fix (#32, §3.7), and the GPU Q4_K embedding port (#33).

• Decode is at parity (essentially closed by #27's q4_K decode matmul port).
• Prefill: the big lm_head over-count (minfer computed [152064×495] logits, llama computes [152064×1] after get_rows(inp_out_ids)) is FIXED (#32) — GPU ~1354 → ~1255 ms, download ~150 → ~0.1 ms. The corrected FLOP accounting (llama ≈ 6.26 TFLOP, NOT 7.0) shows llama ≈ 6.0 TF ≈ the replay's 6.2 TF.
• The residual ~1.2× prefill gap is kernel/graph-level, below static visibility (GPU backend machine-code scheduling / execution environment) and accepted (§3.6). The only remaining research direction is the tensor-API GEMM (mpp::tensor_ops) — llama disables it on M4 by default, so it is a beat-llama option, not a parity fix. The ~40 ms follow-up (last-layer FFN on output rows only, §3.7) is now DONE (#34) — minfer's total graph work (≈6.26 TFLOP) exactly matches llama's, and the pp495 gap closed 1.2× → ~1.16-1.18×.

2. Gap analysis: verified vs inferred

Core finding (2026-08-06 #5, verbatim): "Structural" is an inference, not a proven architectural inferiority. §2.1 is what is strictly VERIFIED; §2.2 is the inference after elimination. The decisive per-kernel measurement was completed 2026-08-13 (§3.3) — see its per-kernel table, which refuted several of the §2.2-§2.4 inferences.

2.1 Strictly verified

1. matmul kernel source is line-for-line identical (nt==1 mul_vec_q_n_f32_impl / block_q*_dot_y translations).
2. dispatch count is comparable (~436 vs ~490-530).
3. dispatch params match (llama ggml-metal-impl.h N_R0/N_SG vs minfer, 2026-08-06 #6).
4. dispatch count is nearly identical (~484 vs ~490-530); the gap is per-dispatch GPU execution time (10.3 µs vs 6.2 µs).
5. prefill GEMM ceiling ~5.4 TFLOPs/s (prefill_gemm_throughput_profile, 2026-08-11 A1); llama ~7 TFLOPs/s effective.

2.2 Closed hypotheses (2026-08-06 #6)

1. Dispatch params — DISPROVEN (see 2.1.3).
2. Attention is the main decode lever — DISPROVEN: llama -fa on/off = 3.64 vs 3.88 ms (only ~0.25 ms); even flash OFF llama beats minfer.
3. Multi-cb is a lever — DISPROVEN: minfer encode is only 0.13 ms; MINFER_SPLIT_CB=N regresses linearly (0.67 → 0.93/1.23/1.62 s).
4. CPU side (encode/sampler) — DISPROVEN: encode 0.13 ms; sampler fixed (2×).
5. Q5_0 scalar dequant is the matmul bottleneck — DISPROVEN: full vectorization only +2.6 %; ~130 GB/s is nt==1 small-grid structural latency.
6. Matmul is the prefill bottleneck (before the attention fix) — 2026-08-11 proved classic attention ~100 ms (48 %) was; matmuls became the main remainder after the fix.

2.3 Per-kernel non-matmul profile (2026-08-10, P0)

src/metal/::tests::non_matmul_bandwidth_profile (batched-cb, median of 3) — a single dispatch is dominated by the ~165 µs cb launch+sync floor, so batch dozens and take the median:

Findings: ① the attention split pair is the dominant non-matmul kernel (44 µs/layer); only a faithful flash port can cut it. ② the small elementwise tail (~300 kernels × 2-3 µs) is structurally cheap per-dispatch latency; rms_norm was the one exception and is now fixed.

2.4 Final gap report (2026-08-06, precise decomposition) — SUPERSEDED

The structural gap is 100 % in non-matmul — minfer's ~340 small kernels run at ~4× llama's efficiency. (Honest uncertainty: minfer's 1.2 ms is reliable; the attention-vs-small split has ±0.2 ms noise; llama's 0.3 ms is inferred.)

SUPERSEDED 2026-08-13 by §3.3's per-kernel trace. The "~0 gap matmuls" and "small ~4×" were subtraction artifacts. Real per-kernel data: matmuls ARE the gap (1.6-3.9× per call) and small-op is at parity. This section is kept as the historical pinpoint that motivated the trace.

2.5 KV-growth component (partially fixed 2026-08-10)

minfer's average decode grows with context (5.05 → 6.7 ms at -n 64→512), llama's stays flat. Two addressable causes were fixed: attention chunk cap 32→16 (avoid over-parallelization) + removal of sync_kv_to_cpu on the pure-GPU paths (an O(nkv)/token copy). Interleaved A/B: 4.65-4.76 → 4.50-4.55 ms/token (~0.2-0.25 ms). The remainder is sub-linear attention KV-read, which llama amortizes natively via its f16 cache + flash.

3. Completed optimizations in detail

3.1 Correctness fixes (Metal backend foundation)

4 early bugs (all affected output correctness, Qwen2-0.5B):

1. RoPE freq_scale not applied (src/metal/ rope_f32 got the param + forward.rs passes hp.rope_freq_scale).
2. output_b not applied (output_norm_gpu adds the bias).
3. softmax max initialization (-INFINITY instead of 0, prevents NaN).
4. attention stack array hardcoded (hd dimension made dynamic).

Q5_K formula + qh index fix (2026-07-31, affects CPU + Metal):

• Formula: Q5_0-style signed dl*(u-16)-ml was wrong → llama's unsigned dl*u-ml.
• qh high-bit index: qh[sub*4+pos/8] bit pos%8 was wrong → qh[pos] bit sub.
• Fixed in quants.rs / kernel.rs / forward.rs (embed).

GQA attention simd_max divergence fix (2026-08-01, 28d4ba2): in a partial KV tile (nkv % 32 != 0), out-of-range lanes exited the loop early → simd_max(dot) ran across divergent lanes with stale registers → corrupted online-softmax running max → repetition loops. Fix: uniform iteration count + valid mask (invalid lanes dot=-INF, e=0). Result: prefill logits cos 0.83→0.999. Regression test tests/gqa_attn_isolation.rs.

GPU-hang safety hardening (2026-08-03, bff73db):

1. submit() bounded 10 s wait + MTLCommandBufferStatus check + MINFER_TRACE dispatch trace (GPU fault errors out instead of freezing the machine).
2. attention kernels never return early before the barrier (prevents nh % nk != 0 deadlock).
3. layer_gpu/output_norm_gpu runtime guards: nh % nk == 0, hd ≤ 256, id % 32 == 0; error-exit (gpu_abort) on violation.

3.2 CPU sampler (2026-08-06)

Root cause: sampler.rs ran a full-vocab O(n·log n) sort per token + 607 KB copy (top_k) + a full sort of 151,936 (usize,f32) tuples (top_p). llama.cpp uses a candidate-list chain (std::partial_sort O(n·log k) → later samplers operate only on the ≤k survivors).

Fix:

• top_k → select_nth_unstable_by (O(n), on a copy to preserve the index→token mapping).
• top_p → softmax + sort only the ≤k survivors (falls back to the full-array path when >1024 survive).
• temp → skip exp() for masked (-INF) logits.
• main.rs moves the logits Vec instead of logits_all[..].to_vec() (607 KB/token).

Measured: -n 128/256/512 all ~2.0×; default sampling ~12.6-14.8 → ~5.5-6.5 ms/token; fixed-seed output byte-identical (7 sampler tests pass).

3.3 Decode optimizations (GPU)

Split-attention design (b3d4c7a):

• Pass 1 kernel_gqa_attn_partial_f32/_f16: grid (nt, nk, n_chunks); each TG computes an online-softmax partial (mx, S, acc) for its KV chunk, same tile/barrier/valid-head structure as the classic kernel.
• Pass 2 kernel_gqa_attn_combine_f32: grid (nt, nh) merges the partials (pure elementwise, no shared mem/barriers).
• n_chunks = clamp((max_pos+1+31)/32, 1, 16) (lowered from /16..32 on 2026-08-10; MINFER_ATTN_CHUNKS overrides). Correctness is invariant to chunk count.

Fused QKV essentials (6f0c847; the "24 %" figure was a GPU-state artifact, corrected to ~5 % in 26b145b): Wq/Wk/Wv and ffn_gate/up are row-major-concatenated at load (concat_rows → blk.{i}.attn_qkv / ffn_gu) when types + input dim match. nt==1 runs ONE matmul/group; rope/store/swiglu read sections via set_buffer byte offsets. MINFER_NO_FUSE_QKV=1 A/B byte-identical; gemm_isolation.rs::qkv_row_concat_layout locks the layout.

Measurement-trap notes (avoid repeating):

• Single-dispatch isolation timing is unreliable (~165 µs cb launch+sync floor) — batch dozens and take the median.
• Each kernel's first timed run has a cold-start/GPU-clock-ramp artifact (~4×) — warm first, measure twice.
• Sustained benchmarking thermally throttles the M4 Pro (extreme: all configs ~1.3 s) — interleave configs, take min/median.

q6_K decode matmul layout port (2026-08-13, §3.3) — the decode gap was matmul-dominated, not attention/small-op (per-kernel trace §3.3). Clean isolation A/B (minfer matmul_bandwidth_profile vs llama test-backend-ops perf -b MTL0 -o MUL_MAT, same od/id/nt) showed q6_K ffn_down (896/4864) 72 → 217 GB/s (3.0× gap) — minfer's stride-64 super-block loop left only 19/64 TG threads busy (~30 % util) with scalar inner loops. Fix: ported llama's kernel_mul_mv_q6_K_f32_impl stride-2 + float4 layout (TG(32, nsg=2) dispatch):

• q6_K isolation 72 → 209 GB/s (llama 217); decode steady GPU 4.27 → 3.72 ms/token (~13 %); wall ~203 → ~220 t/s (Generated: pure-decode).
• byte-identical (git-stash A/B), all tests green. q5_0/q8_0 already at parity.

q4_K decode matmul layout port (2026-08-17, to-do #7, §3.3) — 7B K_M decode is Q4_K-dominated (attn_q/k/output + ffn_gate/up all q4_K; Q6_K only output/ffn_down/attn_v) and carried the same gap q6_K had. Rewrote kernel_q4_k_f32_matmul as a faithful kernel_mul_mv_q4_K_f32_impl transcription: stride-4 super-block loop, float4 acc over the kmask nibble unpack, sc16 scale unpack (verified byte-exact vs llama's get_scale_min_k4, ggml-quants.c dequantize_row_q4_K), TG(32, nsg=2), grid od/4:

• 7B isolation at decode dims: attn_q (3584/3584) 70 → 265 GB/s, attn_k (3584/512) 18 → 146 GB/s, ffn_gate/up (18944/3584) 74 → 243 GB/s.
• 7B steady decode GPU ~49.7 → ~19.3 ms/token (~2.6×); wall 17-19 → 45-49 t/s (llama 50.5). 0.5B/1.5B regression green, all 34 bin tests pass.

flash decode attention port (2026-08-14 hd=64 + 2026-08-17 hd=128, §3.3) — isolation confirmed attention at 42.8 µs/layer (partial+combine) vs llama's ~4-6 µs (flash vec) = ~7-10×, ~0.9 ms of the 3.72 ms decode step (~24 %). Root cause: llama's flash uses dot(float4,float4) + simd_shuffle_down reductions with no threadgroup barrier within a KV tile; minfer's split had 2 threadgroup_barrier per 32-row tile. Ported llama's kernel_flash_attn_ext_vec as kernel_flash_attn_ext_f32/_f16 (NSG=1, DK=DV=64, NE=2, C=32, NL=16, fixed-shape for Qwen2, writes the same {M,S,O} partials so the shared combine merges unchanged). KV-layout check PASSED (minfer [nkv][nk*hd] == llama physical layout — no cache rework). GPU-safety deviations: inline per-lane partial-chunk mask (llama's cross-lane sm[] write is a race), break-only control flow so all 32 lanes reach both barriers, clamped reads + -MINF masking instead of a pad buffer. Verified: tests/flash_attn_isolation.rs (cos vs CPU >0.999 for nkv 1..4097, cos=1.0 vs split through the shared combine), A/B byte-identical (0.5B Q4_K_M f32 + Q4_0 f16 + 7B Q4_K_M), decode GPU ~0.3-1.0 ms/token faster, wall ~10 %. hd=128 decode flash (7B, 2026-08-17): separate kernel_flash_attn_ext_hd128_f32/_f16 (NE=1, DK4=DV4=32, full simd_sum(mqk) per cc-iteration); 7B pp205 steady decode GPU flash ~51.1 vs split ~51.6 ms/token (weight-read-bound; attention was NOT the 7B decode bottleneck — the q4_K port below was).

3.4 Prefill optimizations

simdgroup GEMM (P0/P1, 2026-08-01): faithful llama legacy kernel_mul_mm port (64×32 tile, 4 simdgroups × 32 threads, Q4_0 dequant staged into sa, f32 activations into sb). P0's initial version had 3 bugs (B-staging unclamped rows, store transpose direction, barrier must be mem_threadgroup) → after P1 fixes: +11 % at 30 tok, +34 % at 70 tok. Dispatched for nt ≥ 16; MINFER_GEMM=0 falls back to f32 multi. Isolation test gemm_isolation.rs (nt=12/30/32/33).

Non-Q4_0 GEMMs (2026-08-03, c9f865c/2c03bd1/5b914f0): one simdgroup GEMM per quant — Q8_0/Q5_0/Q5_1 (32-elem blocks) + Q4_K/Q5_K/Q6_K (256-elem super-blocks). K_M prefill 300→650 t/s; 1.5B Q4_K_M 48→442 t/s (~9×). 8 KB threadgroup-memory guard. non_q4_0_gemm_isolation verifies.

Parallel prefill attention (2026-08-11, b2c97fd): classic kernel_gqa_attn_f32 is latency-bound at prefill (grid (nt,nk) sequential KV loop, ~24K barriers, ~100 ms = 48 % of prefill, ~25× llama's attention). Replaced by a 3-pass barrier-free design:

1. kernel_attn_scores: one 256-thread TG per (t,h) row; each thread computes one score.
2. kernel_softmax_attn: masked softmax over the kv axis.
3. kernel_attn_output: softmax·V sum.

GQA via per-head hk = h/gqa (the broadcast-GEMM idea was tried and abandoned — a 2D GEMM can't produce the per-head 3D scores tensor). ⚠️ threadgroup-memory bug: the softmax's shmem[tiisg] writes 32 floats but only 8 were allocated (OOB corrupted adjacent memory → NaN rows) — fixed to 32×4=128 B; rms_norm_256 had the same latent bug.

Measured: pp430 classic 212 → 144 ms (attention 100→30 ms); pp30 44→40 ms; 7B pp230 944→832 ms (attention 169→57 ms); 7B decode unchanged. 34 bin + 6 isolation tests pass, end-to-end byte-identical.

Prefill flash attention port (2026-08-14, §3.4, 5974eb1) — llama's prefill attention is a SINGLE fused kernel_flash_attn_ext_blk (simdgroup_matrix, NSG=4, Q=8, C=64) vs minfer's 3-pass (scores + softmax + output); measured minfer attention 46 ms of the 135 ms pp325 (34 %) vs llama's ~3 ms. Ported kernel_flash_attn_blk_f32/_f16 (fixed-shape NSG=4, DK=DV=64, 7168 B shmem, inline causal mask, kernel_kv_tail_pad for the partial last KV block). Isolation-verified (tests/flash_attn_blk_isolation.rs: cos vs CPU

0.999 across 16 nt/nkv configs incl. partial blocks + GQA, f32+f16, deterministic), A/B byte-identical to the classic gqa_attn_f32 at every layer (f32 AND f16 cache — maxabs 0.0), interleaved MINFER_TIMING prefill GPU ~110 → ~93 ms (~16 %). Bonus: fixes the f16-cache prefill 3-pass bug (the 3-pass kernels read the f16 KV cache as float* → garbage "!!!!!!"; the f16 blk kernel reads half K/V correctly). Default for hd==64 (0.5B/1.5B). Gate: MINFER_NO_PREFILL_FLASH=1.

hd=128 (7B) prefill flash port (2026-08-15, §3.4, 89ac39a) — 7B pp310 prefill GPU 1042 → ~949 ms (~9 %), f32/f16 byte-identical, fixes 7B f16-cache garbage. Default for hd==128 (7B) too.

GEMM partial-tile race + missing Metal memoryBarrier fix (2026-08-19, §3.4, 7253a3b) — (a) all 8 simdgroup mm kernels lacked the threadgroup_barrier BEFORE the partial-tile temp_str stores (temp_str overlaps sa/sb, so a fast simdgroup overwrites them while a slow one still reads → intermittently corrupted last-2-token logits, partial x-tile only); (b) the single prefill encoder had NO memoryBarrier between dispatches → RMSNorm write raced the QKV read of the reused bn buffer (last-2 token slots, huge stale values). Result: 1.5B/7B first-token nondeterminism (~10-30 % wrong tokens) → 24/24 deterministic, output matches CPU byte-for-byte.

mm-kernel hot-loop #pragma unroll (2026-08-19, §3.4, f3a499d) — llama FOR_UNROLLs the staging/ik/load/mac loops; minfer's 8 mm kernels had none. Added the 6 unroll points (llama-parity set): 7B pp495 1438.8 → 1355.6 ms (~5.8 %), 0.5B ~6.9 %, 1.5B ~2.4 %; byte-identical + 24/24 determinism.

ik-loop threadgroup_barrier → simdgroup_barrier(mem_none) (2026-08-19, §3.4 follow-up, 0e756f3) — the pre-unroll corruption was a rolled-loop compiler artifact, not a memory-visibility need (.air diff: the barrier swap changes ONLY the barrier instruction — zero IR scheduling difference). With the unroll in place llama's exact barrier form is now safe: 7B pp495 min 1387.4 → 1370.6 ms (~1.2 %), byte-identical (1.5B×24 / 7B×8 / 0.5B×3). Removes the last structural mm-kernel difference vs llama.

3.5 KV / long context

• KV geometric growth (66f4290): kv_ensure_layer grows ×2 instead of reallocating + copying the whole old KV every token (0.5 ms@KV140 → 4.2 ms@KV2510 → 0.13 ms). ⚠️ an old_v clone typo polluted the V cache (Q4_K_M garbage) — the A/B didn't catch it (both paths share the corrupted KV); found against a known-good reference.
• f16 KV auto-default (387d612 + bff73db + #37): MINFER_CACHE_TYPE=f16/f32 overrides; unset → set_kv_cache_type at load picks f16 for the 7B class (n_layers×n_kv_embd ≥ 8192) and f32 for small models (0.5B measured ~3 % slower on f16 — dispatch-latency-bound). kernel_store_kv_f16 + the _f16 flash / partial kernels serve the f16 path.
• Split-GGUF (cbba68c/34eaf10): multi-part models (7B -0000X-of-0000Y), merged tensor index, 7B verified.

3.6 Prefill GEMM gap investigation — RESOLVED (2026-08-14 → 08-20, decided not to change)

The prefill GEMM throughput gap was investigated exhaustively and is now closed as "not source-addressable" (the mm kernels are proven identical to llama's). Record of the investigation (was §4.3.1-§4.3.10):

Residual status (pre-#32, 2026-08-20): 7B pp495 minfer FULL ~1324-1336 ms vs llama ~1043-1064 ms (~73 %, ≈1.25-1.3×), converging under comparable system load (llama-bench degrades to 1250-1760 ms at load avg 3-4). Updated 2026-08-21 by #32 (§3.7): after the lm_head output-rows-only fix minfer is ~1255 ms GPU (≈82 %, ≈1.2×) with download ~0.1 ms — see §1.6 / §3.7. The only remaining research direction is the tensor-API GEMM (mpp::tensor_ops) — llama disables it on M4 by default, so it is a beat-llama option, not a parity fix (decided-not, §0).

⚠️ 2026-08-21 CORRECTION — the §3.6 FLOP accounting was WRONG. The "7.000 TFLOP" total assumed the output (lm_head) GEMM runs on all N=495 tokens. The graph dump actually shows out=[152064 1 1 1] for the output and the last layer's gate/up/down on N=1 — llama's ggml_get_rows(cur, inp_out_ids) (qwen2.cpp:106-108) reduces to n_outputs rows after the last attention, so the last layer's FFN + final norm + lm_head all run on N=1. Correct llama total ≈ 6.26 TFLOP → llama efficiency ≈ 6.0 TF (NOT 6.71), essentially equal to the replay's 6.2 TF. minfer was computing the full-nt output GEMM [152064×495] (≈539 GFLOP waste) + final norm on all rows + a 301 MB logits download — a large fraction of the measured "gap" was this over-count, not the kernels. Fixed 2026-08-21 (§3.7 / changelog #32).

3.7 lm_head / final-norm output-rows-only (2026-08-21, changelog #32)

Discovery (from the LLAMA_METAL_E2E.md reference): llama's graph reduces the hidden state to n_outputs rows right after the last layer's attention (ggml_get_rows(cur, inp_out_ids) + get_rows(inpSA, inp_out_ids), src/models/qwen2.cpp:106-108), so the last layer's FFN, the final norm and the lm_head all run on 1 row for a single-sequence prefill (graph dump: output GEMM out=[152064 1 1 1], last-layer gate/down [.. 1 ..]). minfer computed the output projection over all nt tokens ([nv×nt]) and downloaded all of it.

Fix: forward() / output_norm_gpu (Metal + CUDA) now take n_out (number of output rows = the LAST n_out tokens, single-sequence row-major [nt][ne] hidden). The final rms_norm + output GEMM + bias + logits buffer + download all operate on n_out rows (n_out=1 for the minfer CLI). Host: src/models/qwen2/forward.rs (GPU output_norm_gpu(…, n_out, …), CUDA path, CPU fallback slices hidden[(nt-n_out)*ne..]), src/metal/runtime.rs + rms_norm(.., off) byte-offset, src/main.rs passes n_out=1.

Measured (7B Q4_K_M pp495, interleaved A/B same window):

• GPU submit-wait: pre ~1354 (noisy 1354-2754) → post 1253-1271 ms (stable)
• logits download: ~150 ms (301 MB) → ~0.1 ms (608 KB)
• 0.5B GPU generation byte-identical pre/post (same seed, same text); 1.5B / 7B greedy generation correct.
• Corrected efficiency: llama ≈ 6.0 TF @ 6.26 TFLOP; minfer post-#32 ≈ 6.46 TFLOP @ ~1.255 s ≈ 5.15 TF — the remaining ~1.2× was the §3.6 kernel/graph residual plus the last layer's FFN still on all nt (llama reduces before it) ≈ ~40 ms follow-up potential.

Follow-up — DONE 2026-08-21 (#34): the last layer's FFN now runs on the output rows only, exactly matching llama's get_rows(inp_out_ids) reduction (qwen2.cpp:106-108 reduces cur + inpSA BEFORE the last layer's FFN). Fix in layer_gpu (src/metal/): the wo matmul stays on all nt (llama build_attn precedes the reduction), then the wo-residual, ffn_norm (byte-offset read of the hidden tail), gate/up/down (dispatched with nt=n_out via the new x_off matmul param) and the final residual all run on the tail n_out rows (add_f32_off); the CPU path in forward.rs mirrors (gate/up/down on n_out rows, residuals on the hidden tail slice).

• Measured (7B Q4_K_M pp499, interleaved A/B same window): GPU submit-wait ~1278 → ~1234 ms (~44 ms, ~3.4 %) — the ~40 ms estimate confirmed.
• minfer's total graph work drops 6.46 → ≈6.26 TFLOP — now exactly llama's total; efficiency ≈ 5.07 TF @ 6.26 TFLOP (llama ≈ 6.0 TF) → gap ~1.16-1.18×.
• Correctness: 7B/0.5B GPU greedy byte-identical pre/post (same seeds) + 0.5B CPU fallback identical; decode untouched (nt==1 ⇒ ffn_nt==nt, the reduced path never triggers); 33/34 bin tests green (attn_parallel_realdata_correctness fails on missing /tmp/dp3 dumps — pre-existing environment dependency).

4. Roadmap status — all items resolved (2026-08-20)

The 2026-08-12 principle — "we do not accept the current state; whatever llama.cpp can achieve, minfer must too" — drove the work below. Every roadmap item is now resolved (done, or decided-not-to-change with measured / source evidence). The detailed records live in §3 ("Completed optimizations in detail", incl. the §3.6 investigation record) — see the §3 links in §0's Progress Overview; decided-not items are in §0's "Decided not to change" table. This chapter holds the remaining research + operational reference, plus the cold-start to-dos added 2026-08-21 (§4.2) (a separate axis from the steady-state gap — not part of the §0 "match llama.cpp" table).

4.1 Remaining research (not a parity fix)

The only direction left to beat llama on prefill is the tensor-API GEMM (mpp::tensor_ops, kernel_mul_mm_id) — llama disables it on M4 Pro by default (PARAMETER_AUDIT A), so it is a research exploration, not a parity requirement. No other prefill path remains open (§3.6 / §0 decided #1).

2026-08-21 (#32, §3.7): the last-layer FFN output-rows-only follow-up (llama's get_rows(inp_out_ids) shrinks the graph before layer-27's FFN, closing the minfer-vs-llama prefill work difference 6.46 vs 6.26 TFLOP) is now DONE (#34) — no prefill work difference remains; minfer's total equals llama's.

2026-08-21 (#41, DEFERRED): quantized KV cache (llama #27390/#27438 — Q4_0/Q4_1/Q5_0/Q5_1/Q8_0 KV: prefill dequantizes to f16 scratch via kernel_flash_attn_ext_kv_f16 when ne01≥32; decode reads quantized blocks directly in the flash vec kernel). Analyzed against llama source (ops.cpp:2805- 2831, ggml-src/metal/kernels//7864-7885) and deferred: the shipped models cap at max_seq=32768 — f16 KV (#37) needs only 1.9 GB at 32K (within the 8.7 GB mmap-era RSS), so Q8_0 KV (0.96 GB) has no validation target on the current models. Revisit when a >32K-context model is the user-facing one.

4.2 Cold-start optimization to-dos (2026-08-21)

Observed: 7B run twice in a row → the second run ≈ 2× faster (Total 1.46 s → 0.77 s). Note the CLI's Total timing starts AFTER model load (main.rs:1442 infer_start), so the 2× is in the inference phases (prefill + decode), not the load. Root causes:

Warm steady-state = run 2's numbers (pp30 ~0.17 s, decode ~46 t/s). Not a bug.

Item 1 removes a genuinely avoidable per-invocation cost; item 2 is mostly a memory + one-copy-pass win (the 4.4 GB disk read is unavoidable either way). Item 3 is the only way to make every invocation fast, at the cost of a daemon. Together they target the cold-start axis only — steady-state inference is unaffected (already at the §1.6 numbers).

4.3 Graph-path (MetalBackend) integration to-dos (2026-08-21+, compute-graph era)

The compute graph is now the inference path (Phase 6). Kernel-level parity work is complete (§0); the remaining gap is wiring already-tested kernels into MetalBackend — see §0.1 for the status table and measured numbers.

4.4 Reference: GPU profiling tooling (done, operational)

• xctrace: /usr/bin/xctrace is a broken stub ("tool not found"); the real binary is /Applications/Xcode.app/Contents/Developer/usr/bin/xctrace.
• Record in Instruments: Metal System Trace, Counter Set → Performance Limiters, Enable Shader Timeline on, Deferred.
• scripts/export_trace.sh <trace> [run]: exports per-kernel durations (metal-shader-profiler-intervals), the limiter profile (gpu-counter-value), and per-forward intervals (metal-gpu-intervals); TRACE_PROC=<proc> filters per-forward intervals.
• Interpretation: percentage counters (Limiter/Utilization/Occupancy) are ratios; bandwidth counters (L1/LLC Read Bandwidth) are cumulative — ignore magnitude.

4.5 Process (was §4.5 backfill)

After each item completes: update the §0 progress table (check, fill in the commit, update measured effect) → record the implementation + verification in the relevant §3 section → update the §1 gap numbers.

5. Historical appendix (reference only)

⚠️ This appendix is a historical record, for reference only — NOT the basis for future plans. Its architecture conclusions, metric calibers, and kernel structures may have been superseded by §1-§4. Future plans follow §0 + §4 only.

5.1 Early phase summary (Qwen2-0.5B, 130→334 t/s, 2026-07-27~08-01)

5.2 Early gap-analysis records

• KV-growth 2.2× (2026-08-01): f32 KV full re-read vs llama f16 (fix in §2.5).
• Per-dispatch encode ~24µs vs llama ~7µs (2026-08-01 conclusion) → REVOKED 2026-08-03: encode measured at only ~1 ms/step; decode is GPU-execution-bound.
• Q4_0 dual dispatch (quantize+matmul) (2026-08-01) → fixed: f32 activations path (§3.1 #4).

5.3 Tested-and-rejected ideas (with commits)

5.4 Stale data tables (do not cite)

The following come from 2026-08-01/03 early measurements with different calibers (blended t/s vs pure decode; old llama baseline) — for historical cross-check only:

• "Q4_K_M/Q5_K_M prefill 7.3×" (35-token table) — before the non-Q4_0 GEMMs existed; now filled in.
• "decode short ~187 / long ~86 t/s" (KV-growth table) — before split attention.
• "pure decode 2.0× / 3.2×" early gap — now 1.1-1.4× (§1.1).

5.5 flash attention: llama.cpp vs minfer — detailed kernel comparison (reference for §3.3)

Line-by-line structural comparison behind the §3.3 ~7-10× attention gap (minfer split attention 42.8 µs/layer vs llama flash ~4-6 µs/layer at nkv=430). Source: llama kernel_flash_attn_ext_vec (ggml-src/metal/kernels/), llama dispatch (ggml-metal-ops.cpp:2959), minfer kernel_gqa_attn_partial_f32 (src/metal/kernels/kv.metal), _f16 (:3127), kernel_gqa_attn_combine_f32 (:3243).

A. Overall design

Key: the cross-TG reduction idea is identical on both sides (online softmax partials + a combine pass). All the difference is inside a single threadgroup.

B. Inside one threadgroup — the core gap

llama (dk64 template: NE=2, NL=16, C=32, 32 threads = 2 rows × 16 cols):

QK^T:  for cc in 0..C/NE:                 // 16 cache columns per simdgroup
         for ii in 0..DK4/NL:             // 16 float4 dots
           mqk[cc] += dot(pk4[..], pq4[..])
         reduce = simd_shuffle_down ×5 + simd_shuffle   // ← NO threadgroup barrier

• QK^T accumulates in registers mqk[]; the reduction is simd_shuffle_down (16 lanes of ONE simdgroup merge via pairwise shuffles — no threadgroup barrier needed).
• softmax update: M = simd_max, S = S*ms + simd_sum(vs) (same simd reduces).
• PV: lo[ii] += float4(pv4)*float4(sst) register accumulate + simd_shuffle_down reduce.
• only 2 simdgroup_barrier per tile (QK→softmax, softmax→PV), and simdgroup-level (much cheaper than threadgroup-level).
• KV read straight as half4 from global; relies on HW cache + barrier-free pipelining.

minfer (partial kernel, 224 threads/TG = 7 simdgroups × 32):

per 32-row tile:
  1. all threads stage the KV tile into threadgroup shmem (k_tile4/v_tile4)
  2. threadgroup_barrier                          // ① make data visible
  3. each lane does one row's QK^T:
       dot = 0; for d4: qv = qhead4[d4]*kj4[d4]; dot += qv.x+qv.y+qv.z+qv.w
       batch_mx = simd_max(dot); online-softmax (corr, acc4 *= corr)
  4. threadgroup_barrier                          // ② before reusing shmem

Gap sources:

• After every QK^T reduce and before every PV accumulate llama only needs simd-level shuffles; minfer must threadgroup_barrier — its 32 rows are processed serially by the same simdgroup's lanes (for j0 loop, tile_sz=32) while different lanes of the 224-thread TG handle different heads (gqa=7). Result: 2 global barriers per tile, 28/layer at nkv=430; each barrier stalls the whole TG to the slowest lane.
• llama's 16-column reduce stays inside one simdgroup (32 threads) — zero cross-simdgroup sync.
• minfer's acc4 is register-accumulated but serialized by the barrier sequence; llama's mqk/lo register accumulation pipelines through the shuffle chain.

C. Why the gap is ~7-10× and not smaller

Of minfer's 42.8 µs/layer (nkv=430), the 28 threadgroup barriers + per-tile shmem stage-in/out dominate: a barrier makes all 224 threads wait, while each 32×64 tile does only ~2048 MACs — swamped by sync overhead. llama's simd reductions drop sync to near zero, and half4 reads with no shmem staging keep the memory pipeline unbroken.

D. Port (option C) — actual change surface

For Qwen2.5-0.5B (hd=64, nh=14, nk=2, gqa=7, nt==1 decode):

1. KV-layout check (prerequisite) — DONE 2026-08-14, PASS: llama's K/V physical layout IS [nkv][nk*hd] (all KV heads packed in dim0 of the cache tensor, token stride nk*hd), and the flash kernel reads it with nb11 = nk*hd*elem (token), nb12 = hd*elem (head), ns10 = nk*hd — the same indexing minfer already uses (k + ki*stride_kv + hk*hd, stride_kv = nk*hd). No cache rework needed; the host passes the stride args mirroring minfer's layout (nb10 = elem, nb11 = nk*hd*elem, nb12 = hd*elem, ns10 = nk*hd). Evidence chain in §3.3.
2. Single-kernel rewrite — DONE 2026-08-14: fixed DK=DV=64, NE=2, C=32, NSG=1 (0.5B/7B decode dims), dropping llama's function-constant system (hardcoded constants; n_chunks is the host-tunable grid depth instead of llama's fixed NWG). Shmem sq4 + ss + so4 (no sm[] — mask inlined), QK^T/PV via float4 + simd_shuffle_down(8,4,2,1) + simd_shuffle(·, NL*ty) broadcast (the reduce routes the full-head sum to lanes 0/16).
3. Cross-TG reduce — DONE: reuse minfer's existing combine kernel — the flash kernel writes the same {M, S, O[hd]} partials (strided C=32-block chunking, interchangeable with the split's contiguous chunking).
4. isolation tests — DONE: tests/flash_attn_isolation.rs (scalar CPU ref, multi-nkv incl. partial/empty chunks, nt 1-2, f32+f16; flash-vs-split A/B through the shared combine) + end-to-end byte-identical A/B (MINFER_NO_FLASH=1).
5. prefill untouched: nt>1 keeps the current 3-pass parallel attention.

Risk points:

• KV layout mismatch was the most likely failure — CLEARED (2026-08-14, see D-1 above); the layout is compatible as-is.
• Function constants → hardcoded constants is NOT a line-by-line translation; shmem offsets must be re-derived (sgitg*SH terms are all 0 for NSG=1, which simplifies).
• simd_shuffle_down lane participation (NE=2 uses only the NE>1/NE>2 branches; simd_shuffle_down(mqk[cc], 16) folds 16 cols onto lane 0) needs careful index alignment.

5.6 Glossary — variables and functions used in this document

This table covers only this document's variable/function symbols. Concept-level terms (quant formats, kernel techniques, methodology) are classified corpus-wide in GLOSSARY.md; the ggml layout symbols below are also collected there (L2).

Common (both kernels, GGUF/model dims):

llama ggml tensor/strided-layout args (used in flash/matmul kernels):

Metal thread variables:

minfer kernels (src/metal/kernels/ unless noted):

minfer env vars: