#2987 metal: prompt chain without MetalPerformancePrimitives (macOS 15 SDK)

说明

### Why

On macOS 15.8 (Apple M2 Max, 96 GB) the engine cannot start the prompt chain at all: every `matmul2d` kernel in `metal/kernels/` is reached through an **unguarded** `#include <MetalPerformancePrimitives/MetalPerformancePrimitives.h>`, and that header only ships with a Metal 4 SDK. `Context::compile()` of `dense_mm.metal` / `moe_mm.metal` / `prompt.metal` / `qsa_prompt.metal` fails, the failure is fatal at init (`exit 3`), and the config has to carry `STRATA_METAL_PROMPT=window` — which caps prefill at ~58 tok/s on this box.

The GPU was never the limit. `simdgroup_matrix<half,8,8>` + `simdgroup_multiply_accumulate` is Apple7+ and has always worked here; the repo's own probe with the guards applied does **7.22 TFLOP/s** at 4096³ (rel L2 2.75e-07 against the fp32 reference) and prints `device Apple M2 Max, metal4 0`. So this is a toolchain question, not an instruction-set question — and it is the same root cause as #2860, one level down.

### What changed

11 files under `metal/`, +399/−28, no kernel semantics touched.

- `#if __has_include(<MetalPerformancePrimitives/...>)` around the MPP include, the `using namespace mpp::tensor_ops;`, and each `matmul2d` kernel body. Four sites get an `#else` twin — `gemm_f16` (dense_mm), `gemm_hc_up` (prompt), `mm_gu2` / `mm_down2` (moe_mm) — each a `simdgroup_matrix<half,8,8>` kernel with the **same signature and the same `[[host_name]]`**, so no host call site changes. The twins reuse the existing threadgroup staging (`stage_w<F, TN, TK>`) and the existing quant tables.
- QSA's three MPP kernels needed no twin: your chain already falls back to `qsa.metal`'s scalar `qsa_score`/`qsa_attn` when `coop_k4` is false, so 17 lines of bare `#if` are enough there.
- `mtl::Context::msl_has_mpp()` asks the **compiler** (one 5-line source that `#error`s when the header is missing) instead of `supports_metal4()`. The GPU family answers "Metal 4" on an M2 running a macOS 26 beta, but the toolchain is the question the kernel's own `#if` answers — keying host and MSL off the same oracle is what keeps variant selection and the compiled library from ever disagreeing.
- `STRATA_METAL_LEGACY=1` forces `#define STRATA_MSL_MPP 0`, so a machine that has both can be measured arm against arm.
- `coop_k4()` (the M5 right-input layout probe, which compiles an MPP-only kernel) returns false without the header instead of aborting on a missing function.
- `MTLGPUFamilyMetal4` is guarded for toolchains that predate the enum.
- Tiles for the twins are measured, not guessed: `mm_gu2_*_16_32_32_2` for gate/up, TN 32 for the dense `gemm_f16` (see below).

### Correctness (all under `STRATA_METAL_LEGACY=1`, M2 Max / macOS 15.8)

