#2323 [Bug]: Mac (M2 Max, macOS 26.3): garbage output — gdn_step for 2+ tokens is never dispatched (pipeline max 448 threads < 512); qsa_prompt.metal fails to compile

Beschreibung

### What happened

Fresh `./setup.sh --model IQ3_S` on an M2 Max (main at `ed4773e`, engine 0.1.44-metal). The engine starts, but every reply is garbage (e.g. `"""ZEROxmlxmlxml...`, `#文哲ntaxntax...`) and MTP accepts 0 drafts. `STRATA_METAL_PROMPT=window`, `STRATA_HC_Q8=0`, `STRATA_METAL_FUSE=0` do not change it.

I tracked it down to two separate problems.

**1. `gdn_step` with `GDN_NTOK >= 2` does not run at all on M2 Max**

`gdn_parity` (built from `metal/tests`) fails on this Mac:

```
gdn_step pipeline: max 512 threads per threadgroup (the kernel needs 512)
gdn parity ntok 1: y rel L2 1.66e-07, ... PASS
gdn parity ntok 2: y rel L2 6.84e-01, state 1.83e+00, history 1.3e+00; ... partial commits bitwise: NO  FAIL
gdn parity ntok 3: y rel L2 8.02e-01, state 2.38e+00, history nan; ... FAIL
gdn parity ntok 4: y rel L2 8.49e-01, state 2.49e+00, history nan; ... FAIL
```

The test prints `max_threads` only for the `ntok = 1` specialization. Per specialization on the M2 Max:

```
GDN_NTOK 1: maxTotalThreadsPerThreadgroup 512
GDN_NTOK 2: 448
GDN_NTOK 3: 448
```

The kernel is dispatched with 512 threads, so the 2+ token variants are silently not executed: a marker written after the token loop stays 0, and there is no command-buffer error. Every verify window / prompt window longer than one token therefore runs on stale outputs. My guess is that M2 has a smaller register budget than M4/M5, where you tested.

Fix: force the compiler to fit 512 (or 1024 for `gdn_step8`) threads:

```diff
--- metal/kernels/gdn.metal
-kernel void gdn_step_t(device float* state [[buffer(0)]], device float* hist [[buffer(1)]],
+[[max_total_threads_per_threadgroup(1024)]] kernel void gdn_step_t(device float* state [[buffer(0)]], device float* hist [[buffer(1)]],
```

After this change, `gdn_parity` reports ALL PASS (ntok 1–4: y rel L2 ~1.7e-07, history exact, partial commits bitwise), and `gdn_prompt_parity` passes too. It would probably also be worth checking `pipeline.max_threads >= tg` at dispatch time and failing loudly instead of silently.

**2. `qsa_prompt.metal` does not compile with the macOS 26.3 MetalPerformancePrimitives headers**

The first longer prompt kills the engine (`exit(3)`) with:

```
mtl: compile failed:
.../metal_cooperative_tensor:389:62: error: cannot initialize a parameter of type 'void *' with an rvalue of type 'const void *'
qsa_prompt.metal:234:16: note: in instantiation of member function 'metal::cooperative_tensor<float, ...>::cooperative_tensor' requested here
```

The cause is the copy `auto tot = acc;` in `qsa_flash`. Getting a second destination tensor instead compiles:

```diff
--- metal/kernels/qsa_prompt.metal
-    auto tot = acc;   // the running output; ...
+    auto tot = opO.template get_destination_cooperative_tensor<__tensor_ops_detail::__remove_addrspace_t<decltype(tPh)>,
+                                                               __tensor_ops_detail::__remove_addrspace_t<decltype(rv)>, float>();   // the running output; ...
```

(`tot` is zeroed right after, so the copy is only there for its type.) Note: `qsa_prompt_parity` then hangs the GPU on M2 (`kIOGPUCommandBufferCallbackErrorHang`), also with `COOP_K4=0`; I did not dig further. The engine itself (no M5 layout → no flash) runs fine.

### Result with both fixes (via `STRATA_METAL_KERNELS=<patched copy>`)

Correct answers, MTP accepting drafts (e.g. 65 of 101). IQ3_S on M2 Max 96 GB, `--spec 2`, temp 0, 3 runs each:

| | decode tok/s | prompt tok/s |
|---|---:|---:|
| short code / RU / story (400 tok) | 31–33 / 27 / 31 | — |
| after 5.6K / 22K prompt | 29–30 / 31–32 | 295–309 / 302–309 |

The prompt runs on the prompt chain with "the older expert / attention kernels (no M5 layout)".

Other parity tests on this Mac all pass: gemv, expert, gr, gr2, qsa, route, sample, fuse, pair, mma, prompt, dense_mm. `moe_mm_parity` reports `no function mm_gu_42_16_64_16_2` (probably a test-only issue).

Both diffs are above (full patch is the two hunks shown).

**Strata version, GPU, OS:** 0.1.44-metal (main `ed4773e`) · Apple M2 Max 96 GB · macOS 26.3.1