Model Compatibility & Performance Tests (M5 Max)

September 5, 2026 · View on GitHub

Compatibility and performance testing for mlxcel models on MacBook Pro M5 Max 128GB, with same-host mlx-lm / mlx-vlm reference measurements and M1 Ultra ratios where available.

Test Environment

ItemValue
HardwareMacBook Pro M5 Max, 128GB RAM
OSmacOS 26.6.2 (build 25G83)
mlxcel version0.7.0-beta.1 (mlxcel_version)
Source revisionb2ff1eee (mlxcel_commit); MLX pin 9a795735 (mlx_commit)
MLX versionupstream main (via mlxcel-core; pinned commit 9a795735)
mlx-lm baseline0.31.3 (dev checkout https://github.com/ml-explore/mlx-lm, commit ed1fca4); not re-run for the 0.6.0 sweep, see note below
mlx-vlm baseline0.4.4; not re-run for the 0.6.0 sweep
Test Prompt"Hello, how are you today?" (text) / "What is in this image?" (VLM)
Max Tokens100
Test Date2026-09-03/04 full text + VLM re-benchmark (0.6.0); prior: 2026-07-11/12 full sweep (0.4.0-rc.1), 2026-06-15 full sweep (0.2.1), 2026-05-27 full sweep (0.1.0)
Benchmark StatusFull text + VLM sweep on mlxcel 0.6.0 using mlxcel-bench-decode: 175 text model dirs via bench_decode.sh all (161 with decode numbers) plus the VLM-mode pass all --vlm (77 with decode numbers). Both runs used --cooldown 30 --big-cooldown 30, which remain required on this host: without them the sweep accumulates enough heat to thermally throttle the mid-sweep Qwen cluster. Time Machine was stopped and its automatic backups disabled for the whole campaign, since a running backup starves both disk and unified-memory bandwidth; a 3.9 TB copy was in fact active when the campaign started. This round also raised the up-front memory guard to a 90 GB weight budget (BENCH_MEM_OVERHEAD_FACTOR=1.209 against the 108.8 GB default limit), so deepseek-v3-4bit (99.96 GiB) now records SKIP:oom_estimate instead of consuming a slot and failing; qwen3-coder-480b-a35b-instruct-4bit (251.54 GiB) was already skipped. Decode is comparable to the 0.4.0-rc.1 sweep; prefill is not comparable for every model. Two correctness fixes landed after 2026-07-12 that change the measured condition rather than speed: #792 (aspect-ratio image processing for Pixtral/Mistral3, merged 2026-07-13) stopped upscaling every image to a fixed square, and the chat-template path now renders Llama's official template exactly. Affected rows therefore run at a much shorter prompt, and prefill_tok_s at the shorter length is arithmetically lower because the fixed per-call overhead is amortized over fewer tokens. See "Condition changes since 0.4.0-rc.1" below before reading any prefill delta. The vs M1 Ultra column is a cross-version ratio this round: 0.6.0 M5 Max decode over the 2026-07-12 0.4.0-rc.1 M1 Ultra sweep (benchmarks/metal_m1ultra_2026-07-12.csv), because that host had not been re-measured at 0.6.0 when this campaign ran. It will be refreshed by the pending 0.6.0 M1 Ultra sweep. The mlxcel vs mlx-lm / vs mlx-vlm percentages further down still carry the 2026-05-18 Python baselines; those sweeps were not re-run this round.

Which version each CSV column records

One column used to carry three different meanings across the corpus, which is why it has been split. Every CSV from this campaign records all of:

ColumnValue hereMeaning
mlxcel_version0.7.0-beta.1this repository's crate version
mlxcel_commitb2ff1eeethe source revision measured, 8 characters, -dirty when tracked files were modified
mlx_commit9a795735the pinned MLX C++ revision the binary links

The sweep ran at b2ff1eee, when Cargo.toml still read 0.6.0. The rows are labelled 0.7.0-beta.1 anyway because the tag (64f5d9b4) is exactly one commit later and that commit changes no compiled source: the diff is manifests, Cargo.lock, CHANGELOG.md, CITATION.cff, README.md, docs/environment-variables.md, the recipes registry and an issue template. Verify with:

git diff --name-only b2ff1eee v0.7.0-beta.1 | grep -E '\.(rs|cpp|metal|cu|h|hpp)$'

which returns nothing. mlxcel_commit pins the exact revision either way, so nothing is lost by the label.

mlx_commit exists because an MLX pin bump changes kernels without moving either mlxcel field, so a sweep taken across one would otherwise look identical to a sweep taken before it.

Legend

  • ✅ Pass: Model works correctly
  • ⚠️ Partial: Loads but output quality problems or low token count
  • ❌ Fail: Does not work

Basic Transformers

ModelTest ModelStatusPrefillDecodevs M1 UltraNotes
llama3Llama-3.2-1B-Instruct-4bit5337.36556.361.32x48 tokens
llama3.1Llama-3.1-8B-Instruct-4bit986.34114.921.08x54 tokens
llama3 (8B bf16)Llama-3.1-8B-Instruct (bf16)⚠️800.5132.810.96x44 tokens; bf16; slow decode
llama4Llama-4-Scout-17B-16E-4bit95.4848.121.35x59 tokens
command-r7bc4ai-command-r7b-4bit240.44113.271.04x100 tokens
aya-expanse-8baya-expanse-8b-4bit228.05111.571.06x100 tokens
aya-vision-8baya-vision-8b (text-only)3245.72108.661.07x45 tokens; text-only
deepseek-r1DeepSeek-R1-Distill-Qwen-7B-4bit436.26125.701.15x100 tokens
internlm2InternLM2-7B-4bit528.62118.341.13x100 tokens
internlm3internlm3-8b-instruct-4bit733.86101.661.21x100 tokens
mimoMiMo-7B-RL-4bit792.79119.631.41x100 tokens
minicpmMiniCPM-2B-sft-bf16-4bit804.30229.721.54x31 tokens
bunny-llama3-8bBunny-Llama-3-8B-V-4bit (text)551.49114.401.14x40 tokens; text-only
llava-1.5-7bllava-1.5-7b-4bit (text)332.78123.311.09x60 tokens; text-only
llava-nextllava-v1.6-mistral-7b-4bit (text)473.14120.601.07x49 tokens; text-only
llava-interleavellava-interleave-qwen-0.5b-bf16 (text)1605.28392.351.23x23 tokens

Gemma Family

ModelTest ModelStatusPrefillDecodevs M1 UltraNotes
gemmagemma-2b-it-4bit1280.60213.761.12x49 tokens
gemma2gemma-2-2b-it-4bit1234.14242.321.41x18 tokens; full-budget raw prompt 245.83 tok/s
gemma3gemma-3-1b-it-4bit1938.99391.421.72x30 tokens
gemma3 (4B)gemma-3-4b-it-4bit850.64183.691.56x84 tokens; full-budget raw prompt 183.77 tok/s
gemma3n (E2B)gemma-3n-E2B-it-4bit782.40157.921.79x72 tokens
gemma3n (E4B)gemma-3n-E4B-it-4bit583.32110.401.65x74 tokens
gemma3n (E4B bf16)gemma-3n-E4B-it (bf16)417.4639.541.14x69 tokens; Gemma3n language MLP bf16 preserved, other bf16 materialized as f16; M5 (Neural Accelerator) uses the split decode path while other Apple Silicon uses the fused path; ~80% of mlx-lm decode
gemma4 (26B MoE)gemma-4-26b-a4b-it-4bit561.93149.411.88x26 tokens
gemma4 (31B)gemma-4-31b-4bit106.4428.521.43x100 tokens
gemma4 (31B IT)gemma-4-31b-it-4bit146.3027.341.43x26 tokens
gemma4 (31B nvfp4)Gemma-4-31b-it-nvfp4⚠️130.1615.651.22x26 tokens; nvfp4 has no fast Metal kernel
gemma4 (E2B 4bit)gemma-4-e2b-it-4bit1466.17210.271.66x72 tokens
gemma4 (E2B 8bit)gemma-4-e2b-it-8bit1264.72149.051.44x79 tokens
gemma4 (E4B 4bit)gemma-4-e4b-it-4bit826.35141.781.66x100 tokens
gemma4 (E4B 8bit)gemma-4-e4b-it-8bit716.1185.781.30x76 tokens
gemma4 (12B)gemma-4-12b-it-4bit320.2243.961.19x27 tokens; NEW (6-13)
gemma4 (26B QAT)gemma-4-26b-a4b-it-qat-4bit547.50139.391.80x26 tokens; QAT; NEW (6-13)
gemma4 (31B IT QAT)gemma-4-31b-it-qat-4bit130.7917.241.10x26 tokens; QAT; NEW (6-13)
gemma4 (E2B QAT)gemma-4-e2b-it-qat-4bit1109.34173.581.53x39 tokens; QAT; NEW (6-13)
gemma4 (E4B QAT)gemma-4-e4b-it-qat-4bit617.0296.111.37x33 tokens; QAT; NEW (6-13)

EXAONE

ModelTest ModelStatusPrefillDecodevs M1 UltraNotes
exaoneEXAONE-3.5-2.4B-Instruct-4bit2300.77284.141.43x43 tokens
exaone4exaone-4.0-1.2b-4bit1977.37422.491.85x10 tokens

Qwen Family

ModelTest ModelStatusPrefillDecodevs M1 UltraNotes
qwen2.5 (0.5B)Qwen2.5-0.5B-Instruct-4bit8467.27660.771.73x39 tokens
qwen2.5 (0.5B bf16)Qwen2.5-0.5B-Instruct (bf16)5814.81400.411.35x37 tokens
qwen2.5 (7B)Qwen2.5-7B-Instruct-4bit880.79124.111.14x41 tokens
qwen2.5 (7B 8bit)Qwen2.5-7B-Instruct-8bit813.5567.691.00x44 tokens
qwen2.5-vl (3B)qwen2.5-vl-3b-4bit1409.70162.041.47x39 tokens; re-downloaded (prior FAIL was a corrupt checkpoint, not a code bug)
qwen2-vl (2B)Qwen2-VL-2B-Instruct-4bit2506.64268.311.58x35 tokens
qwen1.5-moeQwen1.5-MoE-A2.7B-Chat-4bit873.72255.171.71x36 tokens
qwen3 (0.6B)Qwen3-0.6B-4bit3696.62601.122.63x9 tokens
qwen3 (1.7B)Qwen3-1.7B-4bit1747.04380.001.80x38 tokens
qwen3 (4B)Qwen3-4B-4bit1004.44191.191.57x36 tokens
qwen3 (8B)Qwen3-8B-4bit571.59112.241.41x33 tokens
qwen3-30b-a3bQwen3-30B-A3B-4bit415.14174.242.06x34 tokens
qwen3-moeQwen3-MoE-30B-4bit415.67173.612.08x34 tokens
qwen3-vl (2B)Qwen3-VL-2B-Instruct-4bit1353.08370.751.80x59 tokens; text-only
qwen3-vl (4B)qwen3-vl-4b-4bit781.44183.891.56x49 tokens; text-only; NEW (6-13)
qwen3-vl (8B)qwen3-vl-8b-4bit444.34110.701.38x57 tokens; text-only; NEW (6-13)
qwen3-vl (30B MoE)Qwen3-VL-30B-A3B-Instruct-4bit324.18167.902.04x34 tokens; text-only
qwen3-vl (32B)Qwen3-VL-32B-Instruct-4bit122.6427.461.33x30 tokens; text-only
qwen3-next (80B MoE)Qwen3-Next-80B-A3B-Instruct-4bit328.29119.372.01x54 tokens; NEW (0.4.0-rc.1)
qwen3-omni (30B MoE)Qwen3-Omni-30B-A3B-Instruct-4bit408.75167.382.02x42 tokens; text path; NEW (0.4.0-rc.1)
qwen3-coder (480B)Qwen3-Coder-480B-A35B-Instruct-4bit-FAIL-SKIP:oom_estimate
qwen3.5 (0.8B)Qwen3.5-0.8B-4bit3275.49500.092.00x18 tokens
qwen3.5 (2B)Qwen3.5-2B-4bit1671.10327.381.71x31 tokens
qwen3.5 (4B)Qwen3.5-4B-4bit876.75168.711.56x31 tokens
qwen3.5 (9B)Qwen3.5-9B-4bit512.04101.741.44x31 tokens
qwen3.5 (9B bf16)Qwen3.5-9B (bf16)337.0930.120.98x31 tokens
qwen3.5 (27B)Qwen3.5-27B-4bit168.6932.751.37x30 tokens
qwen3.5-35b-a3bQwen3.5-35B-A3B-4bit694.44159.971.96x31 tokens
qwen3.6-35b-a3bQwen3.6-35B-A3B-4bit696.81152.721.91x27 tokens; NEW (5-18)

Phi Family

ModelTest ModelStatusPrefillDecodevs M1 UltraNotes
phi-2phi-2-hf-4bit-mlx⚠️313.07106.081.83x1 token; (likely EOS)
phi-3-miniPhi-3-mini-4k-instruct-4bit591.88210.391.28x25 tokens
phi-3.5-miniPhi-3.5-mini-instruct-4bit582.33203.681.25x40 tokens
phi-3.5-moePhi-3.5-MoE-instruct-4bit126.64110.021.46x35 tokens
phi-3.5-visionPhi-3.5-vision-instruct-4bit872.57203.241.25x36 tokens; text-only
phi-4Phi-4-4bit240.7259.451.03x16 tokens

OLMo Family

ModelTest ModelStatusPrefillDecodevs M1 UltraNotes
olmo-1bOLMo-1B-hf-4bit780.33242.641.19x100 tokens
olmo2-7bOLMo2-7B-4bit621.99117.831.18x27 tokens
olmo3-32bOLMo3.1-32B-4bit470.9129.291.37x100 tokens

MoE (Mixture of Experts)

ModelTest ModelStatusPrefillDecodevs M1 UltraNotes
minimaxMiniMax-M2-3bit185.8173.79-100 tokens
mixtralMixtral-8x7B-Instruct-v0.1-4bit106.2265.561.27x73 tokens
gpt_oss (20B)gpt-oss-20b-MXFP4-Q41231.56173.191.94x100 tokens
gpt_oss (120B)gpt-oss-120b-4bit332.76112.831.90x58 tokens
solar-open-100bSolar-Open-100B-4bit287.5265.511.87x100 tokens
dots.llm1dots.llm1.inst-mixed-4-6bit102.7650.88-39 tokens; mixed 4/6-bit; NEW (6-13)
lfm2-moelfm2-8b-a1b-4bit1088.26338.211.83x37 tokens; NEW (6-13)

DeepSeek Family

ModelTest ModelStatusPrefillDecodevs M1 UltraNotes
deepseekdeepseek-coder-1.3b-instruct-4bit6278.89176.761.11x37 tokens
deepseek_v2DeepSeek-V2-Lite-Chat-4bit389.18210.842.04x44 tokens
deepseek_v3--FAIL-SKIP:oom_estimate

MLA (Multi-head Latent Attention)

ModelTest ModelStatusPrefillDecodevs M1 UltraNotes
minicpm3MiniCPM3-4B-4bit568.98131.781.51x39 tokens

Nemotron Family

ModelTest ModelStatusPrefillDecodevs M1 UltraNotes
nemotron_hNemotron-H-30B-4bit406.57176.011.92x46 tokens
nemotron_nasNemotron-NAS-30B-A3B-4bit406.48176.611.95x46 tokens
nemotron-omniNemotron-3-Nano-Omni-30B-A3B-Reasoning-4bit404.72171.422.01x19 tokens; text path; NEW (6-14)

SSM / Mamba Models

ModelTest ModelStatusPrefillDecodevs M1 UltraNotes
mambaFalcon-Mamba-7B-4bit⚠️274.9866.601.53x2 tokens; chat template EOS
mamba2mamba2-1.3b-4bit959.29163.211.42x100 tokens
mamba2 (130M)mamba2-130m2260.89348.791.34x100 tokens; NEW (6-14)
jambaJamba-v0.1-4bit1086.94216.871.68x100 tokens; raw prompt 215.74 tok/s
falcon-h1falcon-h1-tiny-90m-instruct-4bit1016.95154.820.48x30 tokens; Mamba2 + attention hybrid; NEW (6-13)
plamo2plamo-2-1b381.1781.490.77x100 tokens; Mamba + attention hybrid; NEW (6-13)

Chinese / Asian Language Models

ModelTest ModelStatusPrefillDecodevs M1 UltraNotes
baichuanBaichuan-M1-14B-Instruct-4bit158.1557.661.46x7 tokens
glm4_moe_liteGLM-4.7-Flash-4bit242.30106.112.17x18 tokens
ernie4_5ERNIE-4.5-0.3B-Instruct-4bit7753.121068.372.05x100 tokens
hunyuan_moehunyuan-a13b-instruct-4bit153.1464.611.46x36 tokens; A13B MoE (4-bit), canonical after checkpoint dedup
hunyuan_v1_denseHunyuan-1.8B-Instruct-4bit1098.09330.111.82x42 tokens

Other Models

ModelTest ModelStatusPrefillDecodevs M1 UltraNotes
ministral3Ministral-3B-Instruct-4bit6729.43225.241.57x34 tokens; VLM wrapper
mistral-smallmistral-small-3.1-24b-4bit1000.1439.651.38x20 tokens
molmo2molmo2-4b680.14103.471.76x33 tokens
molmo-7bmolmo-7b372.30123.871.84x59 tokens; text spot-check
internvl3internvl3-1b8271.03664.131.91x37 tokens
smollm-135mSmolLM-135M-Instruct-4bit5105.84926.252.21x100 tokens
smollm3-3bSmolLM3-3B-4bit2315.35231.611.76x19 tokens
stablelm-1.6bstablelm-2-1_6b-chat-4bit2487.59428.361.62x59 tokens
starcoder2-3bstarcoder2-3b-4bit437.34217.031.33x100 tokens
pixtral-12bpixtral-12b-4bit205.4272.601.07x28 tokens; text-only
paligemma2-3bpaligemma2-3b (6-bit)482.68195.86-100 tokens; text-only

Granite Family

Ported 2026-06-13 (dense + Mamba2/attention hybrid + hybrid-MoE). vs M1 Ultra ratios are from the 2026-07-12 0.4.0-rc.1 M1 Ultra sweep.

ModelTest ModelStatusPrefillDecodevs M1 UltraNotes
granitegranite-3.3-2b-instruct-4bit3384.19258.601.39x31 tokens; dense
granite4_h (350M)granite-4.0-h-350m-4bit1682.36130.120.55x19 tokens; Mamba2 + attention hybrid
granite4_h (tiny)granite-4.0-h-tiny-4bit428.6975.050.70x44 tokens; hybrid MoE
granite4.1 (3B)granite-4.1-3b-4bit688.29177.491.37x7 tokens
granite4.1 (8B)granite-4.1-8b-4bit⚠️383.8053.991.90x1 token; likely early EOS, re-check with a code prompt

Recently Ported Families (2026-06-13/14)

New architectures landed in the 06-13/06-14 wave. vs M1 Ultra ratios are from the 2026-07-12 0.4.0-rc.1 M1 Ultra sweep.

ModelTest ModelStatusPrefillDecodevs M1 UltraNotes
apertusapertus-8b-instruct-2509-4bit1297.71112.881.39x40 tokens; xIELU, QK-norm, llama3 RoPE
bitnet (4bit pack)bitnet-b1.58-2b-4t-4bit393.03327.822.21x33 tokens; 1.58-bit ternary
bitnetbitnet-b1.58-2b-4t384.45260.611.91x33 tokens; 1.58-bit ternary
lfm2lfm2-350m-8bit5516.27858.031.53x13 tokens; 8-bit
seed-ossseed-oss-36b-instruct-4bit92.2726.571.35x100 tokens
minicpm-v (4.6)minicpm-v-4.6-bf161465.65276.171.31x100 tokens; text path
youtu-vlyoutu-vl-4b-instruct770.8547.371.07x93 tokens; text path

The following newly-added checkpoints are present in models/ but are not measurable by the text decode harness this round:

  • glm-5-4bit, glm-5.1-4bitFAIL:bench, resolved at the 0.6.0 sweep as a local checkpoint problem, not a load-path defect: glm-5-4bit still holds 21 GB of .incomplete download blobs with no materialized *.safetensors or tokenizer, and glm-5.1-4bit holds only .gitattributes. Both need a re-download before GLM-5 support can be judged.
  • minicpm-v-4.6-mxfp4 — was FAIL:bench at 0.4.0-rc.1; passes at 0.6.0 (1568.60 prefill / 346.72 decode, 100 tokens). See the 0.6.0 section below.
  • diffusiongemma-26b-a4b-it-4bit — block-diffusion generation; decode tok/s is not a meaningful metric for this harness.
  • docling-layout-heron-mlx-bf16 — layout/vision model; no text decode path.
  • granite-speech-4.1-2b-nar-mlx — non-autoregressive speech model; no text decode path.
  • gemma-4-12b-it-assistant-4bit, gemma-4-31b-it-assistant-bf16 — MTP drafter checkpoints, not standalone generators.
  • qwen3.5-0.8b-optiq-4bit, qwen3.5-27b-dflash, qwen3.5-4b-dflash — experimental quant/decode variants; FAIL:bench standalone.

Recently Ported Families (0.4.0-rc.1 / 2026-07-11)

New architectures landed since the 0.2.1 sweep. Numbers are the text-decode path (bench_decode.sh all); the vision-capable models also appear in the VLM table below with their image-prompt numbers. vs M1 Ultra ratios are from the 2026-07-12 0.4.0-rc.1 M1 Ultra sweep (- where the M1 Ultra sweep lacks the model).

ModelTest ModelStatusPrefillDecodevs M1 UltraNotes
gemma2 (9B 8bit)gemma-2-9b-8bit172.2949.84-100 tokens
mistral-small-4 (119B)mistral-small-4-119b-2603-4bit164.7619.051.02x60 tokens; dense 119B
phi-3-smallphi-3-small-8k-instruct-aq4_64662.83111.711.12x100 tokens; aq4_64
llada2.0-minillada2.0-mini-preview-4bit1826.58336.842.15x100 tokens; diffusion LM
deepseek-ocrdeepseek-ocr-4bit979.99699.312.26x21 tokens; text path
deepseek-ocr-2deepseek-ocr-2-4bit965.56700.412.23x20 tokens; text path
deepseek-vl2deepseek-vl2-small-4bit749.35207.181.85x25 tokens; text path
fastvlmfastvlm-0.5b-bf164492.54402.201.36x46 tokens; text path
glm-4.1vglm-4.1v-9b-thinking-4bit298.7766.38-81 tokens; text path
glm-4.5vglm-4.5v-4bit84.4416.74-47 tokens; text path
glm-ocrglm-ocr-4bit2781.34460.041.97x7 tokens; text path
granite4-vision (3B)granite-4.0-3b-vision-4bit1504.10201.951.57x26 tokens; text path
granite-vision (2B)granite-vision-3.2-2b-4bit2473.11257.631.60x20 tokens; text path
idefics2idefics2-8b-4bit317.03121.951.11x100 tokens; text path
idefics3idefics3-8b-llama3-4bit277.43118.321.13x100 tokens; text path
kimi-vlkimi-vl-a3b-thinking-4bit596.60177.21-100 tokens; A3B MoE; text path
lfm2-vllfm2-vl-450m-4bit5423.66927.562.02x19 tokens; text path
llama-3.2-vision (11B)llama-3.2-11b-vision-instruct-4bit517.49113.78-51 tokens; text path
moondream2moondream2183.6943.10-100 tokens; text path
paddleocr-vlpaddleocr-vl-bfloat16931.09157.761.20x100 tokens; text path
smolvlmsmolvlm-instruct-bf16716.61131.02-100 tokens; text path

Newly-added checkpoints present in models/ but not measurable by the text decode harness this round:

  • deepseek-v4-flash-4bit, qwen3-coder-480b-a35b-instruct-4bit: SKIP:oom_estimate on the 128 GB budget.
  • kokoro-82m (TTS), whisper-base (ASR): FAIL:bench; no autoregressive text-decode path.
  • dots.ocr-4bit: loads and runs but emits 0 text tokens on the plain prompt; decode tok/s is not meaningful.

Recently Ported Families (0.6.0 / 2026-09-03)

New checkpoints measured for the first time in the 0.6.0 sweep. Numbers are the text-decode path (bench_decode.sh all). vs M1 Ultra ratios are cross-version against the 2026-07-12 0.4.0-rc.1 M1 Ultra sweep (- where that sweep lacks the model); they will be refreshed by the pending 0.6.0 M1 Ultra run.

ModelTest ModelStatusPrefillDecodevs M1 UltraNotes
qwen3.8-27bqwen3.8-27b-4bit168.0832.88-37 tokens; NEW (0.6.0); qualified on the qwen3_5 path (#1174)
hunyuan (13B)hunyuan-13b89.2864.221.46x31 tokens; NEW (0.6.0)
diffusiongemma (26B MoE)diffusiongemma-26b-a4b-it-4bit529.23114.691.69x27 tokens; NEW (0.6.0)
minicpm-v (4.6 mxfp4)minicpm-v-4.6-mxfp41568.60346.721.56x100 tokens; NEW (0.6.0)
qwen3.5 (0.8B optiq)qwen3.5-0.8b-optiq-4bit3209.41429.251.84x19 tokens; NEW (0.6.0)
qwen2.5 (1.5B)qwen2.5-1.5b-instruct-4bit3445.58383.361.54x31 tokens; NEW (0.6.0)
dots.ocrdots.ocr-4bit⚠️630.780.00-0 tokens; loads and prefills but emits no text on a text-only prompt
glm-5glm-5-4bit-FAIL-FAIL:bench, but not a runtime defect: the local checkpoint is an interrupted download (21 GB still sitting as .incomplete blobs under .cache/huggingface/download/, no *.safetensors and no tokenizer materialized). Re-download before reading this as a GLM-5 support gap
glm-5.1glm-5.1-4bit-FAIL-FAIL:bench, but not a runtime defect: the local directory holds only .gitattributes (8 KB of cache metadata, no config.json), i.e. the download never started. Re-download before reading this as a GLM-5.1 support gap

Duplicate checkpoint directories (not listed separately)

The local models/ store keeps several checkpoints under two names. Each pair below has byte-identical safetensors sizes and an identical tensor layout (same names, shapes, dtypes and offsets, verified by hashing the safetensors headers), so they are the same weights and are benchmarked twice per sweep. Only the name already carried in the tables above gets a row; adding the alias would double-count the model in the Summary Statistics. Teaching the harness to skip them is tracked as issue #1615.

Listed asAlso on disk as
qwen2.5-7b-instruct-4bitqwen2.5-7b, qwen2.5-7b-4bit
qwen2.5-1.5b-instruct-4bitqwen2.5-1.5b-4bit
qwen2-vl-2b-4bitqwen2-vl-2b
qwen3-vl-2b-4bitqwen3-vl-2b
qwen3-vl-4b-4bitqwen3-vl-4b-instruct-4bit
qwen3-vl-8b-4bitqwen3-vl-8b-instruct-4bit
qwen3-0.6b-4bitqwen3-0.6b
qwen2.5-0.5b-4bitqwen2-0.5b
gemma-3-4b-it-4bitgemma3-4b-4bit
llama-3.1-8b-4bitmeta-llama-3.1-8b-instruct-4bit
pixtral-12b-4bitpixtral-12b

VLM (image input) — full sweep

Below table reports the per-VLM-prompt run from bench_decode.sh all --vlm. All entries use the VLM prompt 'What is in this image?' with tests/fixtures/test_image.png.

ModelTest ModelStatusPrefillDecodevs M1 UltraNotes
aya-vision-8baya-vision-8b2660.08109.151.12x29 tokens
bunny-llama3-8bbunny-llama3-8b-4bit2844.29109.721.19x37 tokens
gemma3 (4B)gemma3-4b-4bit549.56134.261.51x9 tokens
gemma3n (E2B 4bit)gemma3n-e2b-4bit2968.80152.581.90x28 tokens
gemma3n (E4B 4bit)gemma3n-e4b-4bit2240.39105.431.73x29 tokens
gemma3n (E4B bf16)gemma3n-e4b-bf162171.4338.081.19x20 tokens; bf16→f16 conversion path
gemma4 (26B MoE)gemma-4-26b-a4b-it-4bit885.66140.392.01x27 tokens
gemma4 (31B)gemma-4-31b-4bit428.4722.041.50x5 tokens
gemma4 (31B IT)gemma-4-31b-it-4bit439.7225.541.40x24 tokens
gemma4 (E2B 4bit)gemma-4-e2b-it-4bit2833.78220.592.04x100 tokens
gemma4 (E2B 8bit)gemma-4-e2b-it-8bit2644.83146.411.59x100 tokens
gemma4 (E4B 4bit)gemma-4-e4b-it-4bit2086.99131.671.76x90 tokens
gemma4 (E4B 8bit)gemma-4-e4b-it-8bit1922.2078.151.32x64 tokens
internvl3 (1B)internvl3-1b6553.50601.002.77x8 tokens
llama4 (Scout)llama-4-scout-17b-4bit397.7248.011.45x100 tokens
llava-1.5-7bllava-1.5-7b-4bit3216.08112.161.14x21 tokens
llava-interleavellava-interleave-qwen-0.5b-bf1615979.88355.351.40x32 tokens
llava-nextllava-next-mistral-7b-4bit2977.26114.611.13x23 tokens
ministral3ministral-3b-4bit5527.54223.311.85x56 tokens
mistral-small (3.1 24B)mistral-small-3.1-24b-4bit1048.4540.021.44x31 tokens
molmo-7bmolmo-7b2348.36122.791.56x65 tokens; mlx-vlm baseline is a 1-token anomaly
molmo2 (4B)molmo2-4b2471.17100.971.72x46 tokens
paligemma2 (3B 6-bit)paligemma2-3b-6bit5246.0696.781.61x2 tokens
phi-3.5-visionphi-3.5-vision-4bit3754.55175.751.49x14 tokens
pixtral (12B)pixtral-12b-4bit1588.4873.581.26x30 tokens; intermittent slow VLM decode reads (~20 tok/s) seen on M5, not consistently reproducible (see Known Issues)
qwen2-vl (2B)qwen2-vl-2b-4bit2452.49248.761.74x12 tokens; EOS-terminate
qwen2.5-vl (3B)qwen2.5-vl-3b-4bit1629.01143.871.35x8 tokens; re-downloaded (prior FAIL was a corrupt checkpoint)
qwen3-vl (2B)qwen3-vl-2b-4bit2159.95273.881.55x100 tokens
qwen3-vl (4B)qwen3-vl-4b-4bit1201.21136.191.43x41 tokens; NEW (6-13)
qwen3-vl (8B)qwen3-vl-8b-4bit991.4781.681.30x38 tokens; NEW (6-13)
qwen3-vl (30B MoE)qwen3-vl-30b-a3b-4bit548.1758.011.43x63 tokens
qwen3-vl (32B)qwen3-vl-32b-4bit297.2619.001.08x49 tokens
gemma4 (12B)gemma-4-12b-it-4bit1408.0343.741.28x25 tokens; NEW (6-13)
minicpm-v (4.6)minicpm-v-4.6-bf16967.29263.501.49x23 tokens; NEW (6-13)
nemotron-omninemotron-3-nano-omni-30b-a3b-reasoning-4bit646.83154.512.21x6 tokens; NEW (6-14)
youtu-vlyoutu-vl-4b-instruct521.8845.521.06x30 tokens; NEW (6-13)
deepseek-ocrdeepseek-ocr-4bit1536.22614.382.45x15 tokens; NEW (0.4.0-rc.1)
deepseek-ocr-2deepseek-ocr-2-4bit1528.38382.592.57x5 tokens; NEW (0.4.0-rc.1)
deepseek-vl2deepseek-vl2-small-4bit859.31172.791.80x8 tokens; NEW (0.4.0-rc.1)
fastvlmfastvlm-0.5b-bf162655.91389.271.42x100 tokens; NEW (0.4.0-rc.1)
glm-4.1vglm-4.1v-9b-thinking-4bit826.9762.57-100 tokens; NEW (0.4.0-rc.1)
glm-4.5vglm-4.5v-4bit212.1615.98-39 tokens; NEW (0.4.0-rc.1)
granite4-vision (3B)granite-4.0-3b-vision-4bit2687.46201.951.64x30 tokens; NEW (0.4.0-rc.1)
granite-vision (2B)granite-vision-3.2-2b-4bit6215.28230.811.97x47 tokens; NEW (0.4.0-rc.1)
idefics2idefics2-8b-4bit929.92114.511.11x12 tokens; NEW (0.4.0-rc.1)
idefics3idefics3-8b-llama3-4bit2101.44118.511.14x100 tokens; NEW (0.4.0-rc.1)
kimi-vlkimi-vl-a3b-thinking-4bit757.42175.48-100 tokens; A3B MoE; NEW (0.4.0-rc.1)
lfm2-vllfm2-vl-450m-4bit5713.971019.852.57x47 tokens; NEW (0.4.0-rc.1)
llama-3.2-vision (11B)llama-3.2-11b-vision-instruct-4bit14.3372.31-69 tokens; NEW (0.4.0-rc.1)
moondream2moondream253.2031.50-4 tokens; NEW (0.4.0-rc.1)
paddleocr-vlpaddleocr-vl-bfloat163374.52124.871.18x12 tokens; NEW (0.4.0-rc.1)
smolvlmsmolvlm-instruct-bf162164.34116.10-5 tokens; NEW (0.4.0-rc.1)
qwen3-omni (30B)qwen3-omni-30b-a3b-instruct-4bit583.4937.541.66x2 tokens; NEW (0.4.0-rc.1)

Condition changes since 0.4.0-rc.1

Two fixes landed after the 2026-07-12 sweep that change what the harness measures, not how fast it runs. Both shorten the prompt, and prefill_tok_s is prompt tokens over prefill milliseconds, so a shorter prompt amortizes the fixed per-call overhead over fewer tokens and reads lower. Decode is unaffected. Do not read either as a slowdown.

1. Image processing no longer upscales (#792, merged 2026-07-13)

fix(vision): aspect-ratio image processing for Pixtral and Mistral3 replaced the fixed-square SigLIP path, which force-resized every image to a square, with a processor that downscales to fit size.longest_edge and never upscales. The VLM fixture is 224x224 and Pixtral's longest_edge is 1024 with patch 16, so the arithmetic is exact:

old (forced square)new (aspect-preserving)
resized to1024 x 1024224 x 224 (unchanged)
image tokens(1024/16)^2 = 4096(224/16)^2 = 196
measured prompt_tokens4099213

VLM prompt_tokens for the affected families dropped accordingly:

Model2026-07-122026-09-04
pixtral-12b-4bit, pixtral-12b4099213
mistral-small-4-119b-2603-4bit304693
mistral-small-3.1-24b-4bit3206253
ministral-3b-4bit3566613

Their VLM decode deltas (mistral-small-4-119b +20.4%, ministral-3b +13.1%, pixtral +9.9%) are context-length effects, not speedups.

2. The chat template renders Llama's official prompt

The text harness applies the checkpoint's chat template. At 0.6.0 the Llama family's rendering of the standard test prompt is 42 tokens; the 0.4.0-rc.1 sweep recorded 98. Tokenizing the canonical Llama 3.1 rendering of "Hello, how are you today?" with the checkpoint's own tokenizer.json gives 42 tokens (and 7 with --no-chat-template), which is exactly what the 0.6.0 binary reports, so the current value is the correct one and 98 carried roughly 56 spurious tokens.

Ten models therefore show a double-digit prefill_tok_s drop that is entirely a prompt-length change, not a regression:

Modelprefill 07-12prefill 09-03prompt tokens
smolvlm-instruct-bf164099.59716.6156 -> 9
llama-3.2-11b-vision-instruct-4bit2132.46517.4998 -> 17
llava-interleave-qwen-0.5b-bf165050.141605.2826 -> 8
llama-3.1-8b-4bit2138.85986.3498 -> 42
idefics3-8b-llama3-4bit580.55277.4318 -> 9
llama-4-scout-17b-4bit195.1995.4869 -> 18
llama-3.1-8b-bf161611.89800.5199 -> 43
llama-3.2-1b-4bit8068.685337.3699 -> 43
granite-vision-3.2-2b-4bit2827.362473.1155 -> 47
meta-llama-3.1-8b-instruct-4bit2142.46987.1198 -> 42

Only three checkpoints lost more than 10% of prefill at an unchanged prompt length, and none is a clean signal on its own: dots.ocr-4bit (-24.4%, already ⚠️ because it emits no text), smollm-135m-4bit (-16.2%, a 135M model at 16 prompt tokens) and stablelm-1.6b-4bit (-13.7%, 26 prompt tokens). All three sit at prompt lengths where per-call overhead dominates. They are the only prefill candidates worth a targeted re-measurement, tracked as issue #1614.

Summary Statistics

Counts reflect the 2026-09-03/04 bench_decode.sh all --cooldown 30 --big-cooldown 30 text sweep on mlxcel 0.6.0, with BENCH_MEM_OVERHEAD_FACTOR=1.209 (a 90 GB weight budget).

StatusCount
✅ Pass (measured decode)144
⚠️ Partial (loads; early EOS, slow path, or no text output)6
❌ Fail / OOM-skip4

How 154 table rows reconcile with 175 benchmarked checkpoints. The sweep walks every directory in models/, but the tables above deliberately do not carry a row per directory:

Count
Text table rows154
...of which map to a row in this sweep's CSV153
...of which have no CSV row (MiniMax-M2-3bit, removed from disk since 0.4.0-rc.1)1
Benchmarked checkpoints with no table row22
...duplicate directories of a listed checkpoint (see the alias table above)12
...non-text checkpoints the text harness cannot decode (whisper-base, kokoro-82m, granite-speech-4.1-2b-nar-mlx, docling-layout-heron-mlx-bf16)4
...speculative drafters and dflash variants, measured in the speculative table instead6
Total benchmarked checkpoints175

161 of the 175 checkpoints produced decode numbers. The 14 non-runs are 2 SKIP:oom_estimate (deepseek-v3-4bit at 99.96 GiB and qwen3-coder-480b-a35b-instruct-4bit at 251.54 GiB, both over the 90 GB weight budget) and 12 FAIL:bench, none of which is a decode-path defect: the 4 non-text checkpoints above, the 6 drafter / dflash variants (not standalone generative models), and the GLM-5 pair, which this round was traced to interrupted local downloads rather than a load-path bug (see Known Issues).

The 6 ⚠️ partials are phi-2-4bit, falcon-mamba-7b-4bit, gemma-4-31b-it-nvfp4, llama-3.1-8b-bf16 and granite-4.1-8b-4bit (early EOS or no fast kernel), plus dots.ocr-4bit, which loads and prefills but emits no text on a text-only prompt.

No decode regression survived verification. 159 checkpoints are comparable with the 0.4.0-rc.1 sweep (both runs produced decode numbers); 156 of them moved less than 10% either way and none moved more than 10% slower. Three moved more than 10% faster, and they are not equivalent: molmo2-4b (+61.7%, 63.99 -> 103.47) generated the same 33 tokens with prefill also up 26% and is a real gain, corroborated by the VLM pass (+56.9% at an unchanged 46 tokens); molmo-7b (+56.5%) changed run length from 24 to 59 tokens, so the two numbers are not directly comparable; and phi-2-4bit (+37.6%) is a single-token run, which is a latency sample rather than a throughput measurement.

Batched serving (B = 1/2/4)

Source: benchmarks/metal_m5max_batch_2026-09-04.csv, produced by scripts/bench_serving_concurrency.py against one fresh mlxcel-server per model, pinned to --parallel 4 --max-batch-prefill 4 at the standard condition (--prompt-tokens 512 --max-tokens 128). Under continuous batching N concurrent streaming clients occupy N decode slots, so the concurrency level is the effective decode batch size.

Reading the TTFT column. Levels run in ascending order and share the synthetic prompt, so B=2 and B=4 start with a warm prompt cache while B=1 pays the cold prefill. That is the intended serving-side condition, and it is why the two dense models show TTFT falling from B=1 to B=2. Compare TTFT across B levels only with that in mind; the aggregate column is the headline number.

qwen2.5-0.5b-bf16 (small dense bf16; isolates scheduler overhead)

Bok/failTTFT mean (ms)TTFT p95 (ms)decode tok/s per requestaggregate tok/sscaling vs B=1
11 / 027.827.8378.4352.21.00x
22 / 014.418.3249.1488.11.39x
44 / 021.733.1291.61118.13.17x

llama-3.1-8b-4bit (canonical dense 4-bit)

Bok/failTTFT mean (ms)TTFT p95 (ms)decode tok/s per requestaggregate tok/sscaling vs B=1
11 / 0173.1173.1104.892.41.00x
22 / 053.771.093.7181.71.97x
44 / 088.0139.478.6300.43.25x

qwen3-30b-a3b-4bit (MoE; batched decode hits the fused-MoE path)

Bok/failTTFT mean (ms)TTFT p95 (ms)decode tok/s per requestaggregate tok/sscaling vs B=1
11 / 0149.1149.1169.6142.51.00x
22 / 0192.0192.1118.3202.31.42x
44 / 0379.2379.565.5220.91.55x

Reading

The dense models scale close to linearly to B=4 (3.17x and 3.25x on aggregate throughput) while giving up 23-25% of per-request decode. The MoE model does not: it reaches only 1.55x aggregate at B=4, per-request decode falls 61% (169.6 -> 65.5), and TTFT rises 2.5x (149 -> 379 ms) despite the warm prompt cache that helps the dense rows.

The mechanism is in the source rather than in the measurement. The fused MoE decode kernel is gated on a single token: qwen3_moe.rs takes it only when array_shape(&x_flat)[0] == 1, where x_flat is [batch * seq_len, hidden]. The scheduler switches away from the per-sequence path at exactly B=2 (dispatch_sync_decode uses it only when seq_ids.len() <= 1) and execute_batched_decode builds its input as [b, 1], so from B=2 the gate fails on every layer of every tick and the model falls back to gather_qmm. That also explains why the scheduler is not the suspect: qwen2.5-0.5b-bf16 is in this trio precisely to isolate scheduler overhead, and it scales fine. Tracked as issue #1616.

No requests failed at any level for any model (0 fail across all 9 cells).

SSM / hybrid / mixed-cache families are intentionally absent: the server serializes their slots, so a B-ladder over them measures nothing.

Attribution was done on M1 Ultra, and the M5 Max re-run is pending. Issue #1616 read these rows as the MoE decode path declining the fused kernel at B>=2. Profiling on M1 Ultra found a different cause: Qwen3MoeModel never overrode forward_batched, so every batching family without that override ran the single-sequence forward once per row and got its aggregate only from overlapping independent graphs. The fix and the full attribution, including an op-level measurement showing the batched fused kernel the issue proposed loses to gather_qmm from n=4, are in moe-batched-decode-m1ultra-2026-09-04.md and in the M1 Ultra document. These M5 Max numbers predate that change and are left as measured; re-running this ladder on M5 Max is what would show its effect here.

Performance vs mlx-lm / mlx-vlm baseline (2026-05-19 benchmark campaign)

Stale baseline. This section is the 2026-05-19 campaign (mlxcel 0.0.28 vs mlx-lm 0.31.3 / mlx-vlm 0.4.4, MLX 0.31.2). It was not re-run for the 2026-07-11/12 0.4.0-rc.1 sweep, so the mlxcel columns here predate several MLX bumps and kernel changes. The current 0.4.0-rc.1 mlxcel numbers are in the per-family tables above; treat the percentages below as the last measured parity snapshot, pending a fresh mlx-lm/mlx-vlm baseline run at the 0.4.0-rc.1 pin.

Source CSVs (same M5 Max host, mlxcel 0.0.28 with --cooldown 15 --big-cooldown 15):

  • mlxcel: benchmarks/metal_m5max_2026-05-19.csv
  • mlx-lm: benchmarks/pylm_m5max_2026-05-18.csv (mlx-lm 0.31.3 dev checkout in https://github.com/ml-explore/mlx-lm)
  • mlxcel VLM: benchmarks/metal_m5max_vlm_2026-05-19.csv, benchmarks/metal_m5max_vlm_2026-05-20.csv (Gemma3n VLM entries)
  • mlx-vlm: benchmarks/pylm_m5max_vlm_2026-05-18.csv (mlx-vlm 0.4.4)

The M5 Max baseline sub-sweeps ran as part of the same continuous benchmark campaign and crossed calendar midnight. For public reporting, this campaign is grouped under 2026-05-19 even though the Python baseline CSV filenames carry 2026-05-18 dates. Numbers are decode tok/s. mlxcel vs mlx-lm is mlxcel / mlx-lm as a percentage; bold = mlxcel >= mlx-lm. FAIL cells are real load/runtime errors on that backend with this configuration. The mlx-lm checkout used for this run is ed1fca4 ("Thread local generation stream"); some text models fail on this snapshot.

Aggregate (text)

  • Comparable text pairs: 66 (models with >=5 generated tokens both sides)
  • mlxcel >= mlx-lm: 27 / 66 (41%)
  • mlxcel >= 90% parity: 62 / 66 (94%, the Phi-3.5, Gemma dense, and Jamba fixes raise four models past 90%)
  • Average mlxcel/mlx-lm: 99% (median 99%, range 72%-127%)

Aggregate (VLM, models with >=5 generated tokens both sides)

  • Comparable VLM pairs: 22
  • mlxcel >= mlx-vlm: 11 / 22 (50%)
  • mlxcel >= 90% parity: 18 / 22 (82%)
  • Average mlxcel/mlx-vlm: 102% (median 101%, range 74%-123%)

Text decode (tok/s)

Modelmlxcelmlx-lmmlxcel vs mlx-lm
Gemma-4-31b-it-nvfp47.17FAIL-
aya-expanse-8b-4bit110.55113.8797%
aya-vision-8b109.24FAIL-
baichuan-m1-14b-4bit64.7364.68100%
bunny-llama3-8b-4bit111.08FAIL-
command-r7b-4bit110.91110.67100%
deepseek-coder-1.3b-4bit178.03FAIL-
deepseek-r1-distill-7b-4bit123.76125.6399%
deepseek-v2-lite-4bit202.25215.0094%
deepseek-v3-4bit-FAIL-
ernie-4.5-0.3b-4bit1053.87FAIL-
exaone-3.5-2.4b-4bit282.35289.0198%
exaone4-1.2b-4bit424.44FAIL-
falcon-mamba-7b-4bit63.19140.1045%
gemma-2b-4bit217.38223.2797%
gemma-4-26b-a4b-it-4bit137.12141.0897%
gemma-4-31b-4bit28.5928.7999%
gemma-4-31b-it-4bit27.3428.7495%
gemma-4-e2b-it-4bit201.90FAIL-
gemma-4-e2b-it-8bit136.69FAIL-
gemma-4-e4b-it-4bit136.68FAIL-
gemma-4-e4b-it-8bit80.88FAIL-
gemma2-2b-4bit241.96241.76100%
gemma3-1b-4bit399.65388.52103%
gemma3-4b-4bit182.16181.66100%
gemma3n-e2b-4bit158.71FAIL-
gemma3n-e4b-4bit110.24FAIL-
gemma3n-e4b-bf1639.0548.7280%
glm4-flash-4bit104.30104.03100%
gpt-oss-120b-4bit114.03110.35103%
gpt-oss-20b-mxfp4172.33168.33102%
hunyuan-1.8b-4bit329.29349.9394%
hunyuan-large-4bit64.43FAIL-
hunyuan-moe-a13b-bf1664.09FAIL-
internlm2-7b-4bit117.25117.9899%
internlm3-8b-4bit101.23FAIL-
internvl3-1b661.48FAIL-
jamba-v0.1-4bit215.84219.3898%
llama-3.1-8b-4bit116.65117.4399%
llama-3.1-8b-bf1633.9334.2999%
llama-3.2-1b-4bit546.81578.6494%
llama-4-scout-17b-4bit48.59FAIL-
llava-1.5-7b-4bit124.52FAIL-
llava-interleave-qwen-0.5b-bf16403.59FAIL-
llava-next-mistral-7b-4bit122.79FAIL-
mamba2-1.3b-4bit184.69FAIL-
mimo-7b-4bit119.66118.85101%
minicpm-2b-4bit233.46228.46102%
minicpm3-4b-4bit131.00FAIL-
minimax-m2-3bit73.7668.94107%
ministral-3b-4bit223.03231.9296%
mistral-small-3.1-24b-4bit41.4141.49100%
mixtral-8x7b-4bit65.2066.0899%
molmo-7b78.74FAIL-
molmo2-4b64.09FAIL-
nemotron-h-30b-4bit177.18178.8099%
nemotron-nas-30b-4bit176.38178.3999%
olmo-1b-4bit243.15FAIL-
olmo2-7b-4bit116.88120.7997%
olmo3-32b-4bit29.1128.99100%
paligemma2-3b-6bit168.83FAIL-
phi-2-4bit79.60FAIL-
phi-3-mini-4bit207.89212.7498%
phi-3.5-mini-4bit204.63207.7998%
phi-3.5-moe-4bit115.20107.56107%
phi-3.5-vision-4bit163.61FAIL-
phi-4-4bit63.8662.28103%
pixtral-12b-4bit76.5674.95102%
qwen1.5-moe-a2.7b-4bit237.73237.50100%
qwen2-vl-2b-4bit273.84381.9872%
qwen2.5-0.5b-4bit682.41637.17107%
qwen2.5-0.5b-bf16404.68402.73100%
qwen2.5-7b-4bit126.36123.59102%
qwen2.5-7b-8bit68.9867.44102%
qwen2.5-vl-3b-4bit156.8398.53159%
qwen3-0.6b-4bit566.50651.1487%
qwen3-1.7b-4bit368.50384.8496%
qwen3-30b-a3b-4bit156.15147.22106%
qwen3-4b-4bit191.04190.94100%
qwen3-8b-4bit112.38113.4099%
qwen3-moe-4bit157.16146.51107%
qwen3-next-480b-4bit-FAIL-
qwen3-vl-2b-4bit365.31382.5096%
qwen3-vl-30b-a3b-4bit151.16146.87103%
qwen3-vl-32b-4bit27.5128.5196%
qwen3.5-0.8b-4bit517.47545.4595%
qwen3.5-27b-4bit32.5134.0595%
qwen3.5-2b-4bit320.84345.5993%
qwen3.5-35b-a3b-4bit151.63152.9699%
qwen3.5-4b-4bit166.56174.4595%
qwen3.5-9b-4bit98.50108.2791%
qwen3.5-9b-bf1629.9832.0993%
qwen3.6-35b-a3b-4bit147.56146.93100%
smollm-135m-4bit905.24711.54127%
smollm3-3b-4bit232.79239.1497%
solar-open-100b-4bit65.3666.3099%
stablelm-1.6b-4bit425.14423.68100%
starcoder2-3b-4bit216.48214.76101%

VLM decode (tok/s)

Modelmlxcelmlx-vlmmlxcel vs mlx-vlm
aya-vision-8b112.09FAIL-
bunny-llama3-8b-4bit112.24FAIL-
gemma-4-26b-a4b-it-4bit134.38136.5798%
gemma-4-31b-4bit23.4139.85-
gemma-4-31b-it-4bit27.2130.2090%
gemma-4-e2b-it-4bit217.32201.70108%
gemma-4-e2b-it-8bit133.74150.5189%
gemma-4-e4b-it-4bit134.10131.24102%
gemma-4-e4b-it-8bit76.2890.0085%
gemma3-4b-4bit159.58FAIL-
gemma3n-e2b-4bit151.36124.63121%
gemma3n-e4b-4bit106.0193.55113%
gemma3n-e4b-bf1636.9549.8874%
internvl3-1b601.50529.33114%
llama-4-scout-17b-4bit48.33FAIL-
llava-1.5-7b-4bit117.70FAIL-
llava-interleave-qwen-0.5b-bf16343.53345.08100%
llava-next-mistral-7b-4bit120.38FAIL-
ministral-3b-4bit195.22FAIL-
mistral-small-3.1-24b-4bit39.62FAIL-
molmo-7b84.9956471.65 (anomalous, 1 token)-
molmo2-4b64.0166.8096%
paligemma2-3b-6bit80.09124.55-
phi-3.5-vision-4bit168.77159.63106%
pixtral-12b-4bit69.71FAIL-
qwen2-vl-2b-4bit247.21279.5588%
qwen2.5-vl-3b-4bit156.83FAIL-
qwen3-vl-2b-4bit281.37FAIL-
qwen3-vl-30b-a3b-4bit56.38FAIL-
qwen3-vl-32b-4bit19.65FAIL-
qwen3.5-0.8b-4bit505.94410.96123%
qwen3.5-27b-4bit32.8433.4498%
qwen3.5-2b-4bit323.00318.14102%
qwen3.5-35b-a3b-4bit151.34128.80117%
qwen3.5-4b-4bit170.78166.46103%
qwen3.5-9b-4bit102.39102.48100%
qwen3.5-9b-bf1630.9231.4598%
qwen3.6-35b-a3b-4bit147.38123.70119%

mlx-lm fail categories (text)

The mlx-lm-side FAILs are unchanged from the 2026-05-18 baseline: unsupported architectures (deepseek-v3-4bit, internvl3-1b, molmo-7b), transformers config schema drift (exaone4-1.2b-4bit, the gemma-4-e{2,4}b-it-{4,8}bit and gemma3n-e{2,4}b-4bit family), tokenizer wrapper bugs (internlm3-8b-4bit), ModelArgs mismatch (mamba2-1.3b-4bit, phi-2-4bit, minicpm3-4b-4bit, Gemma-4-31b-it-nvfp4), VLM-only loaders routed through the text path (aya-vision-8b, bunny-llama3-8b-4bit, llava-*, paligemma2-3b-6bit, phi-3.5-vision-4bit), custom remote code refused (hunyuan-*, deepseek-coder-1.3b-4bit, ernie-4.5-0.3b-4bit), and one runtime crash (olmo-1b-4bit). These are mlx-lm/mlx-vlm regressions in the development checkout under references/, not silent mlxcel wins.

FP32 Promotion Audit

Short prompt A/B runs on 2026-05-18 used origin/main at 5ebc074 as the baseline and the branch as the candidate. Each row used:

mlxcel generate -m models/<model> -p "Hello, how are you today?" -n 20 --profile --no-chat-template

The intent is a hot-path regression/impact check, not a replacement for the 100-token full sweep above. The clearest gains are the MoE rows that still used the nkh,nk->nh expert combine contraction.

Modelmain prefill tok/sprefill tok/smain decode tok/sdecode tok/sDecode change
glm4-flash-4bit5.5415.8254.85108.23+97.3%
solar-open-100b-4bit9.517.9817.0442.72+150.7%
qwen3-vl-30b-a3b-4bit5.475.4660.2458.79-2.4%
gpt-oss-120b-4bit1.471.47114.33115.16+0.7%
qwen3-30b-a3b-4bit5.525.52159.80156.18-2.3%
qwen3.5-35b-a3b-4bit4.824.77136.19135.88-0.2%
qwen3.6-35b-a3b-4bit4.794.79133.75133.02-0.5%
mixtral-8x7b-4bit4.104.1369.3667.33-2.9%
phi-3.5-mini-4bit41.8885.28165.43167.31+1.1%
gemma3n-e4b-bf166.786.8711.2911.25-0.4%
qwen3.5-0.8b-4bit135.49179.83402.99405.63+0.7%
jamba-v0.1-4bit55.6551.80176.60174.16-1.4%
stablelm-1.6b-4bit63.66115.82394.16427.62+8.5%

Reading:

  • glm4-flash-4bit and solar-open-100b-4bit confirm the same FP32-promotion class as in remaining MoE expert-weight combines.
  • Qwen3/Qwen3.5/Qwen3.6 A3B, Qwen3-VL A3B, Mixtral, and gpt-oss are effectively guardrail-neutral in this short run. Their MoE combines now share the same dtype-preserving helper, but the previous contraction was not the dominant measured bottleneck for these rows.
  • Non-MoE guardrails (phi-3.5-mini, gemma3n-e4b-bf16, qwen3.5-0.8b, jamba, stablelm) did not show a decode regression from the compiled activation, softcap, scalar-helper, or intentional-FP32 comments and tests added for this audit.

SolarOpen Decode Sync Audit

removed the accidental FP32 expert-weight combine and raised solar-open-100b-4bit decode into the low 40 tok/s range, but still missed the >=85% mlx-lm decode gate. The remaining SolarOpen-specific difference was the Rust implementation forcing eval_all() after every decoder layer. That is useful for multi-token prefill graph size control, but in single-token decode it adds 48 GPU synchronizations per generated token. mlx-lm does not synchronize at each layer in the decode path.

The branch keeps per-layer eval for prefill and skips it only when the input sequence length is one token. Validation used the same direct real-model command:

mlxcel generate -m models/solar-open-100b-4bit -p "Hello, how are you today?" -n 100 --profile --no-chat-template
BuildPrefill tok/sDecode tok/svs mlx-lm 66.30 tok/s
origin/main after (616c470)9.1941.3562%
branch34.0265.6699%

This is +58.8% decode over current main and +298.4% over the original 16.48 tok/s issue baseline. The issue acceptance gate (>=56 tok/s) is met.

Moderate Gap Triage

The four rows were rechecked on the M5 Max after, and had landed on main. The original falcon-mamba row used a generic chat prompt that exits after <|im_end|> in both mlxcel and mlx-lm, so the useful comparison uses a raw code prompt that generates the full 100-token budget.

ModelTriageRefreshed mlxcel decodemlx-lm decodeResult
glm4-flash-4bitReal regression fixed by the already-merged MoE combine change111.09108.32103%
falcon-mamba-7b-4bitMeasurement artifact from early EOS on the generic chat prompt94.4994.32100%
starcoder2-3b-4bitAlready fixed on current main; dense transformer row now matches mlx-lm214.85213.95100%
qwen3.5-0.8b-4bitReal GatedDeltaNet decode overhead fixed here via fast RMSNorm q/k and gated norm paths535.43555.4396%

For the only row changed by this issue, qwen3.5-0.8b-4bit, the before/after measurement used:

mlxcel generate -m models/qwen3.5-0.8b-4bit -p "def fibonacci(n):\n    " -n 100 --profile --no-chat-template

Five back-to-back runs on current main before the patch decoded at 427.17, 428.86, 423.32, 423.65, and 422.78 tok/s (mean 425.16). The same five-run sequence after the patch decoded at 535.44, 535.26, 535.06, 536.12, and 535.25 tok/s (mean 535.43), a +25.9% decode increase and 96% of mlx-lm's 555.43 tok/s on the same prompt.

Known Issues

ModelIssuePriority
pixtral-12b (VLM)Intermittent slow VLM decode on M5 Max: repeated runs read either ~68 tok/s or ~20 tok/s, roughly 50/50, while the llava VLM control stays steady. The stated trigger no longer applies as written: the issue was attributed to pixtral's large ~4100-token image context, but #792 (aspect-ratio image processing, merged 2026-07-13) stopped upscaling the 224x224 fixture to a 1024x1024 square, so the same run is now a 213-token context and read 73.68 tok/s at 0.6.0. Re-characterize against the new context length before keeping this openLow
glm-5-4bit / glm-5.1-4bitNot a runtime defect. Investigated at the 0.6.0 sweep: both local checkpoints are interrupted downloads (glm-5-4bit = 21 GB of .incomplete blobs, no safetensors and no tokenizer; glm-5.1-4bit = .gitattributes only). Re-download, then re-test before filing anything against the GLM-5 load pathLow (data, not code)
hunyuan-moe-a13b-bf16 (bf16 A13B)Dropped: bf16 weights exceed the 128 GB budget; use hunyuan-a13b-instruct-4bit (4-bit, 64.92 tok/s). The size estimate passed it before it OOM'd at load, so the harness logged FAIL:bench instead of SKIP:oomLow
deepseek-v3-4bit99.96 GiB of weights on a 128 GB host. From the 0.6.0 sweep it is classified SKIP:oom_estimate under the 90 GB weight budget instead of being launched and recorded as FAIL:bench, which is what earlier sweeps did. A capacity exclusion, not a MoE + MLA defectLow (capacity)
qwen3-coder-480b-a35b-instruct-4bitOOM-skip on 128 GB; weights exceed the memory budget (the 480B Qwen3-Next was retired)Medium
qwen3-0.6b-4bitFull-budget raw prompt stays at ~93% of mlx-lm; sub-95% decode gapMedium
gemma-4-31b-it-nvfp4Now decodes at ~15.6 tok/s via the native NVFP4 Metal path (was ~7 tok/s at 0.2.1); still about half the 4-bit rate, so flagged ⚠️Low
falcon-mamba-7b-4bitGeneric chat prompt exits after <|im_end|>; use a non-chat code prompt for perf checksLow
phi-2-4bitGenerates only 1 token — likely EOS handlingLow
llama-3.1-8b-bf16bf16 → f16 conversion path is functional but slowLow

Notes

  • All tests use 4-bit quantized models unless noted.
  • Performance measured with mlxcel-bench-decode (model load, warmup, and measured pass in one process).
  • vs M1 Ultra ratios are M5 Max decode divided by the 2026-07-12 benchmarks/metal_m1ultra_2026-07-12.csv decode (same mlxcel 0.4.0-rc.1 / MLX pin 57c66cac / cooldown-30 conditions). Rows show - where the M1 Ultra sweep did not measure the model.
  • The 2026-07-11/12 sweep used --cooldown 30 --big-cooldown 30. Without cooldowns, heat accumulated over the larger 0.4.0-rc.1 model set thermally throttles the mid-sweep Qwen block (see Test Environment). Re-run full sweeps on this host with cooldowns.
  • Prefill and decode tok/s reported separately.
  • Current per-model values are the 2026-06-15 full sweep on mlxcel 0.2.1 (MLX pin a6ec7123): 151 text models (bench_decode.sh all) + 150 VLM-mode (all --vlm), bare run (pre-warm on, no cooldown). Source CSVs: benchmarks/metal_m5max_2026-06-15.csv and benchmarks/metal_m5max_vlm_2026-06-15.csv.
  • vs the 2026-05-27 sweep: 0 decode regressions among the 93 models measured in both; 11 improved >10% (MoE families from the fused-decode default plus broad MLX-bump gains, e.g. gemma2-2b +16.5%, gemma3-4b +12.5%, qwen3-30b-a3b +12.2%, qwen3-moe +12.1%). The two sweep FAILs first read as regressions both turned out to be environmental, not code (corrupt qwen2.5-vl checkpoint re-downloaded; oversized bf16 hunyuan dropped for the 4-bit); see Summary Statistics.
  • Measurement noise on very fast small models remains high (qwen3.5-0.8b-4bit and similar can span ±15% across back-to-back runs because 100 tokens generate in under 300 ms).

TurboQuant KV cache — M5 Max results

Note: The 2026-04-26 benchmark run (benchmarks/turbo_kv/2026-04-26_Mac.localdomain.csv) was performed on a development machine (Mac.localdomain), not on the reference M5 Max MacBook Pro. The hardware identity is unconfirmed; results may not be directly comparable to the M5 Max decode/prefill numbers above. A dedicated M5 Max run should be appended once available.

PPL evaluation throughput — 2026-04-26 run

The quality gate measures wikitext-2 PPL evaluation throughput (tok/s over a 4K-token window), which is distinct from the decode tok/s reported in the standard model tables above. These numbers characterize TurboQuant overhead on the MLX graph execution path, not peak generation throughput.

ModelKV modePPL eval tok/sWall clock msGate
Meta-Llama-3.1-8B-Instruct-4bitfp16733.76111,617baseline
Meta-Llama-3.1-8B-Instruct-4bitturbo4asym490.32167,034pass
Qwen2.5-1.5B-Instruct-4bit (superseded)fp163205.5425,550superseded —
Qwen2.5-1.5B-Instruct-4bit (superseded)turbo4asym2227.0936,775superseded —

The Qwen2.5-1.5B-Instruct-4bit rows above are retained for historical reference. A later run found that the fixture collapses on raw wikitext without a chat template; the B3 gate now uses the base variant Qwen2.5-1.5B-4bit. Re-run pending.

For the full interpretation and per-model recommendations see docs/turbo-kv-cache.md.

TurboQuant KV cache — M5 Max speed gate readings

First dedicated M5 Max reading of the TurboQuant KV speed gate matrix. Hardware: Apple M5 Max, 128 GB unified memory, macOS 26.4.1 (build 25E253). Model: mlx-community/Meta-Llama-3.1-8B-Instruct-4bit (local dir models/llama-3.1-8b-4bit). Date: 2026-05-03. Binary: mlxcel 0.0.25 post- (fused Sparse-V Metal kernel landed). Reproducer:

./scripts/bench_kv_cache.sh \
  --modes fp16,int8,turbo4-asym,turbo4,turbo4-delegated,turbo3-asym \
  --contexts 4096,16384 \
  --prefill-contexts 8192 \
  --decode-tokens 80 --warmup-tokens 16 \
  --run-cooldown 15 --mode-cooldown 30 \
  models/llama-3.1-8b-4bit

Full CSV at benchmarks/turbo_kv/2026-05-03_Apple_M5_Max_llama-3.1-8b-4bit.csv.

CSV schema note: Rows where stage=prefill record a single-token follow-up to force the KV cache to be populated. The resulting decode_tok_s value (e.g. 1200480 tok/s) reflects a sub-millisecond single-token step and is not a meaningful decode throughput figure; ignore it for prefill rows. Use prefill_tok_s from those rows and decode_tok_s from stage=decode rows.

Decode @ 4K context (80 generated tokens)

ModeDecode tok/s× FP16M5 Max gateVerdict
fp16101.291.000×baselinebaseline
int872.790.719×(no gate; tracking)tracking
turbo4-asym9.150.090×≥0.97×fail
turbo420.760.205×≥0.93×fail
turbo4-delegated27.280.269×≥0.97×fail (partial fix landed)
turbo3-asym6.360.063×(tracking only)tracking

caches the cold-V dequant graph across decode steps; informal in-tree A/B (100-token decode at the same 4K prompt) measures turbo4-delegated at ~41 tok/s post-fix vs ~27 tok/s on v0.0.25, a ~1.5× decode speedup that scales sharply at longer contexts.

Phase-1b (K-side unification): Removes cold_keys and the per-step concat(cold_k, hot_k) graph node. Informal A/B on M5 Max (3 warm runs each, llama-3.1-8b-4bit, 4109-token prompt, 100 generated tokens): fp16 baseline 101.5–102.7 tok/s; turbo4-delegated post- 43.0–43.7 tok/s (~0.43× FP16, up from ~0.41× pre-fix). The modest speedup is explained by SliceUpdate::eval_gpu semantics: MLX copies the full source buffer before writing the update region, so per-step K-side memory traffic is approximately conserved between the old concat layout and the new slice-update layout. The remaining cost is the V-side concat(cold_v_dequant, hot_v) graph node (Phase 2).

Phase-2 (fused dequant + SDPA kernel): Adds a Metal kernel that reads the packed cold V indices directly inside the kernel, removing the earlier cold_v_dequant_cache memo and the per-step concat(cold_v_dequant, hot_v) graph node. The dequantised cold V never materialises in global memory — V-memory budget stays at 4-bit packed.

Measured on llama-3.1-8b-4bit, 4109-token prompt, 100 generated tokens (`benchmarks/turbo_kv/2026-05-04_Apple_M5_Max_issue_528_fused_delegated_sdpa.csv$):

\text{Path}\text{tok}/\text{s}\times \text{FP16}\text{gate}
$fp16`101.761.000×baseline
turbo4-delegated default (no memo)29.600.291×≥0.97× — fail
turbo4-delegated fused kernel (MLXCEL_TURBO4_DELEGATED_FUSED=1)18.900.186×≥0.97× — fail

Removing the earlier memo (per the issue body's "creates dead state, must not remain" requirement) drops the legacy update_and_fetch route from ~0.43× to 0.29×. The fused kernel runs slower than the no-memo legacy route because the host pipeline now composes Q·K + softmax + cold-kernel + hot-matmul + sum out of many small MLX graph ops; the memo path could feed the dequantised cold V into a single steel-attention SDPA call. The 0.97× M5 Max decode gate is not cleared. Bringing the kernel inside the steel-attention envelope is left to follow-up work.

Phase-3 (steel-attention-envelope kernel, M5 Max measurement). Phase 3 lands turbo4_delegated_steel_sdpa — a JIT-compiled Metal kernel that runs the entire post-Q·K SDPA inline (per-Q numerically stable softmax, cold-V dequant + weighted sum, hot-V FP16 weighted sum, all normalised against the same softmax denominator). Bit-parity is gate-1 and was validated at PR-landing time (RMS < 5e-3 over 200 decode steps, two new parity tests in cache::turbo_tests). M5 Max throughput was deferred to a follow-up bench run because the kernel-author agent had no M5 Max access's run.

Measured on llama-3.1-8b-4bit, 4109-token prompt, 100 generated tokens (`benchmarks/turbo_kv/2026-05-06_Apple_M5_Max_issue_531_steel_envelope.csv$):

\text{Path}\text{tok}/\text{s}\times \text{FP16}\text{gate}
$fp16`102.971.000×baseline
turbo4-delegated legacy update_and_fetch + attention() (env unset)29.60*0.291×≥0.97× — fail (reading)
turbo4-delegated cold-only fused kernel (MLXCEL_TURBO4_DELEGATED_FUSED=1, earlier)18.90*0.186×≥0.97× — fail (reading)
turbo4-delegated steel envelope (MLXCEL_TURBO4_DELEGATED_FUSED=1, post-Phase 3)16.230.158×≥0.97× — fail

* cross-referenced from the 2026-05-04 CSV.

The steel envelope runs slower than both the cold-only fused kernel and the legacy fetch route. The likely cause is the kernel's single-thread-per-Q softmax + V accumulation pass — at decode time (Tq=1, B=1, Hkv=8 on llama-3.1-8b) only 8 threads are dispatched per kernel call, each scanning the full T_total range serially. The implementation note acknowledges this design ("single thread; T_total reads << kernel launch overhead, avoids threadgroup tree-reduction barriers") — the assumption held on M1 Ultra at parity contexts but breaks on M5 Max where the threadgroup tree-reduction would actually be faster than the serial scan.

Follow-up readings (Pass 1 parallelization + cold-loop sparse cutoff on top). The follow-up splits the kernel's Pass 1 (per-Q max + sum_exp) across all D threads of each threadgroup with a tree reduction. The same follow-up also precomputes a score-space sparse-V cutoff max + log(threshold * sum_exp), letting the cold loop reject fully-dead tokens before paying the exp + dequant cost. Pass 2's weighted-sum remains D-parallelized exactly as in Phase 3. Measured on llama-3.1-8b-4bit, 4109-token prompt, 100 generated tokens (`benchmarks/turbo_kv/2026-05-06_Apple_M5_Max_issue_534_post_fix.csv$):

\text{Path}\text{tok}/\text{s}\times \text{FP16}\text{gate}
$fp16`103.281.000×baseline
turbo4-delegated steel envelope earlier (reading)16.230.158×≥0.97× — fail
turbo4-delegated steel envelope post- (this fix)19.210.186×≥0.97× — still fail

Pass 1 parallelization plus the score-space cutoff nudges the steel envelope slightly past the cold-only fused kernel at 4K (19.21 vs 18.90) but does not move the needle far enough to clear the 0.97× gate. The residual cost is in Pass 2's per-token T_total scan. A simdgroup broadcast experiment was also measured during and regressed 4K decode, so it was not retained. The broader simdgroup-hybrid pattern from MLX upstream's metal::steel::SDPA (per-simdgroup simd_max / simd_sum plus per-simdgroup partial-sum accumulators in Pass 2) is the proposed next iteration; it would change the per-token per-thread T_total scan into a per-token per-simdgroup scan (4–8× fewer scans on M5 Max for D=128 / D=256).

Post- TurboQuant+ delegated FP16 working-set experiment (4K)

Follow-up on 2026-05-07 after comparing turboquant_plus: the MLX delegated KVCache keeps FP16 K/V in an internal native cache and routes decode through native SDPA, while packed storage is compacted outside the hot path. mlxcel now has an opt-in analogue via MLXCEL_TURBO4_DELEGATED_FP16_FAST_PATH=1. The follow-up handoff compacts the initial packed-V sidecars after prefill and before the first decode forward for max_tokens > 1, matching TurboQuant+'s compact_turbo_cache(...) shape without putting that cost in decode timing. Measured on llama-3.1-8b-4bit, 4109-token prompt, 100 generated tokens (benchmarks/turbo_kv/2026-05-07_Apple_M5_Max_issue_534_fp16_fast_path_predecode_compact_4k.csv):

Pathtok/sx FP16gate
fp16105.231.000xbaseline
turbo4-delegated steel envelope post-19.210.183x>=0.97x — fail
turbo4-delegated FP16 fast path + pre-decode compact104.090.989x>=0.97x — pass

The fast path is 5.4x faster than the post- steel envelope at 4K and 3.5x faster than the legacy update_and_fetch + attention reading (29.60 tok/s). It clears the 0.97x gate because the one-time sidecar compaction is no longer charged to the first decode forward. This remains a speed-path experiment, not the compressed-only memory target, because the full FP16 V working set is retained while the env var is enabled. The handoff is still visible in prefill timing for the decode-stage row: 2462.77 ms vs 1271.07 ms for FP16 at 4K.

Post- lazy sidecar policy experiment (4K)

Follow-up on the pre-decode handoff: MLXCEL_TURBO4_DELEGATED_FP16_SIDECARS=lazy skips foreground packed sidecar folds during generation and compacts missing sidecars only on preservation paths such as detach / prompt-cache donation. Measured on the same llama-3.1-8b-4bit, 4109-token prompt, 100 generated tokens (benchmarks/turbo_kv/2026-05-07_Apple_M5_Max_issue_534_fp16_fast_path_lazy_sidecars_4k.csv):

Pathdecode tok/sx FP16prefill_msgate
fp16104.511.000x1268.97baseline
turbo4-delegated FP16 fast path + pre-decode compact104.090.996x vs this FP16 run2462.77>=0.97x — pass
turbo4-delegated FP16 fast path + lazy sidecars102.150.977x1480.04>=0.97x — pass

Lazy sidecars keep decode within the gate while dropping most of the handoff cost introduced by pre-decode compaction. The remaining prefill delta versus FP16 is now ~211 ms at 4K instead of ~1194 ms.

Decode @ 16K context (80 generated tokens, fewer if early EOS)

ModeDecode tok/s× FP16GeneratedM5 Max gateVerdict
fp1663.581.000×19baselinebaseline
int836.350.572×80(no gate; tracking)tracking
turbo4-asym3.870.061×26≥0.95×fail
turbo46.760.106×80≥0.90×fail
turbo4-delegated3.410.054×21≥0.95×fail (— see below)
turbo3-asym1.850.029×54(tracking only)tracking

The repeated-paragraph prompt hits an EOS early on fp16, turbo4-asym, turbo4-delegated, and turbo3-asym at 16K; the per-token rate is computed over the actually generated tokens. int8 and symmetric turbo4 ran the full 80 tokens.

16K reading (50-token decode, no EOS early-exit)

Measured on llama-3.1-8b-4bit, ~16065-token prompt (`benchmarks/turbo_kv/2026-05-04_Apple_M5_Max_issue_528_fused_delegated_sdpa.csv$):

\text{Path}\text{tok}/\text{s}\times \text{FP16}\text{gate}
$fp16`74.251.000×baseline
turbo4-delegated default (no memo)6.030.081×≥0.95× — fail
turbo4-delegated fused kernel5.120.069×≥0.95× — fail

Same shape as the 4K reading: removing the earlier memo (requirement) regressed the legacy fetch path; the fused kernel is slower still. The gate is wider here because at 16K the dequant cost dominates; the per-step memo materialised ~52 MB / layer of FP16 cold V, which is gone, but the kernel cannot replace the steel-attention SDPA pipeline the memo enabled.

16K reading (early-EOS at 19 generated tokens)

Measured on llama-3.1-8b-4bit, 16163-token prompt, 100 requested decode tokens (early EOS at 19 on both modes — same prompt shape early-exits FP16 and steel envelope at the same point, so the per-token ratio remains fair). Same CSV as the 4K reading (`benchmarks/turbo_kv/2026-05-06_Apple_M5_Max_issue_531_steel_envelope.csv$):

\text{Path}\text{tok}/\text{s}\text{Generated}\times \text{FP16}\text{gate}
$fp16`63.94191.000×baseline
turbo4-delegated steel envelope (post-Phase 3)2.39190.037×≥0.95× — fail

The 16K decode ratio (3.7% of FP16) is the gate's worst-case shortfall in the TurboQuant KV speed gate matrix to date, ~1.4× worse than the cold-only kernel reading (5.12 tok/s, 0.069× FP16). At 16K the per-token serial scan over T_total is ~16K reads × 8 threads, completely dwarfing the tens of milliseconds the FP16 attention path needs for the same step.

16K reading (Pass 1 parallelization, early-EOS at 19–21 tokens)

Same prompt shape as the reading; FP16 early-exits at 19 and the post- turbo4-delegated path at 21 (one extra token before EOS). CSV: benchmarks/turbo_kv/2026-05-06_Apple_M5_Max_issue_534_post_fix.csv.

Pathtok/sGenerated× FP16gate
fp1664.78191.000×baseline
turbo4-delegated steel envelope earlier2.39190.037×≥0.95× — fail
turbo4-delegated steel envelope post- (this fix)2.99210.046×≥0.95× — still fail

The fixes move the 16K ratio from 3.7% to 4.6% of FP16 (a 25% relative improvement) but do not clear the gate. The residual gap is in Pass 2; see the simdgroup-hybrid follow-up note in the 4K subsection above.

Post- TurboQuant+ delegated FP16 working-set experiment (16K)

Same fast-path experiment as the 4K subsection, measured on the 16163-token prompt. Both modes early-exited at 19 generated tokens (benchmarks/turbo_kv/2026-05-07_Apple_M5_Max_issue_534_fp16_fast_path_predecode_compact_16k.csv):

Pathtok/sGeneratedx FP16gate
fp1665.55191.000xbaseline
turbo4-delegated steel envelope post-2.99210.046x>=0.95x — fail
turbo4-delegated FP16 fast path + pre-decode compact70.37191.074x>=0.95x — pass

This confirms the earlier fast-path bottleneck was the first-decode sidecar compaction placement. Once the initial sidecar fold runs during the handoff, decode uses the same unified FP16 K/V native-SDPA hot path as FP16 mode. The 16K run is short because of early EOS, so the >1.0x ratio should be read as FP16-class rather than a stable speedup claim. The handoff cost moved into the decode-stage prefill timing: 11070.22 ms vs 7952.33 ms for FP16 at 16K.

Post- lazy sidecar policy experiment (16K)

Same lazy-sidecar experiment as the 4K subsection, measured on the 16163-token prompt with early EOS at 19 generated tokens (benchmarks/turbo_kv/2026-05-07_Apple_M5_Max_issue_534_fp16_fast_path_lazy_sidecars_16k.csv):

Pathdecode tok/sGeneratedx FP16prefill_msgate
fp1664.15191.000x7852.26baseline
turbo4-delegated FP16 fast path + pre-decode compact70.37191.097x vs this FP16 run11070.22>=0.95x — pass
turbo4-delegated FP16 fast path + lazy sidecars66.63191.039x8155.47>=0.95x — pass

The 16K lazy policy removes nearly all visible sidecar handoff overhead from the decode-stage prefill timing: 11070.22 ms with pre-decode compaction drops to 8155.47 ms, close to the 7852.26 ms FP16 baseline.

Prefill @ 8K context (single-token decode follow-up)

ModePrefill tok/s× FP16M5 Max gateVerdict
fp162444.341.000×baselinebaseline
int82664.411.090×(no gate; tracking)tracking (faster than FP16)
turbo4-asym1680.450.687×≥1.00×fail
turbo41157.400.474×≥1.00×fail
turbo4-delegated2942.941.204×best-effortpass
turbo3-asym1579.360.646×(tracking only)tracking

int8 prefill on M5 Max is 9% faster than fp16 — consistent with the M5 Neural Accelerator's INT8 matmul path. turbo4-delegated keeps the prefill stage at FP16 by design and lands 20% above the FP16 baseline, again likely thanks to the INT8 KV write-back happening only after prefill commits.

M5 Max reading

The Turbo decode gates do not pass on the v0.0.25 binary as of 2026-05-03, on any of the three Turbo modes. The shortfall is largest on turbo4-asym$ (~10 \times \text{off} \text{the} 4\text{K} \text{gate}) \text{and} \text{smallest} \text{on} $turbo4-delegated (~3.6× off). Cross-checking against the 2026-04-29 M1 Ultra reading: M5 Max is roughly 1.4–2.4× faster than M1 Ultra on the same modes, but the headroom from M1's L2-bound regime is not large enough to recover the gates on its own.

A targeted A/B at 4K decode with MLXCEL_SPARSE_V_KERNEL=0 against the default kernel-on path:

ModeKernel ON tok/sGraph fallback tok/sΔ
turbo4-asym9.2318.51graph is 2.0× faster
turbo420.7620.68parity
turbo4-delegated27.2827.09parity

The fused Sparse-V Metal kernel is a measured regression vs. the graph reference for turbo4-asym on M5 Max — likely the per-thread skip path is paying more in kernel-launch and codebook-load overhead than it recovers from skipping below the 1e-6 threshold for an 8B-model decode workload at 4K. turbo4 and turbo4-delegated are at parity because both modes do an FP16 V-side write at decode time anyway, so the sparse-V path is largely inert. Even the faster graph fallback for turbo4-asym (0.183× FP16) is far below the 0.97× gate, so disabling the kernel is not a fix on its own.

M5 Max hardware considerations

  • turbo4-delegated is the only Turbo mode that meets any gate on M5 Max today — the 8K prefill reading (1.20× FP16). Use it when prefill latency matters and the cold-tail compression ratio is acceptable.
  • int8 is the recommended drop-in baseline for memory-constrained long-context workloads on M5 Max. It loses ~28% of decode throughput at 4K and ~43% at 16K against FP16, but prefill is 9% faster and the KV cache halves. No correctness regression has been observed on the Llama-3.1 family.
  • The Turbo decode shortfall is not the L2 wall observed on M1 Ultra. M5 Max has the headroom (M1U turbo4-asym 4K = 3.92 vs M5 = 9.15 tok/s with the kernel on; the graph path on M5 reaches 18.51) but the graph-level dequant cost still dominates. Closing the gates needs either a faster fused kernel or a structural change to fold the V-side dequant into the SDPA inner loop without per-token launch overhead.
  • Avoid turbo3-asym for decode-bound M5 Max workloads. The 3-bit unpack saturates the Metal command-queue overhead; the wall-clock decode rate is only 0.063× of FP16 at 4K and degrades further at 16K.

Acceptance criteria status

CriterionStatus
Decode + prefill numbers measured for 5 KVCacheModes (fp16, int8, turbo4-asym, turbo4, turbo4-delegated) at 4K decode + 8K prefill on M5 Maxdone
16K decode reading on M5 Max (primary M5 Max gate cell)done
turbo3-asym reading on M5 Max (tracking only per epic)done
32K decode readingdeferred — best-effort per epic; useful only after the kernel regression for turbo4-asym is investigated
Cross-hardware consistency check vs. M1 Ultradone — M5 Max numbers are 1.4–2.4× M1 Ultra on Turbo decode
CSV committed under benchmarks/turbo_kv/done
Docs summary in docs/model_tests_m5max.mddone (this section)
Failed-gate perf bug filedfollow-up — file an issue tracking (a) the turbo4-asym fused-kernel regression vs. graph fallback on M5 Max, and (b) the 0.27× ceiling on turbo4-delegated 4K decode

Deferred

  • 32K decode reading on M5 Max. Best-effort per epic; the present 4K/16K gap means a 32K reading would just deepen an already-failed gate without informing kernel work.
  • Multi-model expansion (Qwen 2.5, Gemma 3) — the gate matrix is keyed off Llama-3.1-8B per the epic, but per-family validation is open work.
  • A re-run after the turbo4-asym kernel regression is fixed; the gate matrix should be expected to pass at that point.