Skip to content

Releases: lablup/mlxcel

v0.8.0-beta.1

v0.8.0-beta.1 Pre-release
Pre-release

Choose a tag to compare

@inureyes inureyes released this 21 Oct 17:04
fbfed0a

First beta of the 0.8 line: 416 commits since v0.7.0. The full record is in CHANGELOG.md.

New Features

  • Experimental AMD GPU (ROCm) backend on Linux with --features rocm (#1818). Affine, mxfp4 and mxfp8 checkpoints run natively on gfx1151, NVFP4 converts to affine 4-bit at load (#2030), and the fused kernels with HIP ports (samplers, SSM update, Mamba1 scan, MoE decode, paged attention, add-RMSNorm, RoPE-append and others) run natively (#2098, #2099, #2100, #2101, #2103, #2107).
  • Bundled WebUI behind --webui on mlxcel serve and mlxcel-server: model library, local streaming chat, runtime activity and scoped settings (#1891, #1894, #1928, #1929, #2324). Guide: docs/webui/README.md.
  • mlxcel tag for BIO / BIOES token classification and redaction, including OpenAI Privacy Filter (#2456, #2460).
  • mlxcel extract for GLiNER2.5 span extraction and zero-shot classification on a DeBERTa-v2 encoder (#2461, #2521).
  • Responses API conversation compaction: POST /v1/responses/compact (#2520) and automatic compaction through context_management on /v1/responses (#2522).
  • Streaming tool-call arguments on Chat Completions and Responses for Hermes / Qwen JSON, Mistral, Llama 3 JSON and Qwen3-Coder XML (#2504).
  • /v1/realtime WebSocket for Nemotron VoiceChat full-duplex sessions (#2037, #2039, #2041).
  • Prompt-lookup decoding in mlxcel generate (#2074) and on the server with --draft-kind prompt-lookup (#2262, #2265, #2274).
  • Image generation runtime: mlxcel generate-image, /v1/images/generations and /v1/images/edits (#2497). No model family is registered yet; see Known Issues.
  • New model families: Cohere Compass / North-Micro-Vision (#1735), LLM-jp-4 VL and Jagle-VL (#1737), Laguna (#1739), GOT-OCR 2.0 (#1744), Kimi K3 text and MoonViT3D vision (#1741, #1768), IQuest-Coder Loop (#1757), Nemotron-Parse (#2034), Mage-VL 4B (#2036), nanochat (#2038), EmbeddingGemma 2 (#2316).
  • New speculative drafters: LFM2 / LFM2.5 DSpark (#1751), GLM-4.7-Flash MTP (#1753), Muse Glimmer DFlash assistant (#1762), Laguna DFlash (#1771).
  • Affine 1-bit checkpoints run with fused Metal kernels (#2035); before this they aborted on the first token.
  • Qwen 3.5 vendor fine-grained FP8 checkpoints load (#1742), and vision-stripped VLM checkpoints load on a text-only path (#2031).
  • OpenXLA / IREE paths for Qwen2.5-VL (#920), Qwen3-VL DeepStack (#921), Gemma 3 VLM masks (#917), Molmo (#918), Youtu-VL (#919) and Gemma 3n audio (#922).
  • mlx-lm LoRAEmbedding adapters on embed_tokens, fused and unfused (#2434).
  • Gemma 4 unified accepts video and audio in one prompt (#1740).

Improvements

  • One batch-native engine now drives every decode site. mlxcel run, mlxcel chat and mlxcel generate run on the same engine and request path as the server (#2217, #2224, #2225).
  • Inkling decode on inkling-small-mlx-4bit went from 8.10 to 44.07 tok/s on M3 Ultra (mlx-lm: 48.94 tok/s), with prefill from 122.6 ms to 48.3 ms (#2505).
  • The routed MoE gate/up pair runs as one fused Metal kernel through SwitchGLU::forward (#2514).
  • 11 families project only the sampled row through the LM head at prefill (#1955, #1969, #1971, #1973).
  • Mamba1 selective scan runs as a fused kernel on Metal (Jamba, Mamba, Falcon-Mamba) and on CUDA (Jamba) (#2006, #2009, #2082).
  • KV cache rows are written in place for dense, rotating and paged caches (#1961, #1963, #1965).
  • Model load reads and uploads weights before returning, staging through pinned memory on CUDA and ROCm (#2474, #2485).
  • Penalty, DRY, mirostat and adaptive-p requests stay on the pipelined decode path with token-identical output (#2256).
  • Server prefill chunk size adapts to decode load (#2268).
  • LFM2-VL caches vision features per image (#2453).

Bug Fixes

  • Configs saved by transformers 5 (nested rope_parameters) now load with their real RoPE base and scaling in about 80 config structs (#2448).
  • CUDA sorted gather_mm used the wrong row stride for MoE expert gathers, and captured child graphs raced with concurrent kernels on GB10 (#2477).
  • Batched DFlash on Qwen 3.5 / 3.6 / 3.8 emits the same greedy tokens as one request at a time (#2440).
  • Gated-delta q/k normalization uses the reference eps in Qwen3-Next, Qwen 3.5 and Kimi Linear (#2454).
  • Gemma 3n KV-shared layers, PhiMoE routing and longrope, the Gemma 4 audio mask and Phi-3-small padding ids follow upstream (#2457).
  • Tool-call parsing no longer drops or corrupts Qwen3-Coder, Mistral and pythonic calls, and model-supplied tool-call ids are returned (#2487).
  • Streamed and non-streamed chat output agree on 8 previously divergent cases (#2458, #2500).
  • Llama 3.x <|python_tag|> calls no longer leak into streamed content, and /v1/responses returns Harmony reasoning as a reasoning item (#2496).
  • SentencePiece-fallback tokenizers encode non-special added tokens as single ids (#2319).
  • A bf16 model on an fp16 TurboQuant KV cache no longer aborts on CUDA when attention carries an array mask (#2417, #2435).
  • Gemma 4 MTP exactness restored for 31B (#1987) and 12B long context (#2315); MTP verify is byte-identical to decode on CUDA (#2185, #2203).
  • Prompt cache: a whole-prompt hit no longer duplicates the last token (#2136), and completion snapshots are keyed by the tokens their state holds (#2139).
  • Paged KV blocks are reclaimed during decode under budget pressure (#2077, #2095), and the fused paged-attention kernels use 64-bit pool offsets (#2186).
  • Prefill state is correct for Kimi MLA, RWKV7, Bailing MoE Linear, RecurrentGemma and AFMoE (#2296, #2297, #2299, #2300, #2301).
  • Nemotron-H SSD prefill memory is bounded and requests are admitted against a checked budget (#2290, #2294).
  • Generated reasoning is re-injected into content-only history on Chat, Responses and Messages (#2116, #2121, #2125).
  • Array readback is fallible, and decode readback failures no longer lose sampler state (#2269, #2284).
  • 32 ROCm fixes from the backend bring-up, including thread safety of the JIT cache and BLAS handles and launch grids past the HIP cap (#2196, #2202, #2208, #2237, #2309).

CI/CD Improvements

  • Self-hosted gfx1151 runner builds, links and generates on ROCm (#1989), and make verify-rocm runs the gate locally (#2008).
  • Superseded CI runs are cancelled through a concurrency group, except on main (#1780, #1783).
  • The shared GB10 runner is skipped for changes it cannot build (#1993).
  • docs/environment-variables.md is checked against the source in CI (#2486), and the tests/ Python suites run in make verify and hosted CI (#2223).
  • Binary assets are no longer committed and a guard rejects new ones (#1904); license headers are gated in CI (#1728).

Technical Details

  • Engine design is recorded in ADR 0007 (unified batch-native engine), ADR 0008 (KV attention dispatch in the cache) and ADR 0009 (engine step API) under docs/adr/.
  • MLX pin advanced to bff10220, each overlay declares its base commit, and MLX-side overlays and custom kernels live under src/lib/mlxcelverse (#2298, #2304, #2305, #2505).
  • recipes/registry/0.8.0-beta.1.json is the architecture registry snapshot for this version.

Dependencies

  • rustls patched for the TLS handshake boundary vulnerability (#1896), and rustls-webpki and webpki roots moved past RUSTSEC-2026-0381, 0382 and 0383 (#2510).
  • The unmaintained YAML parser is replaced (#2279).
  • Minor and patch group updates (#1941, #2003, #2308) and tokio (#2450).
  • actions/upload-artifact 4 to 7 (#1940) and actions/setup-node 5 to 7 (#1879).

Breaking Changes

  • mlxcel generate decodes on the batch-native engine. PromptLookupGenerator and CxxGenerator are removed from the public API; use MlxInferenceSession or mlxcel_core::engine::DirectEngine (#2225).
  • The per-token pipeline diagnostics MLXCEL_PROFILE_PIPELINE, MLXCEL_PROFILE_PIPELINE_DETAIL, MLXCEL_TRACE_ASTYPE, MLXCEL_EXPORT_DECODE_DOT and MLXCEL_CAPTURE_DECODE are removed (#2225).
  • For Hermes, Mistral, Llama 3 and Granite, object-valued tool-call arguments are now the model's own JSON text instead of a re-serialization (#2504).
  • Model load now uploads weights before returning; MLXCEL_WEIGHT_MATERIALIZATION=lazy restores first-use materialization (#2474).

Known Issues

  • Image generation has no registered model family yet. mlxcel generate-image and mlxcel-server --image-model recognize image checkpoints and then refuse them (#2348 tracks the families).
  • The classic draft-model speculative path (--draft-model without --draft-kind mtp or dflash) prints a deprecation notice and is still present; removal is planned before v0.8.0.
  • On Metal, the sorted gather_mm numeric tests fail at 2.6 ulp against their tolerance; CI does not run them because it has no GPU.
  • The ROCm backend is experimental and validated on gfx1151 only.

What's Changed

  • fix(solar_open): pair the GPTQ biases dtype with the scales dtype by @inureyes in #1731
  • feat(bench): hold the comparison preconditions in one tool by @inureyes in #1732
  • fix(bench): move the store guards ahead of the output file, again by @inureyes in #1733
  • feat(vlm): add the Cohere Compass (North-Micro-Vision) VLM by @inureyes in #1735
  • feat(llmjpvl): port LLM-jp-4 VL and Jagle-VL by @inureyes in #1737
  • fix(internlm2): add family rope_scaling tests and NTK rescale log by @inureyes in #1736
  • feat(gemma4_unified): accept video and audio in the same prompt by @inureyes in #1740
  • fix(server): report an audio-only server healthy once its worker is ready by @inureyes in https://github.com/la...
Read more

v0.7.0

Choose a tag to compare

@inureyes inureyes released this 09 Sep 14:50

Two themes. llama-server b10621 compatibility is complete, with every one of 376 pinned options, routes and request fields carrying a terminal state and nothing deferred. And a re-measurement campaign across two Apple Silicon hosts found a dtype-promotion class that had been widening half-precision hidden states to f32 in six places, worth up to 3.26x on the checkpoints it reached.

Embeddings and reranking, DeepSeek-V4, and the Inkling family across text, image, video and audio also land in this release.

New Features

  • llama-server b10621 compatibility, complete (epic #1431). The pinned release (commit c1d0e7a0, published 2026-08-25) was reduced to 376 manifest entries across 18 shards under compat/llama-server/b10621/: 169 supported, 22 aliased onto an mlxcel spelling, 35 recorded as a deliberate divergence with a written rationale, 150 not applicable to a runtime that does not execute GGUF. A CI gate fails the build if an entry regresses. Covers HTTP transport, TLS, CORS and multi-key authentication (#1462, #1465), the native completion, embedding, tokenization, template and infill routes (#1467, #1480), the sampler remainder and a GBNF grammar engine (#1487, #1506), prompt cache, batching, context retention and YaRN (#1468, #1503, #1522), router mode (#1495, #1501), unfused LoRA adapters with live scales (#1497, #1520), mp3 and flac transcription streamed per token (#1523), embedding and reranking mode and pooling (#1483), speculative flags mapped to MTP and DFlash (#1488), the Vertex AI predict adapter (#1493), /props, /slots, /metrics, /health and slot persistence (#1492, #1524), idle sleep (#1528, #1530), and assistant prefill, echo and streamed thinking tags (#1526, #1527). Full map in docs/llama-server-compat.md.
  • Embeddings through mlxcel embed and POST /v1/embeddings (#1408, #1410-#1416). BERT/XLM-RoBERTa, ModernBERT, SigLIP text, EmbeddingGemma, Qwen3 and Qwen3-VL, bidirectional Llama/Nemotron/LFM2.5, Llama-Nemotron-VL, and ColBERT-style multimodal late-interaction models. Pooling, L2 normalization, Matryoshka dimensions, token and image inputs, bounded worker admission and per-checkpoint length limits are shared by the CLI and server paths, and the server can load an embedding checkpoint alone or beside a chat model with --embedding-model.
  • Reranking through mlxcel rerank and the Cohere/Jina-compatible POST /v1/rerank (#1417). One-label BERT/XLM-RoBERTa/ModernBERT cross-encoders, Qwen3 generative rerankers and Qwen3-VL multimodal rerankers.
  • DeepSeek-V4 (deepseek_v4) (#523). HyperConnections replace the plain residual stream, attention runs one shared 512-wide KV head across 64 query heads inside a 128-token rotating window with per-head softmax sinks, per-layer compression picks local, pooled or sparse-compressed attention, sparse layers add a HiSA indexer, and MoE routing hashes the first three layers by token id. Validated against mlx-community/DeepSeek-V4-Flash-4bit (43 layers, about 151 GB on disk) with a strict weight-coverage check at load and three real-checkpoint generation gates.
  • The Inkling family across four modalities (#1532, #1535, #1540, #1546, #1548). Hybrid sliding/global NoPE attention, learned banded relative-position logits, four f32 short-convolution states per layer, and logsigmoid-normalized routed and shared experts, loading original bf16/f32 weights, native ModelOpt NVFP4 expert planes, and pre-converted affine MLX 4-bit experts. On top of it: HMLP image input with the reference 40x40 tiler, adjacent-frame video under one request-wide 16-pair budget, dMel audio, and a native chained-MTP drafter.
  • IQuest-Coder (iquestcoder) and text-only Youtu-LLM (youtu, youtu_llm) (#1594, #1593), both new ModelType arms on existing loaders with no new architecture code.
  • Locally typical sampling (typical_p) and top-n-sigma logit filtering (#1482, #1479), applied as row-wise graph transformations before the single fused sampling dispatch so requests stay on the batched path.
  • mlxcel inspect --json and mlxcel arch --json (#1605, #1606, #1608). The first emits the memory estimator's byte fields as a stable contract; the second derives a per-family registry of runtime, modality, backend, distributed, drafter and KV facts. make recipes-registry writes a versioned snapshot under recipes/registry/.
  • A live runtime settings endpoint (#1516), per-request speculative acceptance on the response (#1588), tool_choice accepts required and a named function (#1319), Qwen-VL video input and Responses-native image parts (#1510, #1310), and LFM2-VL image splitting (#1405).
  • Chat Completions reasoning under both common field names (#1308) and OpenAI-shaped reasoning controls driving chat-template thinking (#1307). reasoning_effort, Responses reasoning.effort and compatible extra-body values derive enable_thinking.
  • A generation that goes entirely to the reasoning channel says so (#1721). mlxcel generate suppresses the <think> channel by default, so a reasoning model whose generation ends before the channel closes printed a blank line, which is indistinguishable from a broken checkpoint. Both the one-shot flow and the chat REPL now say how many tokens went there. scripts/ab_output_equality.sh is the output half of an A/B built on the same reading.

Improvements

  • Half-precision paths keep their dtype. A helper that builds its constants as f32 returns f32 for a half-precision input; the widened activation enters the residual stream at the first layer that uses it, and every later matmul promotes its own weight to match. Nothing is arithmetically wrong, which is why unit tests, output comparisons and perplexity all pass over it. Five bridge activation helpers were doing this (#1709), which took gpt_bigcode-santacoder from 28% to 90% of mlx-lm (51.23 to 164.63 tok/s) and pythia-1b from 31% to 97%; the multimodal and 3D rotary paths dropped the restore upstream performs (#1714), worth up to 2.45x on M1 Ultra and 3.26x on M5 Max; Gemma3n held two dtypes in one checkpoint and now loads uniformly bf16, moving decode from 78% of mlx-vlm to 97%; and Mistral 4, Ministral 3 (#1711) and the GraniteMoeHybrid gated norm (#1723) each carried a local copy.
  • Pre-Ampere CUDA gets three arms it did not have. bf16 weights convert to f16 at load (#1542), since no pre-Ampere part has a bf16 ALU. qmm_naive gets a Volta tensor-core path (#1543): prefill goes from 138 to 536 tok/s on qwen3.8-27B-4bit and 326 to 1579 on gemma-4-12B-it-4bit, against scalar-path ceilings of 259 and 583. qmv blocks output rows per warp, dividing activation traffic while leaving weight traffic and per-row arithmetic untouched: gemma-4-12B-it-4bit decode 44.7 to 57.4 tok/s, generated text byte-identical at --temp 0. Separately, qmv accumulates bf16 in float below Ampere at every bit width (#1539), which is 1.87x on qwen3.8-27B-4bit and 1.73x on gemma-4-12B-it-4bit on a V100, and qmm_naive sizes its CTA tile against the device's real shared-memory budget instead of an architecture proxy (#1541).
  • apply_token_bias no longer synchronizes the device on every decode step. It read the pre-bias argmax id back to the host to feed two suppression counters, which breaks the async lookahead pipeline. Measured on M1 Ultra as the delta between --ignore-eos and no flag: +65.4% decode time on qwen3-0.6b-4bit, +30.4% on llama-3.2-1b-4bit, +24.0% on qwen3-30b-a3b-4bit, +13.6% on llama-3.1-8b-4bit. Behind MLXCEL_LANG_BIAS_COUNTERS=1 the same deltas are within 1.2% of zero. Only requests using logit_bias, ignore_eos or --lang-bias were affected.
  • Constrained decoding keeps the grammar mask packed (#1578). The token bitmask stays in u32 form from the matcher to the GPU. On the Qwen3.8-27B geometry that removes two 248k-element host loops and 97% of the per-step host-to-device copy, on every tick of every constrained sequence.
  • The GQA repeat is gone from the fused SDPA paths (#1686), Qwen2-VL takes the text-only path when no image is in play, grouped softcap-GQA decode is on by default, Qwen3 MoE batches the decode forward instead of one forward per row (#1616), and the Mamba2 hybrids drop two M5-gated eval boundaries (#259, #266, #1685).
  • Chat templates are compiled once and cached (#1518), response stores are bounded by bytes rather than entry count (#1519), and the JSON body limit is derived from the image budget (#1515) instead of a fixed number that either truncated valid image requests or admitted oversized ones.
  • mlxcel serve carries the help sections that only mlxcel-server had (#1628), and a test asserts that every flag named under src/commands appears in some --help (#1131).
  • OpenXLA compiles a set of context capacities and routes each request to the smallest that fits (#1302).

Bug Fixes

  • Half-precision reductions in the bridge accumulate in f32 (#997). max over a bfloat16 array returned NaN for a finite input on M5 Max, about one call in six. Reading the same logits row back to the host found every element finite on 60 of 60 runs while the device reduction over that row reported non-finite on 28, with no overlap. Measured one variant per process over 150 forwards: 23 non-finite for the plain reduction, 31 reshaped to 1-D, 20 made contiguous, 24 reduced over an axis, 0 cast to f32, so element width is the trigger and contiguity is not.
  • Four RoPE scaling defects across shared decoder paths. Qwen3 ignored a configured rope_scaling (#1398), the shared Llama 3 path dropped the scaled frequencies, Gemma 3 skipped the linear factor on global-attention layers, and InternLM doubled positions under dynamic NTK. Phi-3 selects its LongRoPE table by whole-prompt position (#1580) rather than by the current chunk.
  • **nemotron_nas resolves llama3 RoPE scaling through the...
Read more

v0.7.0-beta.1

Choose a tag to compare

@inureyes inureyes released this 04 Sep 00:49

Pre-release. 158 commits since v0.6.0. Full record in the changelog.

The headline is epic #1431: the llama-server b10621 compatibility surface is now completely classified, with no deferred entries left. Alongside it, the Inkling family arrives across four modalities, and the pre-Ampere CUDA path gets three correctness and performance fixes measured on a Tesla V100.

New Features

  • llama-server b10621 compatibility, complete (#1431). llama.cpp release b10621 (commit c1d0e7a0, published 2026-08-25) was extracted into a machine-checked manifest under compat/llama-server/b10621/: 249 help entries, 323 long-option spellings, 136 environment variables, 53 routes and 74 native request fields, reduced to 376 entries across 18 shards. Every entry carries a terminal state: 169 supported, 22 aliased onto an mlxcel spelling, 35 recorded as a deliberate divergence with a written rationale, 150 not applicable to a runtime that does not execute GGUF. None deferred, and a CI gate fails the build if an entry regresses.
  • Native routes and utilities (#1462, #1465, #1467, #1480, #1491, #1492). HTTP transport, TLS, CORS and multi-key authentication; /completion, /embedding, /tokenize, /detokenize, /apply-template and /infill; resumable streams; /props, /slots, /metrics, /health and slot persistence.
  • The sampler remainder and a GBNF grammar engine (#1487, #1506). Mirostat v1 and v2, dynamic temperature, adaptive-p, min_keep, logit_bias, DRY breaker-string semantics, n_probs and post_sampling_probs, the uint32 seed fold; --grammar, --grammar-file, --json-schema, --json-schema-file, and the native json_schema field with its grammar alias, lazy grammars, triggers and preserved tokens.
  • Locally typical sampling (typical_p) and top-n-sigma logit filtering (#1482, #1479), both applied as row-wise graph transformations before the single fused sampling dispatch.
  • Router mode and multi-adapter LoRA (#1495, #1501, #1497, #1520). Checkpoint discovery from a model directory, a model cache and INI presets, served from an in-process model pool; multiple adapters stay unfused for per-request or hot-swapped scale changes.
  • mp3 and flac transcription, streamed per token (#1523). The container is identified from the clip's own magic bytes rather than the client's format string.
  • Vertex AI custom-container routing (#1493) behind AIP_MODE=PREDICTION, and idle sleep (#1528, #1530): --sleep-idle-seconds drops the serving worker and its weights, and the next request reloads.
  • The Inkling family, across four modalities (#1532, #1535, #1540, #1546, #1548). Hybrid sliding/global NoPE attention, learned banded relative-position logits, four f32 short-convolution states per layer, and logsigmoid-normalized routed and shared experts; plus HMLP image tiling, adjacent-frame video, dMel audio, and a native chained-MTP drafter.
  • IQuest-Coder (iquestcoder) and text-only Youtu-LLM (youtu, youtu_llm) (#1594, #1593). Both checkpoints previously failed with Unsupported model type; each is now a new ModelType arm on an existing loader with no new architecture code.
  • DeepSeek-V4 (deepseek_v4) (#523). HyperConnections replace the plain residual stream, one shared 512-wide KV head serves 64 query heads inside a 128-token rotating window, per-layer compression picks local, pooled or sparse-compressed attention, and MoE routing hashes the first three layers by token id. Validated against mlx-community/DeepSeek-V4-Flash-4bit (43 layers, about 151 GB on disk).
  • Embeddings and reranking (#1408 through #1417). mlxcel embed, mlxcel rerank, POST /v1/embeddings and POST /v1/rerank, covering BERT/XLM-RoBERTa, ModernBERT, SigLIP text, EmbeddingGemma, Qwen3 and Qwen3-VL, bidirectional Llama/Nemotron/LFM2.5, Llama-Nemotron-VL, and ColBERT-style late-interaction models.
  • Reasoning field compatibility (#1307, #1308). An OpenRouter-compatible reasoning alias alongside reasoning_content, and OpenAI-shaped reasoning_effort / reasoning.effort values driving chat-template thinking behavior.
  • Qwen-VL video input, Responses-native image parts, LFM2-VL image splitting (#1510, #1310, #1405).
  • A live runtime settings endpoint (#1516) and per-request speculative acceptance on the response (#1588): drafter kind, verify rounds, and drafted and accepted token counts.
  • mlxcel inspect --json and mlxcel arch --json (#1605, #1606, #1608). The memory estimator's byte fields as a stable contract, and a per-family registry of runtime, modality, backend, distributed, drafter and KV facts. make recipes-registry writes a versioned snapshot under recipes/registry/.

Improvements

  • CUDA qmv accumulates bf16 in float below Ampere, at every bit width (#1539). On a Tesla V100-PCIE-32GB, qwen3.8-27B-4bit decodes at 117.83 ms/token against 219.96 before (1.87x) and gemma-4-12B-it-4bit at 71.66 against 124.26 (1.73x). In an nsys profile at identical launch counts qmv runs 2.01x faster and accounts for 99.4% of the end-to-end gain.
  • CUDA qmm_naive sizes its CTA tile against the device's real shared-memory budget (#1541). A tile larger than the 48 KB a launch gets without cuFuncSetAttribute now opts in first, which f32 activations at group size 128 need and never had; a tile that does not fit is refused at the launch site instead of reaching the driver. The wide tile was measured on Volta and refused: 1.50x slower at a 106-token prompt with no occupancy gain.
  • Constrained decoding keeps the grammar mask packed (#1578). The bitmask stays in u32 form from the matcher to the GPU. On the Qwen3.8-27B geometry that removes two 248k-element host loops and 97% of the per-step host-to-device copy from the scheduler thread.
  • Chat templates compiled once and cached (#1518), response stores bounded by bytes (#1519), and the JSON body limit derived from the image budget (#1515).
  • OpenXLA compiles a set of context capacities and routes each request to the smallest that fits (#1302).

Bug Fixes

  • Youtu-VL restores raster order after the windowed vision tower (#1600). reverse_window_indices reproduced the window permutation rather than its inverse, so 192 of a 512x512 image's 256 merged tokens reached the language model out of place.
  • Qwen2.5-VL runs its vision tower norm in f32 and restores raster order correctly (#1601), with the raw Hugging Face patch-embed layout normalized at load (#1582).
  • CUDA grouped GEMM no longer tells CUTLASS a Volta part is a Turing part (#1544). The pre-Ampere arm named the m16n8k8 MMA that Turing introduced and sm_70 does not have. Nothing was computing wrong, and the retag moves no device code on any architecture: byte-identical SASS at compute_70, compute_80 and compute_121.
  • Four RoPE scaling defects across shared decoder paths. Qwen3 ignored a configured rope_scaling (#1398), the shared Llama 3 path dropped the scaled frequencies (#1385), Gemma 3 skipped the linear factor on global-attention layers (#1386), and InternLM doubled RoPE positions under dynamic NTK scaling (#1389).
  • Phi-3 selects its LongRoPE table by whole-prompt position (#1580), and Falcon-Mamba applies the B/C/dt RMS norm once (#1574).
  • mlxcel_core::is_gpu_available() is renamed default_device_is_gpu() (#1421), with gpu_backend_available() answering whether a GPU backend exists. The old name read as a hardware query but reported the current default device, so it turned false after any set_default_device(false).
  • Recurrent decode state is persisted across requests (#1513), KV modes reach model-owned caches (#1400), and symmetric Turbo4 stops reading past its sign vectors on MLA latent caches (#1395).
  • Sampling filters on the untempered distribution and scales by temperature last (#1391).
  • The prompt cache no longer donates or adopts KV-less shadow paged entries (#1390), and preemption victim selection has a total order (#1301).
  • LoRA adapters that do not map onto the loaded model are refused (#1576) instead of loading and contributing nothing.
  • Chat-template failures surface: a rejection is no longer swallowed (#1511), a missing map key returns None (#1394), and tojson matches json.dumps (#1382).

CI/CD Improvements

  • The Volta sm_70 baseline is recorded and gated in CI (#1556), so the pre-Ampere arms cannot regress unmeasured. The Volta first-token fixed cost is instrumented rather than guessed (#1545).
  • The Python client is tested across CPython 3.9 through 3.13 (#1572).
  • CI links an OpenXLA binary so a link-only regression cannot reach main (#1305), and the OpenXLA warning backlog is cleared with warnings denied (#1381).
  • Nightly verify-test failures are stabilized (#1256), ThinLTO is disabled in test-fast (#1407), and the TP crossover and distributed pipeline LRU eviction assertions are made deterministic (#1512, #1502).
  • webpage-build is guarded against missing MkDocs sources (#1570), and cross-repo ref boundaries are derived live rather than pinned (#1397).

Technical Details

  • Every workspace crate carries 0.7.0-beta.1: mlxcel, mlxcel-core, mlxcel-mlx-pin, mlxcel-surgery, mlxcel-xla. make verify-versions asserts it.
  • The b10621 manifest lives in 18 JSON shards with a schema_version of 4 and a pinned upstream archive SHA-256, so the compatibility claim is checkable rather than asserted.
  • recipes/registry/0.7.0-beta.1.json is the architecture registry snapshot for this release, with recipes/registry/CURRENT pointing at it.
  • The MLX pin is unchanged from v0.6.0.

Dependencies

  • Added symphonia 0.5 with only the mp3 and flac codecs enabled, ...
Read more

v0.6.0

Choose a tag to compare

@inureyes inureyes released this 22 Aug 07:46

Speculative decoding is the theme: MTP now measures the things it used to assume, and reports what it measured. The OpenXLA backend graduates to alpha and rejoins the workspace version line.

New Features

  • GET /v1/internal/mtp-policy reports whether the B=1 MTP burst runs and why, as a versioned read interface (#1257, #1268). States are settled, profiling, forced, unavailable and exactness_declined. The same state rides /health under observability.mtp_policy. See docs/mtp-policy-api.md.
  • Draft trees behind MLXCEL_MTP_TREE, off by default (#1204, #1212, #1214). Includes the per-node RoPE position fix a branching tree needs. The measured verdict is negative and the flag ships off: a linear tree costs 1.9% to 4.1% against the chain, branching 8.1% to 10.1%, against an oracle ceiling of about 6.5%.
  • MLXCEL_MTP_BLOCK_CONTROLLER selects the verify-width controller: the measured-throughput comparator (default), proxy for the previous acceptance gate, requested to pin a width for sweeps.
  • scripts/bench_speculative.sh measures speculative decoding under the protocol the published numbers use: baked-in prompts, a quiet-host gate, ABBA alternation, and a spread limit above which a run is reported as untrustworthy rather than averaged (#1215).

OpenXLA backend, now alpha

The OpenXLA / StableHLO backend leaves the experimental shelf and ships as alpha. It stays off by default behind xla-backend and xla-iree; alpha means the surface is real enough to run and report bugs against, not to deploy.

Its crate was exempted from the workspace version check as a default-off backend outside the release contract, so mlxcel-xla sat at 0.4.0 while 23 pull requests landed in it and none reached a release note. That exemption is gone and the work is collected in the changelog: multimodal execution through IREE (LLaVA #913 / #897, Qwen2-VL #915, Phi4MM audio #914, Gemma3n #892, sparse DeepStack #893, multimodal RoPE state #894, Molmo2 indexed attention pooling #916, image requests in CLI and continuous-batch serving #895, image context floors #1280), session and prefill plumbing (#879, #880, #888), operator numeric contracts and probes (#934, #937, #938, #940, #941, #942, #943, #944), and fixes (#945, #968, #1169, #1273, #1275). CI compiles the feature combinations (#1282), which is what keeps the list from going stale again.

Improvements

  • Gemma 4 MTP is gated on a measured exactness probe instead of an unconditional yes (#1188, #1258). On Apple GPU generation 15+ the default kernel selection breaks temperature-0 byte-identity; the gate now probes, and where the probe fails it selects the exact kernel instead. Measured on M5 Max: 93.2 tok/s against 43.8 classic (2.13x) with byte-identity, where the fast kernel reads 120.4 (2.76x) without it.
  • The B=1 MTP burst is gated on GPU generation, not the Neural Accelerator (#1217). M3 Ultra, which the old proxy lumped with M1 Ultra, measures 1.95x to 2.65x and now enables by default.
  • The MTP verify width is decided by measured throughput (#1207). The old proxy waited for the configured prefix to be fully accepted, which the Gemma 4 12B pairing never reaches. Measured on M5 Max at pinned widths: 89.15 tok/s at 4, 94.91 at 5, 85.96 at 8; the controller adopts 5 and refuses 8.
  • Speculative rollback stops rewriting the whole KV cache every round (#1209). A dense tail trim moves an offset instead of copying the live window: +0.9% at 1k context, +2.5% at 8k, +5.3% at 24k, growing with the conversation. Output is byte-identical.
  • The MTP accept hook builds its paired-hidden block with one slice, and multi-token drafter forwards stop materializing a causal mask per call (#1185).

Bug Fixes

  • /v1/internal/mtp-policy no longer reports profiling forever when the exactness probe vetoes MTP (#1298). The veto is now exactness_declined, carrying the probe's reason.
  • A --lang-bias-config YAML file with two or more languages steers deterministically, and a repeated language code is rejected instead of silently taking the last one (#1267).
  • Five more silent HashMap iteration-order defects: RT-DETRv2 sanitization (#1281), distributed registry ordering (#1284), pipeline cache eviction (#1288), the server's four LRU min_by_key calls (#1299), and synthetic test fixtures (#1266).
  • The MTP round-loop diagnostics report the width the loop actually drafted at, not the requested one (#1206, #1208).
  • macOS release signing is gated on a Developer ID certificate (#1216), so a build cannot ship assets signed with an identity Gatekeeper blocks.

CI/CD Improvements

  • The Metal test gate runs single-threaded like the CUDA one (#1210), removing a workspace-parallel SIGSEGV that made the nightly red with nothing wrong in the code.
  • The test harness pins f32 GEMMs to full precision (#1259, #1260), fixing 17 algorithmic-equivalence tests that failed on generation 17 because MLX dispatches f32 matmuls to a TF32-class kernel by default.
  • CI compiles the OpenXLA feature combinations (#1282) and gates clippy at PR time (#1285, #1292).

Technical Details

  • Every workspace crate now carries the root version; scripts/ci/check_crate_versions.py has no exemptions left, and a new member fails CI until someone decides which line it is on.
  • The early-exit verify walk was evaluated and not built (#1179). The LM head is weight-read-bound, so a sequential walk re-reads the full 566 to 715 MB weight per position and costs 1.2 to 1.6x the batched projection at production widths. MLXCEL_ENABLE_MTP_DEFERRED, which gated a path that deferred nothing, is removed. Record: docs/benchmark_results/mtp-verify-early-exit-decision-m5max-2026-08-22.md.
  • The qmv_wide narrow pin's collateral on batched decode is priced in docs/benchmark_results/qmv-wide-pin-tax-m3ultra-2026-08-22.md (#1261, #1278). #1289 tracks the order-preserving streamed kernel that would remove the speed-versus-exactness tradeoff.

Dependencies

  • h2 to 0.4.16 for RUSTSEC-2026-0258 (#1213).
  • The minor-and-patch group across 11 updates (#1162), plus a cargo package refresh.

Breaking Changes

  • Gemma 4 MTP throughput on Apple GPU generation 15+ drops about 23% against v0.5.2, in exchange for temperature-0 output that is byte-identical to classic decode. Set MLXCEL_MTP_ALLOW_INEXACT=1 together with MLXCEL_QMV_WIDE=1 to restore the previous behavior and forfeit the contract.
  • MLXCEL_ENABLE_MTP_DEFERRED is removed (#1179). It gated a verify path with identical compute and one extra bridge crossing; nothing replaces it.
  • mlxcel-xla jumps 0.4.0 to 0.6.0. The crate is not published, so this affects only builds that pin the path dependency by version.

Known Issues

  • The Gemma 4 31B + bf16 assistant pairing declines MTP on generation 15+ (#1279). Its exactness probe fails under both kernel selections on every host measured (M3 Ultra and M5 Max), so the pairing serves classic decode: correct output, and a measured 2.25x left on the table. The offline CLI reports this as an error; the server logs it at startup and reports exactness_declined at /v1/internal/mtp-policy. The mechanism is not the qmv_wide split the gate was built around, since the divergence survives the narrow retry and its byte count does not vary with block width.
  • RWKV7, RecurrentGemma and KimiLinear rebuild their recurrent state every decode step (#1220), so generation runs with no carried history. Pre-existing, not a regression in this release.
  • A --draft-block-size above the qmv batch limit forfeits byte-identity on every generation, because the matrix-matrix kernel takes over. The probe catches it and declines.

What's Changed

  • fix(speculative): report the block width the MTP round loop actually used by @inureyes in #1208
  • chore(deps): bump h2 to 0.4.16 for RUSTSEC-2026-0258 by @inureyes in #1213
  • fix(ci): run the Metal test gate single threaded, like the CUDA one by @inureyes in #1210
  • feat(speculative): draft and verify a token tree, behind a flag, default off by @inureyes in #1212
  • feat(speculative): branch the draft tree, and verify it at its own positions by @inureyes in #1214
  • fix(ci): gate macOS release signing on a Developer ID certificate by @inureyes in #1216
  • feat(bench): measure speculative decoding properly, and record what it measures by @inureyes in #1215
  • docs: update SECURITY.md supported versions and bug-report placeholder by @DebayanSen96 in #1229
  • deps(deps): bump the minor-and-patch group across 1 directory with 11 updates by @dependabot[bot] in #1162
  • perf(speculative): single-slice the MTP accept hook and drop the drafter's materialized causal mask by @inureyes in #1253
  • perf(cache): make the dense KV tail trim a logical rewind instead of a buffer rewrite by @inureyes in #1255
  • fix(speculative): probe Gemma 4 MTP exactness instead of assuming it by @inureyes in #1258
  • test(harness): pin f32 GEMMs to full precision, and accept minijinja's Python-style bools by @inureyes in #1260
  • perf(speculative): gate B=1 MTP on GPU generation, not Neural Accelerator (#1217) by @inureyes in #1254
  • docs(make): record the TF32 pin at the verify-test-cuda gate definition by @inureyes in #1262
  • fix(make): correct the TF32 note, klear reproduces too and is flaky by @inureyes in https://github.com/...
Read more

v0.5.2

Choose a tag to compare

@inureyes inureyes released this 17 Aug 17:39

Two correctness fixes for M5-class hardware, and MTP speculative decoding now works on the generation where it used to decline.

If you run a hybrid or recurrent model on an M5, upgrade. Qwen 3.5, Mamba, Jamba, RWKV, Nemotron-H, Falcon-H1, Kimi Linear, LFM2, Plamo2 and GraniteMoeHybrid were served with a corrupted recurrent state whenever the prompt was not already a multiple of 32 tokens. Nothing errored. The output just changed.

New Features

  • MLXCEL_METAL4_ATTENTION=0 turns off the M5 neural-accelerator fused attention route (#1065). An off-switch only: the hardware test still gates the route, so a stray value on an M1 is inert. Both copies of that hardware test now go through one shared predicate, because a half-applied switch produces a run that looks like the route is off and is not. Measured on M5 Max across llama-3.2-1b-4bit, gemma-3-4b-it-4bit and qwen2.5-7b-instruct-4bit: the route engages on all three, and disabling it changes neither the generated text nor throughput beyond noise.

Improvements

  • The MTP drafter is quantized at load: 810 MiB to 228 MiB (#1185). The drafter is read once per drafted token, so its cost is weight traffic. Measured on M5 Max against qwen3.8-27b-4bit: draft_block goes from 10.40 and 10.71 to 2.70 and 2.69 ms per round, the accept hook from 10.35 and 10.60 to 2.74 and 2.71, and the verify forward is unchanged. Throughput goes from 37.90 and 35.57 to 49.71 and 50.33 tok/s, about 1.19x to 1.5x against classic decode. Acceptance does not move (0.6831 to 0.6601 at 120 tokens, 0.6500 to 0.6589 at 300, disagreeing in sign). Output stays byte-identical to classic decode, structurally: the target verifies every proposal, so drafter numerics cannot reach the output. Applied at load rather than through a converted checkpoint, so existing bf16 drafters get it without being re-downloaded. MLXCEL_MTP_QUANTIZE_DRAFTER=0 opts out.
  • MTP engages on Apple GPU generation 15 and later instead of always declining (#1187). MLX sends M >= 2 affine quantized matmuls to a kernel that reduces along K differently from the one M == 1 takes, so a verify block was never byte-equal to the single-token chain there and the fail-closed gate declined every time. The gate now re-probes with that kernel disabled and keeps it off when that is what restores exactness. The switch costs 17 to 20 percent on the verify forward; measured over four alternating reps the result is 1.04x classic decode against 1.00x for declining. MLXCEL_QMV_WIDE=1 pins the faster kernel and declines MTP instead.

Bug Fixes

  • A VLM wrapper no longer re-enables a padded prefill its own text backbone disabled (#1201). supports_padded_prefill defaults to true and every hybrid family overrides it to false, because a tile-aligned prefill appends up to 31 pad positions and a recurrent state that has absorbed them cannot be rewound by trimming the KV caches. Qwen35VLModel and MiniCPMV46VLModel forwarded four other capability predicates and left this one defaulted. qwen3.8-27b-4bit is affected because its architectures field routes even a text-only run through the wrapper.
  • The server gates tile-aligned prefill on the model, not on the hardware alone (#1201). execute_full_prefill and both chunked-prefill paths padded whenever the host had a neural accelerator, without asking the model. Unlike the wrapper case this needed no wrapper, so it reached every hybrid family directly. Confirmed by A/B on one endpoint: the same /v1/completions request at temperature 0 returns different text before and after the guard.
  • The MTP exactness probe compares three inputs before reporting equality (#1186). One input is not enough to conclude two arms took the same kernels: mxfp4 group 32 at 5120 to 5120 moves 1 to 11 bytes of 10240 depending only on the operand draw, while the affine row at the same shape moves about 39%. The probe short-circuits on the first divergence, so the added cost lands only on the passing case at worker startup.
  • MTP is gated on a measured block-vs-chain probe rather than a static check (#1186). Whether a verify block is byte-identical to the equivalent single-token steps depends on which MLX kernel each quantized projection dispatches at M = K versus M = 1, which varies by GPU generation, quantization mode, operand size and block width and cannot be predicted from the checkpoint.
  • mlxcel video uses -fps_mode instead of the removed -vsync, and the tests no longer hide that (#1183).
  • A partial or unpinned Muse Glimmer checkpoint skips its tests instead of panicking (#1172, #1179).

CI/CD Improvements

  • A source-level test fails when a vision wrapper holds a backbone that refuses padded prefill and does not forward the predicate, and a second fails when an align_to_na_tile call site is not guarded by it. Both derive the refusing-backbone set by scanning src/models rather than hard-coding it, so a hybrid family added later is covered on the day it lands. Both were verified to fail by removing a guard.

Technical Details

  • The MLX pin moves to upstream 9a795735, 168 commits on from 2c46b953. Upstream touched four of the 21 in-tree overlay targets in that range, so four were three-way merged and the rest left alone. One conflict, in GatherMM::eval_gpu, where upstream widened an empty-input guard on the same lines the overlay had renamed two variables; the resolution takes upstream's added condition under the overlay's names. Each merged overlay's delta against the new base is the same size it had against the old one.
  • mlx/backend/metal/quantized.cpp is now overlaid to expose the qmv_wide off-switch the exactness gate needs. Upstream has no equivalent knob and the predicate is unchanged at this pin.
  • The three in-tree fused Metal kernel launchers re-validate at RMS 1.7263e-4 and 1.5259e-4 against their 5e-3 contract.

Dependencies

  • MLX C++ pinned at 9a795735 (2026-08-17).

Breaking Changes

None.

Known Issues

  • Seventeen numeric-agreement tests fail on M5 Max and three of them on CUDA, all under the same cause (#1065). Every one passes with MLX_ENABLE_TF32=0. The failures predate this release: the M5 set is identical at the previous MLX pin, verified over three isolated runs at each pin. The issue's original framing as an M5-specific defect is wrong, since the same switch explains both backends.
  • Whole-workspace test runs are not a stable comparison basis on either platform (#1092). On macOS the mlxcel-core binary takes an intermittent SIGSEGV under --workspace and is clean over five isolated runs; on CUDA one surgery target failed once and has passed on every re-run since, standalone and in-workspace. --test-threads=1 bounds threads within a binary and does not serialize across them, so it is not the lever here.
  • CUDA validation for this release is make verify-test-cuda at fa260a5d on NVIDIA hardware. It found no defect attributable to the MLX pin bump: the three failures are the TF32 class above, and the intermittent surgery target does not reproduce. The four rebased overlays are CUDA-only and this is the only gate they get, since the CUDA backend has no CI test coverage.

v0.5.1

Choose a tag to compare

@inureyes inureyes released this 15 Aug 01:56

Multi-turn prompt-cache reuse now works on the families that could not use it. A conversation's second turn on qwen3.5-0.8b-4bit went from 0 cached tokens to 194 of 227.

New Features

  • History-boundary snapshot during prefill (#1143). Snapshot-only families reuse a prompt cache only through an exact-prefix match, and the vector donated at end of generation is prompt + generated, which fails to prefix the next turn for three reasons that all live past the history boundary: generation-prompt-only scaffolds, <think> blocks dropped when a reply is re-rendered as history, and sampled tokens not being the canonical tokenization of their own text. A second snapshot is now taken during prefill, keyed by the add_generation_prompt = false render clipped to the longest common prefix it shares with the live prompt, which is a prefix of every follow-up turn by construction. MLXCEL_DISABLE_BOUNDARY_SNAPSHOT=1 restores the previous prefill.
  • Background warm-up of the next turn's history prefix (#1144). After a healthy completion the server renders the expected prefix and, when it has nothing else to do, restores the conversation's snapshot and prefills only the delta. Dispatched only from the idle tick with an empty active batch, an empty prefill queue, and no parked chunked prefill. Tool-calling turns are skipped. MLXCEL_DISABLE_CACHE_WARMUP=1 turns it off.
  • Session-chain supersede and operator control of the snapshot budget (#1146). A newer snapshot removes every stored snapshot that is a strict prefix of it within the same session and the same model/lora/template/multimodal bucket, so a conversation holds one entry instead of evicting its own earlier turns. --prompt-cache-snapshot-capacity-bytes, --prompt-cache-snapshot-max-entries, and --prompt-cache-snapshot-ttl are on both binaries with MLXCEL_* fallbacks, and /v1/cache/stats reports snapshot_supersedes apart from snapshot_evictions_lru.
  • Gemma 4 longest-prefix snapshot restore (#1145). A snapshot that shares a long prefix but diverges before its own end is adopted at that prefix rather than refused, following the upstream mlx-lm can_trim_prompt_cache precedent. Recurrent families keep exact-prefix semantics.
  • snapshot_diverged reject classification (#1147). A same-session candidate that is not a prefix of the request used to be indistinguishable from an empty store. /v1/cache/stats now carries the divergence geometry as last_reject_context_len and last_reject_entry_len, and /metrics gains a reason="snapshot_diverged" series.
  • -m/--model accepts --revision <REV> on generate, run, serve, inspect, and mlxcel-server (#1113), matching mlxcel download --revision. The HuggingFace cache probe is revision-aware; the legacy ./models/<name> directory and the mlxcel store are keyed without a revision component, so a revision-qualified request skips them rather than answering with an unknown revision.
  • mixed_4_8 checkpoints load on all 16 families sharing the fused QKV projection (#1090). A layer whose q/k/v planes cannot be concatenated keeps them separate in exactly the layout the checkpoint stored them in: nothing is dequantized, nothing is requantized, and the extra memory cost is zero.

Improvements

  • LocateAnything loads through the shared path (#1090). The per-family workaround added in #1070 dequantized 18 of the released 3B checkpoint's 36 layers' q/k/v planes to bf16, about 190 MB. Those planes now stay packed. Helium's pre-flight weight validator no longer rejects a mixed_4_8 attention block either.
  • A zero-block --kv-cache-budget names what consumed it (#1091). The paged decode v2 workspace reserve is charged to the requested budget first and is device-derived, about 16 MiB for the common 8-kv-head, 128-head-dim geometry. The warning blamed model size and available memory; it now reports the requested budget, the reserve, and the smallest budget that would mint one block. No budget resolves to a different block count than before.
  • One SSE keepalive constant and one constructor for all five streaming routes (#1105, #1107). The 15-second interval was declared three times and the compile-time assertion that gives it meaning guarded one of them, so /v1/responses and the Anthropic-compatible surface could be raised past a reverse-proxy idle timeout without the build failing. Routes now build through sse_response, which takes the keepalive by value, so a sixth route cannot forget it. Behavior is unchanged; every value was already 15.
  • Documentation: the six paged-decode v2 environment variables and a corrected MLXCEL_PAGED_ATTENTION_NATIVE row (#1104), why CLI DRY runs without sequence breakers (#1108), the three tracked shared functions in utils.rs (#1110), and stale cross-references left behind by earlier renames (#1106).

Bug Fixes

  • mlxcel detect returned boxes that matched no page content (#1089). RT-DETRv2's pred_logits and pred_boxes come out bf16 for the shipped checkpoints, because the first conv against a bf16 weight settles the graph back into bf16 after pixel_values is cast to f32, and the host readback parsed the raw buffer as 4-byte f32. That fused adjacent bf16 pairs into bogus floats and returned half the elements, and the decode then indexed into buffers half the advertised size, misaligning every query, label, and box without erroring. The coordinate mapping the issue suspected was already correct, and the aspect-ratio dependence was a coincidence of which garbage indices survived.
  • Server token streams at temperature <= 0 with DRY and sequence breakers (#1102). The greedy branch of build_sampling_config dropped a per-request dry_sequence_breakers and replaced it with an empty vector, so the backward match ran past the intended boundary and the penalty came out at or above what the request asked for. Output for those requests changes, toward the requested configuration.
  • --dry-sequence-breaker reaches the sampler instead of being parsed and discarded (#1103). Breaker strings are tokenized at startup and used as the default for any request that does not send its own, /props reports the resolved token IDs, and a breaker that does not encode to exactly one token fails startup with a message naming it. The escapes \n, \t, \r, and \\ are interpreted, and the tokenizer's prepend normalizer is discounted so a breaker resolves to the intended token on SentencePiece-derived checkpoints.
  • Four server flag spellings that worked on only one of the two binaries now work on both (#1109). mlxcel serve accepts --parallel and --predict; mlxcel-server accepts --n-parallel and --adapter. --dry-sequence-breaker is the primary name on both, with the plural kept as an alias.
  • The "not a model" error names all three forms -m/--model accepts (#1114), including the bare name resolved against $MLXCEL_DEFAULT_ORG, and states the character class that sends a bare name down that path.
  • The 2D parallelism guard named --tensor-parallel-size, a spelling no binary accepts; it now names --tp-size (#1112). No user-visible change: the guard's condition is the exact negation of an early return above it, so the message has never been emitted. The dead condition is filed as follow-up.

CI/CD Improvements

  • Homebrew tools are on PATH for the macOS release jobs, which is what the v0.5.0 release run needed.
  • The docs-* Makefile targets are guarded and the manual documentation split is written down (#1111).
  • Two load-sensitive test failures fixed: the pinned muse_glimmer checkpoint tests are skipped when the checkpoint is absent and the MLXCEL_BACKEND reader race is locked (#1155), and prompt_cache_e2e derives its slack from the paged block size instead of a hard-coded constant (#1156).

Technical Details

  • Measured on Apple M1 Ultra, 2026-08-14, through the production /v1/chat/completions path with counters read from /v1/cache/stats. Arm attribution is by counter in every table: snapshot_inserts advances once per turn without the boundary snapshot and twice with it, and snapshot_warmups_run can only move when a warm-up actually restored, prefilled, and stored.
  • Boundary snapshot, qwen3.5-0.8b-4bit, identical three-turn conversation: turn 2 cached tokens 0 to 150 of 189, turn 3 0 to 184 of 214. A llama-3.2-1b-instruct dense-KV control is identical in every cell across arms.
  • Warm-up, same fixture: turn 2 cached 150 to 194 of 227, uncached tokens down 57%. Under 12 requests over 2.6s of continuous load with 2 clients at --parallel 4, 0 warm-ups ran and 10 were skipped; the 2 runs landed after the load stopped.
  • Gemma 4 partial restore, gemma-4-e2b-it-4bit with a 512-token sliding window: a diverging turn adopted 66 tokens of a 168-token stored entry. The qwen3.5-0.8b-4bit negative control logged no partial restore and kept snapshot_hits at 0, so the recurrent path is provably unchanged.
  • Total tokens forwarded by a boundary-snapshot prefill are unchanged; the cost is one extra graph launch plus a model-state copy. The segment forward evaluates only its final logits position.
  • Snapshot cost tracks model width, not prompt length: on qwen3.5-0.8b-4bit one conversation snapshot is 13246464 bytes at 72, 98, 124, and 151 tokens alike. Size the budget by measuring snapshot_bytes and multiplying by the conversations expected resident.
  • Epic-level verification on the large checkpoints: gemma-4-31b-it-4bit turn 2 reached 147 of 171 cached tokens and qwen3.5-4b 94 of 124, with a dense control showing no regression.

Dependencies

None. The only lockfile change is the four workspace crates moving to 0.5.1.

Breaking Changes

None to any API or flag. Two output changes are worth knowing about before upgrading:

  • Requests that set dry_multiplier and sequence breakers generate diffe...
Read more

v0.5.0

Choose a tag to compare

@inureyes inureyes released this 12 Aug 09:55

All changes since the last stable release, v0.4.3. The v0.5.0-beta.1 entries are consolidated here.

New Features

Meta Muse Glimmer

  • Meta Muse Glimmer (muse_glimmer) vision-language support (#1101). The 52-layer mixed-cache decoder and the 50-layer vision and fusion path drive single-image and multi-image generation through both the CLI and the continuous-batching server, with expanded prompt accounting and fail-closed unsupported modes. The pinned Muse chat template, reasoning strengths, and bounded ATEM parsing and replay are wired across the Chat Completions, Responses, and Anthropic-compatible routes, including streaming cleanup, so reasoning is routed to reasoning_content rather than leaking into the answer.
  • mlx-community/Muse-Glimmer-30B-4bit loads through the same path (#1116), preserving the canonical bf16 path and the dense vision tower. Warm text qualification on NVIDIA GB10 reached 12.43 prefill and 13.34 decode tok/s, about 3.1x the recorded 4.25 tok/s bf16 baseline on the same box; the image path inserted 64 vision tokens and decoded at 13.15 tok/s. Checkpoint-format support is fail-closed to the pinned affine-Q4 layout, with conflicting quantization contracts, invalid group-size and bit combinations, orphan sidecars, missing affine biases, and quantized vision-tower tensors all rejected before kernel selection.

Florence-2

  • Florence-2 (florence2) encoder-decoder vision-language support, across the five sub-issues of epic #850: the BART seq2seq engine and text core (#1060), the DaViT vision backbone (#1063), vision-language fusion and full weight loading (#1064), the processor with the fifteen task markers and <loc_N> coordinate parsing (#1069), and CLI integration with real-checkpoint validation (#1071).
  • Quantized Florence-2 checkpoints at 3, 4, 6, and 8 bits, both base-ft and large-ft (#1082).
  • mlxcel-server serves Florence-2 through a dedicated single-stream seq2seq worker (#1083). message.content is byte-identical to the CLI answer, and parsed coordinates arrive as JSON in message.florence2_result.

Other new model families

  • LocateAnything (locateanything) grounding VLM (#1070): a MoonViT tower reused from Kimi-VL, an MLP connector, and a Qwen2 decoder emitting <ref> / <box> markers and the 1001 coordinate tokens.
  • Falcon-OCR (falcon_ocr) early-fusion document OCR (#1075), the first family in the tree with no vision tower: 16x16 RGB patches go through one linear projector straight into the token stream. mlxcel generate --layout-detections <FILE> runs the layout-aware second stage.
  • Jina VLM (jvlm) (#1076): a 27-layer SigLIP-so400m-class tower feeding a Molmo-style connector into a Qwen2-class decoder on an OLMo tensor layout.
  • Eight text families: Ling / Bailing MoE (#838), OpenELM (#839), TeleChat3 (#841), DBRX (#835), Ling / Ring linear-attention MoE (#840), Phixtral (#844), Arcee AFMoE (#845), and Klear (#846).

Decode and sampling

  • Fused paged-attention decode v2, with a CSR page table, a cross-CTA split-KV partial kernel, and a variable-length merge kernel (#898), routed for batched decode in the server behind measured floors (#899).
  • Fused sparse-attention decode by page indirection over the v2 kernel (#904). MiniMax-M3 block-sparse decode measures 2.06x at 32x sparsity.
  • Cascade (shared-prefix) decode (#903), default off.
  • Fused residual-add RMSNorm and fused RoPE + KV-append decode kernels (#905), both default on.
  • Sorting-free top-p sampling by dual-pivot rejection kernels (#901), 1.28x to 2.35x on M1 Ultra, and a softmax-free Gumbel-max categorical sampling kernel (#900).
  • Shape-bucketed kernel autotuner and a cold-cache microbenchmark harness (#906); matrix-absorbed MLA decode over a compressed-latent KV cache (#907), default off.
  • Opt-in acceptance-optimal rejection sampling for the classic speculative path (#902), default off.
  • MLXCEL_EXTRA_CA_CERTS accepts a PEM bundle so checkpoints can be fetched from behind a TLS-inspecting proxy (#912).

Improvements

  • Chunked GLA prefill is the bailing_moe_linear default (#1062). On mlx-community/Ring-mini-linear-2.0-4bit and an M1 Ultra, chunked wins perplexity in every configuration measured, by 1.50% at a 128-token window rising to 13.75% at 512 across 32 windows, and prefill is 2.06x to 2.39x faster from about 512 to about 8192 prompt tokens. Decode is unchanged. MLXCEL_BAILING_LINEAR_CHUNKED_PREFILL=0 restores the sequential recurrence.
  • Loop detection is less trigger-happy (#967). At the old threshold a markdown table alignment row tripped the detector at its fourth column, so a correct answer stopped mid-table and the client saw a normal stop with no error.
  • The paged block-pool slab size is a per-pool value sized from the reserved KV memory instead of a constant 32 (#899).
  • Per-expert tensors are stacked from borrows instead of copies, removing 5376 Copy graph nodes for ling-lite-1.5 (#948).
  • A parked chunked prefill's wait is bounded by a fairness grant, cutting worst-case admission latency from 50.5s to 23.8s (#1011).
  • Fused decode-MoE byte-identity is documented as checkpoint-dependent rather than as a general property (#1061). Over 96 calls per checkpoint the fused kernel sits at 1.65e-3 normalized RMS from an all-f32 ground truth against the fallback's 1.02e-2, closer to truth in 96 of 96 calls on two families.
  • The modelsize / modeltype / arch label taxonomy is documented (#1038).

Bug Fixes

  • Four model families ran a fully bidirectional prefill. deepseek_v2 (#991), then internlm3, hunyuan, and gemma2 (#999) passed mask == None through to attention instead of building a causal mask, so every prompt token attended to every later one. Output stayed fluent, which is why it survived.
  • CUDA JIT launches read their buffers through the wrong pointer type when one process ran a kernel at two input dtypes. MLX's CUDA backend omits input dtypes from the JIT cache key while generating buffer parameter types from the runtime dtype, so the first dtype to compile won for the life of the process and later calls returned numbers unrelated to their inputs, with nothing thrown. Fixed for the paged decode v2 kernels (#1058) and the rest of the class, including both samplers (#1059). Metal was never affected.
  • mlx::core::dequantize returned wrong values whenever its biases input was not row-contiguous; the FFI shim now forces contiguity (#1078). The window opened with the MLX pin bump in #1046 and was two days wide.
  • Speculative decoding lost about half its throughput to a draft KV cache rewind arithmetic error (#994). Run-to-run spread goes from 3.09x to 1.06x.
  • Two Molmo v1 seams that corrupt image-conditioned generation (#1099): an omitted rope_impl now defaults to the checkpoint's effective LLaMA rotate-half layout, and image preprocessing pads raw black before normalization and preserves fractional mask coverage.
  • Queue admission, image cardinality, and usage accounting on the dedicated single-stream workers (#1098). DiffusionGemma, LLaDA-2 MoE, and Florence-2 now reserve before streaming SSE opens, so a request losing the final race returns HTTP 503 at the route boundary instead of mid-stream.
  • The Gemma 4 server reasoning splitter consumes the full <|channel>thought opener, so reasoning_content no longer starts with a leaked channel argument (#1094).
  • Quantized Jamba MoE experts are built through the shared SwitchGLU (#974); deepseek::ModelArgs no longer drops the nested quantization block (#975); QuantizedMultiLinear infers its mode and validates the bias plane (#1028); the declared quantization.mode is bounded by an allowlist at load (#973); quantization parameters are bounded in the family-local MoE expert loaders (#958, #929).
  • The MLA kv_b_proj sanitizers refuse a scales-without-biases plane (#1026), and mamba / mamba2 treat a scales-without-biases embedding as non-quantized (#976). GLM4 MoE Lite loads its own kv_b_proj layout (#1029).
  • rope_traditional is honored on the shared Llama path (#931); topk_group is bounds-checked (#947); KV-cache estimation recognizes the n_layer / n_embd and multi_query spellings (#927); LoRA adapter weight fusion is Conv1D-layout aware (#925).
  • Grammar-only requests are excluded from loop detection (#977) and generation stops when the structured matcher completes (#978). Disaggregated router chat guards are enforced (#979) and client token budgets are clamped (#980).
  • CUDA: the broken RMSNorm overlay is dropped and DeepSeek-V2 graph capture is restored (#831); the fused RoPE kernel calls libdevice cosf / sinf (#1049).
  • XLA: the terminating EOS token is no longer emitted as output (#963), the Qwen2-VL prefill export is gated on the feature that calls it (#961), and the loader contract is pinned (#966).

CI/CD Improvements

  • The quality gate covers all workspace members (#1007). The previous bare cargo test and cargo clippy resolved to -p mlxcel and never built mlxcel-core, mlxcel-surgery, or mlxcel-xla, leaving 1754 tests and their whole lint surface ungated.
  • Tests build under a new [profile.test-fast] rather than [profile.release] (#1000), which was making the nightly spend its entire 180-minute budget in codegen and linking without starting a test.
  • make bump-version VERSION=x.y.z rewrites every version-tracking manifest and syncs Cargo.lock in one step, and make verify-versions derives the member list from [workspace] members at runtime so it cannot go stale (#1056).
  • make verify-kernel-dtype-keys and a CI job require every template_args initialiser in a file containing a cuda_kernel( call to name at least one input dtype (#1059). There is no allowlist.
  • The CUDA merge gate is runnable and serialized (#1048). The rust-toolchain pin is restored and excluded from dependabot bumps (#952).
  • Two flaky suites are made ...
Read more

v0.5.0-beta.1

Choose a tag to compare

@inureyes inureyes released this 06 Aug 12:56

First beta for v0.5.0. 70 commits since v0.4.3.

The bulk of it is the decode-kernel epic (#909): paged-attention decode v2 becomes the production path, sparse and shared-prefix decode reduce onto it, and sampling stops sorting. Eight text model families are added, four families are fixed for having run a bidirectional prefill, and the quality gate stops skipping three quarters of the workspace.

New Features

  • Fused paged-attention decode v2 with a CSR page table, a cross-CTA split-KV partial kernel, and a variable-length merge kernel (#898). v1 splits the KV inside one threadgroup, so one CTA serves one (batch, query head) pair and adding context adds no parallelism. v2 splits across CTAs: parallelism is num_chunks * kv_heads * q_groups and grows with context. The CUDA bodies are structural transliterations of the Metal ones and have never been compiled or run.
  • Batched paged decode routes through v2 in the server (#899), gated by measured floors rather than switched unconditionally. On M1 Ultra: 1.41x at batch 4 and 1.47x at batch 8 with 1024 tokens per request, 1.08x median for a single request at 4096 tokens, and 0.91x for a single request at 1024, which is why the single-request floor is 4096 visible tokens and the batched floor is 512 per request. MLXCEL_PAGED_V2_MIN_KV_TOKENS and MLXCEL_PAGED_V2_MIN_KV_TOKENS_PER_REQUEST move the floors; MLXCEL_PAGED_ATTENTION_NATIVE=0 pins the previous gather path.
  • Fused sparse-attention decode by page indirection over the same kernel (#904). A sparse selection becomes a page_size = 1 page table, so no gathered copy is materialized at any point. MiniMax-M3 block-sparse decode is routed onto it above a sparsity gate (default 8x, MLXCEL_SPARSE_PAGED_MIN_SPARSITY): 0.67x at 2x sparsity and 0.77x at 4x, then 1.22x at 8x, 1.17x at 16x, and 2.06x at 32x. DeepSeek Sparse Attention is not routed onto it; docs/sparse-paged-decode.md records why and what kernel generalization it needs.
  • Cascade (shared-prefix) decode (#903). A whole-page prompt prefix shared by several sequences in a decode batch is attended once for the subgroup and merged into each member's suffix state. Default off (MLXCEL_CASCADE_ATTENTION): correctness is verified, throughput is not.
  • Sorting-free top-p sampling by dual-pivot rejection kernels on Metal and CUDA (#901). One softmax and one custom kernel replace the argsort + cumsum nucleus filter and the trailing random::categorical, with no sort anywhere. 1.28x to 2.35x on M1 Ultra across vocab {32K, 64K, 152K} and batch {1, 4, 8}. Routing is restricted to what measured faster: top-k alone, min-p alone and top-k with min-p keep the stock chain, and top-k with top-p routes only up to vocab 32768.
  • Softmax-free Gumbel-max categorical sampling on the no-filter stochastic path (#900), which skips the normalization pass over the 32K to 152K entry vocabulary.
  • Eight text model families: Ant Group Ling / Bailing MoE (#838), Apple OpenELM (#839), TeleAI TeleChat3 (#841), Databricks DBRX (#835), Ling / Ring linear-attention MoE (#840), Phixtral (#844), Arcee AFMoE (#845), Kuaishou Klear (#846).
  • Matrix-absorbed MLA decode over a compressed-latent KV cache (mlxcel_core::mla), wired to deepseek_v2 behind MLXCEL_MLA_ABSORBED (default off) (#907).
  • Shape-bucketed kernel autotuner (mlxcel_core::autotune) and a cold-last-level-cache microbenchmark harness (#906). The v1 paged-decode NumSplits launch shape is the first consumer, validated on Apple Silicon; the two CUDA knobs are wired but unvalidated.
  • Fused residual-add RMSNorm and fused RoPE + KV-append decode kernels, both default on, with MLXCEL_FUSED_ADD_RMSNORM=0 and MLXCEL_FUSED_ROPE_APPEND=0 as kill switches (#905).
  • Opt-in acceptance-optimal rejection sampling for the classic speculative path, with the proof, regression guards, and a closed-form acceptance diagnostic (MLXCEL_SPECULATIVE_STOCHASTIC_ACCEPT=1, default off) (#902).
  • MLXCEL_EXTRA_CA_CERTS accepts a path to a PEM bundle added to the default TLS trust roots, so checkpoints can be fetched from behind a TLS-inspecting corporate proxy (#912).

Improvements

  • Loop detection is less trigger-happy: LOOP_DETECTION_RECOMMENDED moves from min_count = 4 to 12, and the Gemma 4 family default-on is narrowed (#967). At the old threshold a markdown table alignment row tripped the detector at its fourth column, so a correct answer stopped mid-table and the client saw a normal stop with no error.
  • The paged block-pool slab size is a per-pool value instead of a constant 32, sized by the server from the configuration it already reserved KV memory for. MLXCEL_PAGED_SLAB_BLOCKS=0 pins the old default (#899).
  • Per-expert tensors are stacked from borrows instead of copies, removing one Copy graph node per per-expert tensor: 5376 of them for ling-lite-1.5 (#948).
  • A parked chunked prefill's wait is bounded by a fairness grant, cutting worst-case admission latency from 50.5s to 23.8s at the default --prefill-grant-interval of 16, at 1.60x mean stream ITL during admission (#1011).
  • MLXCEL_MIXED_STEP prototype and ADR 0005, which records why model-level ragged mixing and kernel-level mixing on MLX were rejected (#908).
  • The modelsize / modeltype / arch label taxonomy is documented (#1038).

Bug Fixes

  • Four model families ran a fully bidirectional prefill. Generation calls forward with mask == None and expects the model to build its own causal mask; deepseek_v2 (#991), then internlm3, hunyuan and gemma2 (#999) passed None through to attention, so every prompt token attended to every later one. Output stayed fluent, which is why it survived. gemma2 was the mitigated case (its sliding window bounded the leak); PaliGemma 2 was checked and is unaffected.
  • Speculative decoding lost about half its throughput to a draft KV cache rewind arithmetic error (#994). Aggregate decode goes from 41.8 tok/s with a 3.09x run-to-run spread to a 1.06x spread.
  • Quantized Jamba MoE was entirely broken, dense checkpoints included, not just quantized ones. Experts now build through the shared SwitchGLU, and the router moves to UnifiedLinear so a checkpoint shipping 8-bit routers under a 4-bit default loads (#974).
  • Quantized checkpoints fail at load instead of aborting at the first forward. The declared quantization.mode is bounded by an allowlist (#973), MoE expert loaders bound their parameters (#929, #958), QuantizedMultiLinear infers its own mode and validates the bias plane (#1028), and the MLA and Mamba sanitizers refuse a scales-without-biases plane (#976, #1026). deepseek::ModelArgs had been dropping the nested quantization block outright, so the expert loaders never saw the declared values (#975).
  • GLM4 MoE Lite loads its own layout: kv_b_proj is decomposed into the per-head embed_q / unembed_out pair (#1029).
  • rope_traditional from config.json is honored on the shared Llama path (#931).
  • topk_group is bounds-checked inside group_mask_scores (#947).
  • KV-cache estimation recognizes the n_layer / n_embd and multi_query config spellings (#927).
  • LoRA adapter weight fusion is Conv1D-layout aware (#925).
  • Grammar-only requests are excluded from loop detection (#977), and generation stops when the structured matcher completes (#978).
  • Disaggregated router chat guards are enforced (#979), and client token budgets are clamped (#980).
  • CUDA: the broken RMSNorm overlay is dropped and DeepSeek-V2 graph capture is restored (#831); the fused RoPE kernel calls libdevice cosf / sinf instead of the inaccurate variants (#1049).
  • XLA: the terminating EOS token is no longer emitted as output (#963), and the Qwen2-VL loader contract is pinned with all three outcomes documented (#966).

CI/CD Improvements

  • The quality gate covered one workspace member out of four. The workspace root is itself the mlxcel package, so a bare cargo test or cargo clippy resolved to -p mlxcel and never built mlxcel-core, mlxcel-surgery or mlxcel-xla: 1754 tests, 1354 of them mlxcel-core's, plus those members' entire test-target lint surface. make verify-test and make verify-clippy now pass --workspace (#1007).
  • Tests build under a new [profile.test-fast] (#1000). It inherits from release and keeps opt-level = 3, so the optimized MLX numerics the suite depends on are unchanged, but drops fat LTO and codegen-units = 1, which had the nightly spending its entire 180-minute budget in codegen and linking without ever starting a test. An unfinished nightly run is now also reported instead of skipping the red-main step because a timeout records as cancelled. release.yml still builds and links under [profile.release].
  • A second mlxcel-core test binary sharing the GPU fails loudly instead of aborting the run (#1008).
  • The fused MoE GeGLU parity bound is derived from the reference's measured jitter (#964); the mllama ragged test tolerates f32 reassociation and gained a nightly gate (#939); the autotune rep-scaling test no longer pins an exact ceiling (PR #990); a dead-callsite race behind the drafter log assertions is removed (#1023); the test harness honors CARGO_TARGET_DIR (PR #996).
  • The CUDA merge gate is runnable and serialized (#1048).
  • The rust-toolchain pin is restored and excluded from dependabot bumps (#952).
  • Homebrew tap formula stanzas are edited by name instead of line position (PR #949).

Technical Details

  • The v2 kernel resolves a visible token through indices[indptr[r] + (fpo + i) / page_size] and assumes nothing about the rows being adjacent, ordered or complete. That is why a sparse selection needs no kernel change: it is a shorter page list, and at page_size = 1 the page indirection becomes a per-token gather executed inside the attention loop.
  • Both #903 and #904 consume #898's `p...
Read more

v0.4.3

Choose a tag to compare

@inureyes inureyes released this 27 Jul 04:54

The headline fix: 4-bit decode on CUDA is now deterministic. Every 4-bit model produced different output run to run at temperature 0, and occasionally a garbage token, because the qmm_sm80 quantized-matmul kernel reused shared memory still being written by in-flight async copies. This release also fixes three Gemma 4 output-correctness bugs, adds four text model families, and expands the opt-in OpenXLA multimodal paths.

New Features

  • GPT-2 (#924), GPT-BigCode (#926), GPT-NeoX (#928), and Kyutai Helium (#930) text models.
  • OpenXLA backend (opt-in, xla-iree) multimodal expansion: Gemma3n text runtime and dense PLE prefill (#892) and audio (#883); token-exact Phi4MM audio (#887) with per-slot adapters (#914); LLaVA vision through IREE (#913); the Qwen2-VL vision path (#915); sparse DeepStack prefill (#893); multimodal RoPE position state (#894); image requests admitted in the CLI and continuous-batch serving (#895); bounded audio preprocessing and serving plumbing (#896); a prefill embeddings entry point (#879) with IREE sessions seeded from prepared embeddings (#888).
  • --show-reasoning CLI flag to print the reasoning channel that is otherwise hidden (#889).

Improvements

  • OpenXLA operator numeric contracts and a bounded numeric-oracle probe harness covering dense matmul, affine Q4 dequant, prefix scan, and core ops (#934, #937, #938, #940, #941, #942, #943, #944).
  • Parameterized static context capacity on OpenXLA (#880).
  • VLM host prefill preprocessing refactored to owned buffers (#881).
  • Intra-MoE phase profiling for nemotron-h (#832).

Bug Fixes

  • 4-bit CUDA decode was non-deterministic at temperature 0 and could emit a stray token. The qmm_sm80 kernel reused shared memory still being written by in-flight cp.async copies; it now drains outstanding copies before the epilogue store, so greedy decode is byte-identical run to run on every 4-bit model (#910).
  • The fused decode-MoE kernel rounded each of the K per-expert partials to bf16 before summing, which corrupted long multi-turn Gemma 4 output. Partials are now accumulated in f32 and rounded once (#886).
  • Gemma 4 chunked-prefill continuation dropped the sliding-window attention mask when the caller mask was shorter than the returned keys, collapsing output to reserved <unused> tokens under concurrent load. The mask is now sized to the keys each attention family returns (#891).
  • The Gemma 4 reasoning channel and its chain-of-thought no longer leak into CLI output as raw tokens (#889).

CI/CD Improvements

  • OpenXLA diagnostic local-task thread count is bounded to keep the diagnostics harness within resource limits (#945).

Technical Details

  • The IREE compiler and runtime are pinned to 3.12.0rc20260721 on both CUDA and macOS. Wheels are fetched from the official GitHub release and verified by sha256, the CUDA setup script selects the wheel per Python ABI and host architecture, and the toolchain cache directory is versioned so a pin bump cannot reuse a stale build (#882).
  • License update for the FlashInfer project reference (562aadc).

Dependencies

  • Minor and patch version bumps across 11 packages (#828).

Breaking Changes

  • The Gemma 4 reasoning channel (<|channel>thought ... <channel|>) is hidden from CLI output by default. Scripts that read the raw reasoning from mlxcel run / mlxcel generate stdout will no longer see it; pass --show-reasoning to restore the previous output. Server reasoning_content behavior is unchanged (#889).

Known Issues

  • The standard 4-bit gemma-4-26b-a4b-it-4bit checkpoint deterministically emits an occasional stray token in some long contexts. This is inherent standard-4-bit (RTN) weight quantization of that checkpoint, not a runtime defect; the quantization-aware gemma-4-26b-a4b-it-qat-4bit checkpoint is unaffected (#933).
  • For Gemma 4 channel-format models, server reasoning_content retains a leading thought prefix from the channel marker (#890).

What's Changed

  • deps(deps): bump the minor-and-patch group across 1 directory with 11 updates by @dependabot[bot] in #828
  • update(nemotron-h): add intra-MoE phase profiling (#284) by @inureyes in #832
  • feat(xla): parameterize static context capacity by @inureyes in #880
  • feat(xla): add prefill embeddings entry (#858) by @inureyes in #879
  • refactor(vlm): add owned host prefill preprocessing (#859) by @inureyes in #881
  • build: pin IREE 3.12.0rc20260721 and harden CUDA toolchain setup by @inureyes in #882
  • feat: add token-exact Phi4MM audio reference path (#874) by @inureyes in #887
  • feat(vlm): add Gemma3n audio reference path (#875) by @inureyes in #883
  • feat(xla): seed IREE sessions from prepared embeddings (#860) by @inureyes in #888
  • fix: hide Gemma 4 reasoning channel from CLI output by @inureyes in #889
  • feat(xla): add Gemma3n text runtime and dense PLE prefill by @inureyes in #892
  • feat(xla): add multimodal RoPE position state by @inureyes in #894
  • feat(xla): admit image requests in CLI and continuous-batch serving by @inureyes in #895
  • feat(xla): add sparse DeepStack prefill by @inureyes in #893
  • test(xla): validate a LLaVA reference architecture end to end by @inureyes in #897
  • fix(gemma4): size chunked-prefill masks to returned keys per family by @inureyes in #891
  • feat(xla): add bounded audio preprocessing and serving plumbing by @inureyes in #896
  • fix(moe): keep fused decode-MoE partials in f32 until the final K-sum (#886) by @inureyes in #911
  • feat(xla): execute LLaVA vision through IREE by @inureyes in #913
  • feat(xla): support Phi4MM audio and per-slot adapters by @inureyes in #914
  • feat: add Qwen2-VL OpenXLA vision path (#865) by @inureyes in #915
  • fix(cuda): drain cp.async before qmm_sm80 epilogue smem store (#910) by @inureyes in #923
  • feat(models): add GPT-2 (gpt2) text model support by @inureyes in #924
  • feat(models): add GPT-BigCode (gpt_bigcode) text model support by @inureyes in #926
  • feat(models): add GPT-NeoX (gpt_neox) text model support by @inureyes in #928
  • feat(xla): version auxiliary dtype contracts by @inureyes in #934
  • feat(xla): bind operator numeric contracts by @inureyes in #937
  • feat(xla): add bounded numeric oracle harness by @inureyes in #938
  • feat(xla): run dense matmul numeric probe by @inureyes in #940
  • feat(xla): expand core numeric probes by @inureyes in #941
  • feat(xla): share contract-sensitive numeric emitter helpers by @inureyes in #942
  • feat(xla): add explicit CUDA prefix scan contract by @inureyes in #943
  • feat(xla): add affine Q4 dequant micro-oracle by @inureyes in #944
  • feat(models): add Kyutai Helium (helium) text model support by @inureyes in #930
  • fix(xla): bound diagnostic local-task threads by @inureyes in #945

Full Changelog: v0.4.2...v0.4.3

v0.4.2

Choose a tag to compare

@inureyes inureyes released this 20 Jul 05:31

New Features

  • MiniMax-M3 text model: a hybrid dense/MoE architecture with block-sparse attention (#799).
  • MiniMax-M3-VL multimodal support (#800).
  • Unlimited-OCR: long-document OCR on the DeepSeek-OCR stack with a per-layer ring sliding decode cache that keeps the full prefill and rotates only the most recent decode window, so decode does not grow the cache unbounded (#801).
  • XTC (Exclude Top Choices) sampling, read and applied end to end on the OpenAI-compatible chat, completions, and responses routes, with xtc_probability / xtc_threshold range validation (#802).
  • thinking_mode chat-template kwarg is injected when the template references the identifier and thinking is enabled (#811).
  • Per-reason prompt-cache reject metrics and an APC trace log (#810).

Improvements

  • The speculative slice slot rotates across waiting requests instead of always going to the window head (#816).
  • The effective-input text check is allocation-free (#815).
  • Decode benchmark sweeps get a monotonic OOM backstop and a persistent OOM record (#808).

Bug Fixes

  • DeepSeek-V2-Lite generated repeated tokens on GB10 CUDA. An upstream MLX 0.32.1 RMSNorm kernel regression is overlaid with its last-good version, and CUDA graph capture is disabled for the DeepSeek-V2 family (as it already is for Gemma 4), so output is coherent again (#829).
  • Gemma 4 audio: the CLI rendered the <|audio|> placeholder before the prompt text, which flipped the 12B unified model from transcription into answering the perceived content on acoustically hard clips. The placeholder now follows the prompt text (#798).
  • Requests with no effective input (empty or whitespace-only prompts and messages) were dispatched to the model; they now return a 400 before dispatch on /v1/chat/completions, /v1/completions, and /v1/messages (#803, #813, #814).
  • Long-lived speculative serving on CUDA hit a fatal MLX "Cache thrashing" abort once the graph cache filled. CUDA builds raise the MLX_CUDA_GRAPH_CACHE_SIZE default so the abort no longer fires (#818).
  • An MLX evaluation throw in the batch decode loop aborted the whole worker and dropped every in-flight request; it now fails the affected request instead (#825).
  • An empty paged-state length list from a not-yet-populated sequence returned an Err that spammed a "Failed to sync paged state" warning; it is treated as a benign no-op (#826).

CI/CD Improvements

  • Added a fast test profile to cut edit-test iteration time (#812).

Technical Details

  • The DeepSeek-V2 regression was pinned to upstream MLX #3792 ("Fix CUDA RMSNorm small-row dispatch") by a deterministic bisect run under MLX_USE_CUDA_GRAPHS=0: DeepSeek-V2's kv_a_layernorm (axis 512, N_READS 8) lands on the affected small-axis dispatch config. The MLX pin is kept at the latest commit; only mlx/backend/cuda/rms_norm.cu is overlaid back to its pre-#3792 version. Reporting the regression upstream and dropping the overlay once fixed is tracked in #830.
  • DeepSeek-V2 also shares the Gemma 4 CUDA graph-capture decode hazard (#688), so it is added to the MLX_USE_CUDA_GRAPHS=0 startup lever. Re-enabling capture for both families is tracked in #831.

Dependencies

None.

Breaking Changes

None.

Known Issues

  • DeepSeek-V2 runs with CUDA graph capture disabled on GB10 (as Gemma 4 does), which costs roughly 8% decode throughput. Output is correct; the throughput is recovered once the upstream graph-capture hazard is fixed (#831). Set MLX_USE_CUDA_GRAPHS=1 to override.

What's Changed

  • fix(cli): render Gemma 4 audio user turn after the prompt text by @inureyes in #798
  • feat(models): add MiniMax-M3 text model (hybrid dense/MoE with block-sparse attention) by @inureyes in #799
  • feat(vlm): add MiniMax-M3-VL multimodal support by @inureyes in #800
  • feat(vlm): add Unlimited-OCR support with ring sliding decode cache by @inureyes in #801
  • feat(sampling): implement XTC sampling end to end by @inureyes in #802
  • fix(server): reject requests with no effective input before model dispatch by @inureyes in #803
  • feat(bench): monotonic OOM backstop and persistent OOM record for decode benchmark sweeps by @inureyes in #808
  • feat(server): per-reason prompt-cache reject metrics and APC trace log by @inureyes in #810
  • build: add a fast test profile to cut edit-test iteration by @inureyes in #812
  • feat(server): derive thinking_mode chat-template kwarg from template by @inureyes in #811
  • fix(server): reject /v1/messages requests with no effective input by @inureyes in #813
  • fix(server): reject whitespace-only prompts on /v1/completions by @inureyes in #814
  • perf(server): allocation-free effective-input text check by @inureyes in #815
  • perf(server): rotate the speculative slice slot across waiting requests by @inureyes in #816
  • fix(server): raise the CUDA graph-cache default to survive long-lived shape-diverse decode (#818) by @inureyes in #820
  • fix(server): survive MLX decode-eval throws instead of aborting the worker by @inureyes in #825
  • fix(cache): treat empty paged length list as a benign no-op by @inureyes in #826
  • fix(cuda): restore DeepSeek-V2 output by reverting the regressed RMSNorm kernel by @inureyes in #829

Full Changelog: v0.4.1...v0.4.2