Mac Qwen3.6-27B Metal Q4_K perf diagnosis (2026-06-26)

Run on node-macos-a (Apple M3 Pro / Mac15,7, 12 CPU core = 6P+6E, 18 GPU core, 36 GB unified, macOS 26.5, Metal 4), fresh shallow public clone at HEAD 20fdf16, built -tags fakmetal (CGO Metal). Driven over Tailscale SSH from a Windows box. The llama-server (com.fak.qwen36-model, the 7.29 tok/s reference) was stopped for the clean arm (launchctl bootout → run → launchctl bootstrap restore) so fak had the whole 36 GB GPU/unified pool; an EXIT trap guaranteed the service came back.

TL;DR

fak’s resident-Q4_K Metal decode path is CORRECT but does not reach the speed bar. The 27B’s clean decode is 1.2 tok/s (vs the CPU resident-q4k baseline 0.9, the 3× goal 2.7, and the llama.cpp-Metal reference 7.29). The cause is not a wrong kernel — the Q4_K GEMV matches the CPU f32 reference to cosine 1.000000 and the greedy decode token sequence is bit-identical. The cause is orchestration: each decode token runs ~336 separate Metal command-buffer GEMVs (≈7 projection/MLP matmuls × ~48–64 layers), and an isolated GEMV is ~360 µs launch/sync-bound on top of ~98 µs of actual bandwidth-limited work. The lever is the one-command-buffer-per-token GPU-resident decode forward (the decode twin of the already-shipped resident prefill forward), plus a kernel-efficiency pass on q4k.m.

1. Clean measurement (the headline)

FAK_Q4K=1 FAK_METAL=1 FAK_GPU_LEASE_NOWAIT=1 FAK_QPROFILE=1 fakchat -gguf ~/.cache/fak-models/gguf/Qwen3.6-27B.q4_k_m.gguf -tok ~/.cache/fak-models/tokenizers/qwen3.6 -p "Write three sentences about why the ocean is important." -n 64 -quiet

backend banner: fak in-kernel Gated-DeltaNet (resident Q4_K decode + Q8 fallback, cached) [Metal q4_k prefill]. Qwen3.6-27B is a Gated-DeltaNet hybrid (cfg.IsQwen35Hybrid, 48 GDN layers + periodic full-attention), so fakchat routes it through runHybrid, which sets s.MetalQ4K = q4k && FAK_METAL.

metric clean (llama stopped) contended (llama up, swap) CPU resident-q4k llama.cpp-Metal 3× goal
decode tok/s 1.2 (64 tok / 54.12 s) 0.8 0.9 7.29 2.7
prefill tok/s 0.6 (29 tok / 48.27 s) 0.6 — 51.55 —
load ~40 s 190 s — — —
max RSS / peak footprint 26.0 GB / 53 GB 13.1 GB / 53 GB — — —
swaps 0 0 (but 15.4M page reclaims) — — —

The clean arm fits (26 GB RSS, 0 swaps). The contended arm has BOTH 27B models resident (peak 53 GB footprint, 15.4M page reclaims) — its 0.8 tok/s is swap-noise, not a fak number. Either way decode lands ~1 tok/s: fak is at 0.16× of the 3× decode goal and 0.16× of the llama.cpp-Metal bar.

2. Where the time goes (FAK_QPROFILE phase profiles)

Decode (total 54122 ms for 64 tokens = 845 ms/token):

phase ms % calls ms/call
mlp_decode 29200 54.0% 4096 7.13
qwen35_linear_step_in_proj 8697 16.1% 3072 2.83
qwen35_linear_step_out_proj 3453 6.4% 3072 1.12
qwen35_linear_step_recurrent 3199 5.9% 3072 1.04
full_attn_qkv_proj 2900 5.4% 1024 2.83
lm_head_q8 1934 3.6% 64 30.2
full_attn_o_proj 1066 2.0% 1024 1.04
rest (attn, conv, norms, gate) ~3700 6.8% — —

Prefill (total 48273 ms for 29 tokens): mlp_gate_up_proj 54.7% (412 ms/call), mlp_down_proj 18.2%, qwen35_linear_in_proj 11.2%.

The matmuls (MLP + projections) are ~85% of both phases. mlp_decode alone is 54% at 7.1 ms per single-token matmul — a Q4_K MLP weight read is ~40 MB, which at the device’s ~150 GB/s should be ~0.27 ms. The 26× gap is per-call command-buffer overhead and CPU↔GPU round-trips, not arithmetic.

3. The kernels are correct but the per-call overhead dominates (microbench)

go test -tags fakmetal on the same box:

So the q4k.m kernels are correct but slow in two compounding ways: (a) per-call command-buffer overhead (~360 µs), multiplied by ~336 matmuls/token in the live decode loop; (b) low in-kernel utilization (GEMV ~21% bandwidth, GEMM ~5% FLOP) even ignoring launch cost.