| harness | result |
|---|---|
| `dense_mm_parity` | 67 PASS / 0 FAIL, worst rel L2 1.5e-06 vs fp16 operands |
| `prompt_parity` | ALL PASS (mixed 2.35e-04, inject 2.22e-04, swiglu_h 4.77e-04) |
| `gdn_parity` | ALL PASS, ntok 1–8 |
| `moe_mm_parity --v2` | PASS for every format this pack routes: IQ3_XXS 3.18e-05, IQ2_S 2.51e-05, IQ4_XS 2.61e-05, IQ3_S 2.28e-05, Q2_0 4.42e-06, Q8_0 2.55e-05 |
| same, upstream tile vs legacy tile | **identical error numbers** ⇒ the tile choice is speed-only |
| K-quant `16_16_128_2` under legacy | Q4_K 1.71e-05, Q5_K 9.57e-06, Q6_K 2.65e-05 PASS (speed unmeasurable here: this pack's routed experts carry no K-quant) |

### Numbers — one binary, one environment variable apart

1023-token prompts, five alternating runs each:

| arm | prefill tok/s | decode tok/s (64 tokens) | TTFT for 1023 tokens |
|---|---|---|---|
| `STRATA_METAL_PROMPT=window` | 58.58 / 57.90 / 57.42 / 58.86 / 59.10 | 26.81, 26.38 | 20.6 – 26.4 s |
| prompt chain on the simdgroup_matrix twins | 182.34 / 177.02 / 173.42 / 179.41 / 173.61 | 28.98, 27.27 | 8.8 – 14.4 s |

**3.02× on prefill** (medians 177.0 vs 58.6), decode unchanged to slightly better. Server-level over HTTP (`/v1/chat/completions`), a 4252-token prompt: 55.3 → 159.4 tok/s median. Short prompts (27–34 tokens) are unchanged because they stay under `_PROMPT_GEMM_MIN` = 64 and never reach the chain. The first prompt after boot is slower on both arms — one pass over the pack is cold storage: 84.3 tok/s cold, 177 once the file cache holds it.

Attribution on one 1024-token chunk (`prompt_chain --breakdown`), same session: upstream dense tile `64_64` = 58 tok/s (gdn proj alone 8.73 s) → dense twin at TN 32 = 140 (gdn proj 1.18 s, qsa proj 2.44 → 0.24, hc gemm 0.80 → 0.36) → expert gate/up `16_32_32_2` = 169 (gu 4.15 → 2.93). So the dense projections carry 2.4× and the expert tile 1.2× of the total 2.9×.

### What this does not claim

It does not overtake llama.cpp on the same weights. Same machine, same prompt texts, the **identical** GGUF file, `llama-server` b11408 (`-ngl 999`, `--parallel 1`, `-c 65536`, MTP draft):

| engine on the identical pack | 4220-token prompt | 1609 | 875 |
|---|---|---|---|
| window path | 58 | – | – |
| this branch | 160–172 | 167–172 | 153–162 |
| llama.cpp b11408 | 285.6 / 254.9 | 267.6 | 267.6 |

The branch recovers ~3× of this engine's own 4.9× deficit on the old SDK and still leaves ~1.6× on the table. The deficit is spread rather than concentrated — llama.cpp's whole 1024-token chunk costs ~3.6 s at that rate, while gate/up (2.93 s) plus gdn proj (1.21 s) here are already 4.14 s — so it points at how the chain schedules these kernels, not at raw `simdgroup_matrix` throughput (the dense twin runs at 6.89 of the 7.22 TFLOP/s the probe measures). Decode is a wash: llama.cpp's 1596 completed production requests median at 30.1 tok/s, this at 27–31.

### Merge state / what I would want reviewed

- **This branch is based on `296b59d2` (v0.1.44.1) and will conflict with `main`.** main is 264 commits ahead, 13 of them under `metal/`; `prompt_chain.inc:302` is the hunk that moved, and the `maxBufferLength` wrap (8b3ec44b, which bounds the MoE kernels per layer) touches `moe_mm.metal` where the twins live. I can rebase and re-run the gates — flagging it as a draft rather than pretending it merges clean. It is also worth noting that main currently does not build on a Mac at all (#2860, compile at `477ac13` / link at `a2de744`; #2895 looks like the fix), so the rebase would need that too.
- The `msl_has_mpp()` oracle + the `__has_include` guards are the part I think is worth merging even if the tile choices stay a follow-up; happy to split it that way if you prefer two reviews.
- The twins are ~350 lines of MSL that only exist for SDKs older than Metal 4 and are dead code on a 26 toolchain. If you would rather not carry them, the alternative is documenting that old-SDK Macs get the window path — but on this machine that is a 3× prefill cliff, and we would rather they lived upstream than in a local patch.
- Open question for you: should `mm_gu2_*_16_32_32_2` also become the MPP default? It is faster for the twins here, and the guard is only about which tile the legacy branch asks for.

Environment: macOS 15.8, Xcode Command Line Tools' macOS 15 SDK, Apple M2 Max (96 GB), engine 0.1.44.1-metal.