4. The lever (and the bandwidth ceiling)

A 27B Q4_K decode reads ~15 GB of weights per token. At ~150 GB/s that is a ~100 ms/token floor ≈ 10 tok/s ceiling; llama.cpp-Metal achieves 7.29 there. fak’s path leaves ~85% of the wall-clock in launch overhead and round-trips, so it sits at 1.2.

Primary lever — a one-command-buffer-per-token GPU-resident decode forward. The repo already ships the prefill analog: internal/metalgemm/forward.m mg_prefill runs the whole fresh prefill in ONE command buffer with the activation resident on-GPU (prefillMetalResident in internal/model/metal_prefill.go). The decode needs the twin: a mg_decode_step that, per token, keeps the f32 activation on-GPU across all projection/MLP matmuls, reads the GPU-resident KV cache, and submits ONE command buffer — paying the ~360 µs once per token instead of ~336 times. The kv.go MetalQ4K doc already names this: “a lone GEMV is occupancy-bound; the decode bar needs the one-command-buffer forward, a tracked follow-up.” Complication: the GDN recurrent scan + periodic full-attention must run on-GPU (or in a tight hybrid that does not round-trip per matmul).

Secondary lever — a kernel-efficiency pass on q4k.m. Even amortized, the GEMV at 21% bandwidth and the GEMM at 5% FLOP leave throughput on the table: revisit threadgroup/grid sizing, vectorized/coalesced dequant, simdgroup_matrix for the GEMM tile, and larger per-launch work so a single launch saturates more of the 18 GPU cores.

4b. Proof the lever works — batched GEMV witness (BenchmarkMetalQ4KGemvBatch)

To confirm the wall is the per-op command buffer (not the kernel), a new primitive metalgemm.Q4KWeight.GEMVBatch encodes N GEMVs of the same weight into ONE command buffer (one commit/waitUntilCompleted) and a benchmark times it against N separate calls. Clean run (llama-server stopped, M3 Pro, 50× benchtime):

path per-GEMV effective BW % of ~150 GB/s
BenchmarkMetalQ4KGemv — single (one command buffer each) 864 µs 17.1 GB/s ~11%
BenchmarkMetalQ4KGemvTiny — 256×256 (overhead floor) 188 µs — —
BenchmarkMetalQ4KGemvBatch — 64 in one command buffer 165 µs 89.1 GB/s ~59%

Batching is 5.2× faster per GEMV and lifts effective bandwidth from ~11% to ~59% of the device ceiling. Two effects compound: the ~188 µs fixed submit/sync overhead (the tiny-GEMV floor) is paid once instead of N times, and the GPU pipelines the N dispatches, hiding the memory latency a lone GEMV can’t. (The batch GEMVs the same weight N times — a measurement, not the decode access pattern — so the 89 GB/s is a per-dispatch-pipelining datum, not a model-decode rate.)

Projection for #67. A one-command-buffer-per-token resident decode forward reads ~15 GB of weights/token; at the measured batched 89 GB/s that is ~169 ms/token ≈ 5.9 tok/s, and ~8 tok/s if a kernel pass closes the remaining ~40% to bandwidth — right at the llama.cpp-Metal 7.29 bar and 5–7× over today’s 1.2 tok/s, clearing the 3× goal. The q4k_gemm prefill kernel stays at 4.8 GB/s / 377 GFLOP/s for the short P=22 panel (its 128-token tile leaves most threads idle at small P) — a separate kernel-occupancy fix for prefill.

5. Honest fences

Reproduce

# on node-macos-a, fresh clone built -tags fakmetal -> ~/fak-3xbench/{fakchat,modelbench}
# stop the launchd-managed llama-server for a clean GPU, restore after (EXIT trap):
launchctl bootout gui/$(id -u)/com.fak.qwen36-model
FAK_Q4K=1 FAK_METAL=1 FAK_GPU_LEASE_NOWAIT=1 FAK_QPROFILE=1 ~/fak-3xbench/fakchat \
  -gguf ~/.cache/fak-models/gguf/Qwen3.6-27B.q4_k_m.gguf \
  -tok ~/.cache/fak-models/tokenizers/qwen3.6 -p "..." -n 64 -quiet
launchctl bootstrap gui/$(id -u) ~/Library/LaunchAgents/com.fak.qwen36-model.plist
# kernel microbench:
go test -tags fakmetal -run NONE -bench 'BenchmarkMetalQ4KGemv|BenchmarkMetalQ4KGemm' -benchtime 30x ./internal/model

Raw artifact: experiments/benchmark/runs/by-machine/node-macos-a/20260626T055239Z-q4k-metal-decode-27b/score.json.