The real review inputs used in issue #142 are publicly downloadable from the R2 bucket mlx-node-benchmarks. Each manifest pins an HTTPS URL, SHA-256 hashes, model identity, and the measured protocol. No Cloudflare account, credentials, or Wrangler installation is needed:
oxnode scripts/benchmark-fixture.ts fetch --fixture gemma4
oxnode scripts/benchmark-fixture.ts fetch --fixture qwen38
oxnode scripts/benchmark-fixture.ts verify --fixture qwen38 --tokenizer /path/to/prepared-model/tokenizer.json| Fixture | Measured input tokens | Recorded protocol |
|---|---|---|
| Gemma 4 QAT Q4_0 | 7,733 / 40,528 / 66,904 | 256 output tokens; BF16 KV in both runtimes; 512-token prefill chunks; fresh process per sample. |
| Qwen3.8 UD-Q4_K_XL | 6,219 / 18,754 / 32,488 | 512 output tokens; BF16 MLX / F16 llama.cpp KV; 2,048-token chunks; three requests per runtime/context process. |
The default selection is gemma4. Downloads and existing files must match the manifest before use. fetch and offline verify need no model or native addon; only the optional tokenizer check uses the addon to compare current rendering against every pinned token ID and prompt. Qwen verification sets preserve_thinking=true in a temporary tokenizer copy to match its native session; the model files are untouched. Its package contains only measured cases, with the original messages and saved rendered prompts; warmup used each same context for 16 output tokens. The unused synthetic warmup in the historical preparation file is excluded.
In GitHub Actions, after checkout, Node.js setup, and yarn install --immutable, restore both fixtures with:
- name: Download pinned benchmark fixtures
run: |
yarn oxnode scripts/benchmark-fixture.ts fetch --fixture gemma4
yarn oxnode scripts/benchmark-fixture.ts fetch --fixture qwen38This step requires no secrets and verifies the archive, decoded payload, and all token-ID hashes before saving. The manifest's storage.url also supports ordinary HTTPS clients. The public r2.dev endpoint is rate limited; a custom domain can be added if CI download volume grows. The same payload and hashes must be retained when changing its host.
Keep each fixture version unchanged and record runtime, cache, timing, and load conditions for new measurements. Historical messages contain tool results, but replay must not execute tools. These fixtures do not assert generated-answer parity or universal speed thresholds. See the Gemma research reference and each manifest for restoration and protocol details.
Per-generation profiling is exposed from @mlx-node/lm:
import { enableProfiling, disableProfiling } from '@mlx-node/lm';
enableProfiling();
// run generations...
disableProfiling();The store lives in crates/mlx-core/src/profiling.rs (global PROFILING_STORE: Mutex<Vec<GenerationProfile>>, gate via PROFILING_ENABLED: AtomicBool). NAPI exports: setProfilingEnabled, isProfilingEnabled, getProfilingData, resetProfilingData.
The per-generation profiler (crates/mlx-core/src/decode_profiler.rs) records:
- TTFT (
time_to_first_token_ms) - Phase breakdown:
forward,sample,eval_token,extract,async_eval - Memory snapshots before / after each generation
Note: MLX lazy evaluation means
prefillMsmeasures only graph construction (~1 ms). UsetimeToFirstTokenMsas the real prefill latency indicator.
| Var | Effect |
|---|---|
MLX_PROFILE_DECODE=1 |
Auto-enables profiling at startup |
MLX_PROFILE_DECODE + trace |
Per-cycle timeline; set it with MLX_METAL_COMMAND_TRACE=1. A DSpark/DFlash2 turn dumps [dspark-span] JSON lines at turn end: a span per phase per cycle (dspark_cycle, _propose, _verify, _accept, _commit, _emit_loop, _cache_clear, _eval_boundary) on the [metal-command] clock, and a cycle marker (depth, keep). See tools |
MLX_NODE_LOG |
Tracing-level filter |
MLX_INFERENCE_TRACE_FILE |
Path for inference trace dump |
MLX_DEBUG_GEMMA4_DUMP |
Diagnostic dumps for Gemma4 |
MLX_DFLASH2_PHASE_TIME |
Per-cycle stderr timing of the DFlash2 draft-build / draft+selector / verify-build / verify phases. Forces extra evals to isolate the phases — diagnostic only, never set in production. |
MLX_DFLASH2_VERIFY_LAYERS |
stderr timing per 15-layer block inside the DFlash2 verify forward (forces mid-graph evals — diagnostic only). |
MLX_DFLASH2_TF_RECORD=<dir> |
Measurement tooling, not an A/B switch: a DSpark/DFlash2 turn writes its generated ids to <dir>/ref-<prompt key>.json (kept when an existing file is longer). |
MLX_DFLASH2_TF_DIR=<dir> |
Measurement tooling: teacher-forced replay. The reference ids replace the drafts as the verify input (same compiled verify path); the turn commits only reference ids and stops when they end. Appends [pos, len, a_ref, a_live, flip, wall_us] per cycle to <dir>/cycles-<label>.jsonl. Missing reference or anchor mismatch errors. See research/splash-qwen38.md. |
MLX_DFLASH2_TF_LABEL |
Label for cycles-<label>.jsonl (default run). |
| Var | Effect |
|---|---|
MLX_NO_COMPILE=1 |
REMOVED (dead). The Qwen3.5 compiled C++ forward path this var gated was deleted in the chat-engine refactor (ee88b92b). The former switch is unrelated to the current eligible DFlash compiled verifier. Nothing reads this var today. |
MLX_EVAL_ALL_CACHES=1 |
REMOVED (dead). Token-only eval is the only strategy; nothing reads this var today. |
MLX_WEIGHT_MATERIALIZE_CHUNK_MB |
Weight-loading chunk size |
MLX_GDN_KERNEL=perstep|chunked_ops |
Controls the retained device-agnostic/CUDA recurrence path: eligible long, unmasked non-Metal prefill uses chunked ops by default; perstep selects the serial ops reference. Decode and masked calls retain per-step behavior. Metal always uses its per-step kernel; the native Metal chunked experiment and its chunked selector were removed. |
MLX_LFM2_PAGED_PREFILL_PAGED_ATTENTION |
Opt-in (default off): multi-turn LFM2 cache-hit prefill tries the graph-native paged-attention bridge (gather_kv_for_prefill_chunk) before read_kv_range, skipping the forced per-layer pool sync. Held opt-in pending a stable-checkpoint paged-vs-flat gate (fused-vs-masked-SDPA ~1-ULP divergence) |
MLX_LFM2_FUSED_BATCH_DECODE |
Opt-in (default off): LFM2 multi-row decode waves run one fused [N,1] forward instead of the default row-exact wave (per-row N=1 forwards, bit-identical to serial). Fused M=N GEMM tiles round differently than M=1 and the carried ShortConv state amplifies the ULP diff into occasional greedy near-tie flips |
MLX_BONSAI_HOIST_METADATA=1 |
Experimental, default off: materialize rotated PQ2 projection scales/biases in FP32 once at load. Rotated projections promote their transformed inputs to FP32 (the 16-bit residual stream a normal Prism load produces), so outputs propagate FP32; adds about 0.8 GB for Bonsai 2 27B. No stable speedup is claimed. |
MLX_BONSAI_SHARE_HADAMARD=1 |
Experimental, default off: reuse identical signed Hadamard transforms within one paged Qwen3.5 decode step. Cache entries retain input/sign owners and are cleared on return or error; flat inference and prefill do not retain entries. |
MLX_QMM_SPLITK_MIN_M |
Diagnostic override for the quantized-matmul vector/GEMM boundary: M ≥ this many rows takes qmm/split-k instead of qmv/qmv_wide. Read by MLX's affine/fp path (fork patch, mlx-fork.md) and by the K-quant and mixed-affine dispatchers. A/B tuning only; the golden gates refuse it. |
MLX_KQUANT_SMALL_M_BENCH=1 |
Body-level opt-in for the #[ignore]d kquant_small_m_bench integration test: exact-shape K/IQ small-M quantized-matmul timings (MLP + lm_head shapes, M sweep). Run MLX_KQUANT_SMALL_M_BENCH=1 cargo test -p mlx-core --test kquant_small_m_bench -- --ignored --test-threads=1 --nocapture (serial — the two benches share the GPU). |
MLX_METAL_COMMAND_TRACE=1|2|3 |
Diagnostic (fork patch); any value ≥ 1 enables it. 1: [metal-command] per command buffer (with dispatches, primitives), [mlx-evaluation] per eval, [mlx-compiled] per compiled call. 2: plus [metal-op] <primitive> per primitive. 3: plus dtypes. Perturbs timing. See trace interpretation. |
Three Metal routes added with the tiled K-quant work are fixed: 2-D K-quant weights
with N % 64 == 0 and K % 256 == 0 run in the 64-row tiled layout (@t64 mode
suffix). The GGUF importer now stores eligible linears tiled on disk and marks each
one "layout": "t64" in its quantization entry (the mode string stays bare, see
convert-quantize.md), so such a checkpoint loads
with no permute; a row-major (legacy or secondary-output) K-quant tensor is still
permuted at load on a Metal host. 3-D expert weights, embeddings,
router gates and odd widths stay row-major on the row-major kernels (the
sorted MoE expert matmul at prefill widths, B / E >= 4, takes a tensor-op
gather_qmm_rhs_nax kernel with one expert per 32/64-row tile on gen-17+
GPUs and aligned N % 64, K % 64; the simdgroup gather_qmm_rhs stays for
the rest, and gather_qmv serves decode and the 8-row verify);
a tiled K-quant matmul of 8 rows (the DFlash2 verify block, MTP verify) takes
the tensor-op kernel's 8-row tier (qmm_m8_nax_t64) and, row-major at M = 8,
Q3_K and IQ4_NL only, with the K split count derived from the IORegistry GPU
core count — the Qwen3.8 GGUF route is what it shipped as; the tiled affine
contracts (a4g64@t64, a8g64@t64) take the 8-, 16- or 32-row tier at
M = 8..32 (qmm_m8/m16/m32_nax_t64: the weights are dequantized once per tile
whatever the row count, x staged per step with zero rows past M), their split
count also bounded by the fp32 partials the weight bytes can absorb, and the
tier is skipped (qmv_wide / the GEMM) when that grid falls under 6 threadgroups
per core at M = 8 or 2 above; and a DFlash2 verify attention block
takes a block kernel (tensor-op where the device has one, else the simdgroup-matrix
tile kernel) from a key count calibrated once per process at first use (GPU-timed
vector vs block at 256..4096 keys, block must win by 5%; ~35 ms), below which the
bit-exact vector kernels serve it. The one-time [kquant route] and [sdpa route]
log lines appear under MLX_METAL_COMMAND_TRACE.
The tiled kernels also carry MLX's affine quantization as modes (a4g64,
a8g64: the same LSB-first codes, one bfloat16 scale and bias per group,
mlx_kquant.h Mode::A4G64 / Mode::A8G64; the kernels are parametric, so
other (bits, group) combinations are instantiation lines away), in the @t64
layout only. On a Metal host every 2-D MLX-affine linear with bf16 companions at
4/64 or 8/64 (the pairs mlx convert emits: bodies, routers, lm_heads) of at least
[N = 6144, K = 1536] (kquant_affine_tiled_shape) is tiled at load by the same
arm that tiles K-quants (try_build_affine_quantized_linear_tiled in the Qwen3.5
dense/MoE, MTP, Gemma4, Muse-Glimmer, LFM2/K2 loaders) and retagged a4g64@t64 /
a8g64@t64 — same bytes, the K-quant kernel family, the row-major sources released
from the map as layers install. The shape rule is where the tiled kernels stop
paying off against MLX's route (M5 Max, dependency-chain and DRAM-cold
kquant_small_m_bench, t64/MLX): M = 1 ties from N = 6144 up (0.96-1.09x) and
trails below it (N = 2560..4096: 0.92-0.96x); the M >= 8 tiers win from N = 4096
up (1.1-1.6x) and lose at N <= 3072 (0.8-1.0x). So Qwen3.5-4B tiles its
gate|up (9216 x 2560), in_proj_qkvz (12288 x 2560, with the padded
in_proj_ba) and lm_head, and keeps o_proj, out_proj and down_proj (N = 2560)
on MLX's route; Gemma-4-E2B tiles gate|up (6144 x 1536) and the per-layer
embedding/lm_head widths and keeps its attention projections (N = 256..2048) and
down_proj. The DFlash2 draft tiles every tileable projection (gated on its own
numbers, tile_kquant_layout_any_shape). What stays on MLX's row-major affine
route besides the rule: f16 companions (GGUF affine imports, also g32), other
(bits, group) pairs (the 2/3/5/6-bit tensors of a mixed checkpoint, 8/32), shapes
that are not whole tiles and super-blocks (N % 64, K % 256; the GDN in_proj_ba
pads to whole tiles so it still merges with the tiled in_proj_qkvz), calibrated
activation-fp8 sites (input_amax), Prism hadamard sites, Gemma4 Q4_0 decode
sidecars, tied lm_heads / packed embeddings (gather), 3-D experts (gather_qmm),
and non-Metal hosts. The retag is inference-only: KQuantMatmul has no VJP, as
for every tiled K-quant (the training paths read the dense .weight params, not
QuantizedLinear). Measured against MLX's own route on an M5 Max (DRAM-cold,
kquant_small_m_bench): M = 1..7 at parity (±5%), M = 8 1.2-2.4x where the
tensor-op grid fills the GPU (tie on the small shapes the grid rule leaves to
qmv_wide), M = 9..16 1.2-4x, M = 17..32 0.8-3x (behind only on the wide-N short-K
shapes, N ≥ 12288 at K = 2560, where MLX's GEMM reaches 220 GB/s). The M = 1
qmv_t64 owns 8, 16 or 32 rows of a tile per threadgroup, the largest whose grid
still holds 16 threadgroups per core (the narrow Gemma-4-E2B shapes would otherwise
leave the GPU 1-2 threadgroups per core and the decode step, a dependency chain,
pays the matvec's latency: Gemma-4-E2B AR decode measured 5% behind MLX's route on
32-row threadgroups, within noise on 8/16-row ones). Every tiled route is
bit-deterministic across evaluations (kquant_affine_tiled
affine_tiled_routes_are_deterministic; the chunk partials are summed in a fixed
order through threadgroup memory). Gemma4, Muse-Glimmer and K2-Horizon greedy
decode is likewise reproducible run to run: their grouped attention partition
counts (D512, D128) follow one rule of the context, the query heads and the GPU
core count (below), not a timing sweep or a context table. The remaining
decode-tuning sweep (MLX_NODE_LOG=mlx_core::decode_tuning=info) picks only the
submission depth, which changes when graphs are submitted and never their
arithmetic.
Eligible Qwen/DFlash projection merges, fused GDN preparation/window convolution,
fused draft convolution/top-16/greedy selection, segmented one-call verifier attention,
compiled verification, eligible two-column GDN recurrence and D256 full prefill
attention now use their optimized paths without manual rollback switches.
Runtime shape, dtype and device guards retain required fallbacks; D256 FP32
attention still requires TF32 support.
GDN snapshots alias immutable state; prefill chunks submit asynchronously.
DFlash uses its checkpoint proposal width (seven for this companion); the shared
mtpDepth override remains a native-MTP option. Q8 head reuse is unconditional
within Q8 and does not change the BF16 precision default.
The experimental K/IQ residual/SwiGLU epilogues, empty-finalize shortcut,
eager-commit submission hook, direct-convolution replay shortcut and conditional
pressure-cache trimming have been deleted, along with the E48 four-column
Qwen override. Qwen4's separate four-column route remains. Ordinary boundary
submission, lazy context publication and periodic cache clearing are retained.
Historical evidence remains in the research reports.
Full end-to-end screening also rejected dense Q4, adaptive fallback, and the
forced matrix/command-buffer presets. The Splash/Qwen3.8 reference
records completed validation and final performance: exact output/acceptance
parity, with no established throughput change from cleanup.
| Var | Effect |
|---|---|
MLX_PAGED_DECODE_CACHE_CLEAR_INTERVAL |
Override decode-time clear_cache cadence |
MLX_PAGED_PREFILL_EVAL_INTERVAL |
Override prefill eval cadence |
MLX_PAGED_PREFILL_CHUNK_SIZE |
Prefill chunk size |
MLX_PAGED_NATIVE_KV_WRITE=0 |
Disable the default-on graph-native paged K/V write for every paged family. Whole-turn decode falls back to the synchronous write (also taken automatically on a native-write error). Batched and ragged decode instead hard-error, so this switch turns continuous batching off for every paged family |
MLX_PAGED_GRAPH_DECODE_GATHER=0 |
Disable the default-on graph-native paged decode gather for every paged family. Whole-turn decode falls back to the synchronous gather (also taken automatically when the graph gather cannot serve the inputs). LFM2 and Nemotron-H batched decode instead hard-error. The synchronous gather reads the same pages (decode_read_span: sliding groups from their first live block) and takes the same planned partitions, so it gives the graph route's bits wherever it has the same kernel (generic V2, grouped D512); grouped D128 is graph-only, so Muse-Glimmer's and K2-Horizon's global layers run generic V2 here and the transcript differs |
MLX_PAGED_GROUPED_QWEN35=0 |
Disable the default-on BF16 D256/BS16 grouped decode kernel (24Q/4KV above 16K q1 / 8K q2; 16Q/2KV above 32K q1 / 16K q2) for same-binary A/B or driver rollback |
MLX_PAGED_GROUPED_D512=0|force |
Disable the direct-read BF16 D512/BS16 kernel or force a supported geometry for diagnostics. Ordinary single-row Gemma AR takes it for every capability-confirmed geometry beyond 512 tokens; other paths retain their existing policy. Metal pipeline and shape guards always apply. |
MLX_PAGED_GROUPED_D512_STRIPES |
Diagnostic partition override (4|8|16|32|64|128|256: the counts the D512 dispatcher admits, so it never exceeds the rule's 256 cap; a count above the context's page count runs empty stripes). It does not enable grouped attention by itself; pair with MLX_PAGED_GROUPED_D512=force when testing an unmeasured geometry. Gemma AR otherwise uses the smallest power of two whose query-head SIMD groups reach 16 per GPU core (IORegistry gpu-core-count), bounded by the 16-token work tiles and 256; the bf16 result changes with the partition count, so a forced value changes the greedy transcript. |
MLX_GEMMA4_PAGED_DECODE_ROUTE |
Diagnostic Gemma decode route (auto|sdpa|paged). Auto takes the grouped kernel when available; other routes keep their geometry and live-memory guards. |
MLX_GEMMA4_DECODE_EARLY_EVAL_LAYERS |
Diagnostic override for how many leading eager decode layers are submitted early. Set before process startup; default depth is learned from completed token timings rather than a chip-name table. |
MLX_MUSE_GROUPED_STRIPES |
Diagnostic Muse-Glimmer partition override for single-row paged decode in either lane (0 = generic V2, or 4|8|...|1024). It replaces the device rule below but not its hard cap: a larger value runs at the live limit from mlx_paged_grouped_d128_max_stripes (kernel limit, work tiles, temporary storage, buffer length), and generic V2 when that limit is below 4. The bf16 result changes with the partition count, so a forced value changes the greedy transcript. Set before process startup. |
MLX_MUSE_DECODE_EARLY_EVAL_LAYERS |
Diagnostic override for how many leading Muse-Glimmer decode layers are submitted early. Set before process startup; default depth is learned from completed token timings rather than a chip-name table. |
MLX_GEMMA4_MIXED_QMV=0 |
Restore stock QMM for single-row affine Q4/group32 BF16 decode. The default Metal kernel reads original FP16/FP32 sidecars directly and derives a symmetric bias only when checkpoint metadata authorizes it. Set before loading the model. |
MLX_TEST_PAGED |
Test-only paged-path toggle |
MLX_PAGED_CACHE_INITIAL_MB |
Initial pool size (u32 MiB) for the qwen3_5 / qwen3_5_moe grow-on-demand paged pools. Env wins over the paged_cache_initial_memory_mb config field; a set-but-unparseable value falls back to config, and unset everywhere keeps the historical fixed pool (initial == max). The pool doubles on exhaustion up to the max (paged_cache_memory_mb / the auto one-full-context default). The agent override manager writes paged_cache_initial_memory_mb: 2048 into qwen3_5 clones by default |
MLX_PERSIST_PAGED_CACHE |
Override the per-model persistPagedCache config for the SSD-backed cold KV tier. Precedence: env > config alias > off. Parsed leniently (1/true/on/yes → on, 0/false/off/no → off; unset/other falls through to config). Off by default at library level |
Gemma's and Muse-Glimmer's submission-depth calibration is a bounded search over
layer-count depths, not a guarantee of the global optimum. It excludes first-use
timings and requires wins above measured jitter. Decisions stay within a loaded
model (up to eight context scales); batched and speculative paths keep their
existing scheduling. Calibration that occurs during a request contributes to
that request's latency.
Attention partitions are deliberately not searched: the grouped kernel's stripes
accumulate online-softmax partials in bf16 over different page subsets, so two
partition counts differ by a few bf16 ulps in the logits and a timing-based pick
made greedy transcripts differ run to run. Gemma4 D512, Muse-Glimmer D128 and
K2-Horizon D128 (ForceD128 without a plan) use
engine::decode_tuning::grouped_partition_stripes: the smallest power of two
whose query-head SIMD groups reach 16 per GPU core, bounded by the 16-token work
tiles and a cap (D512: the reducer's 256; D128: the live limit from
mlx_paged_grouped_d128_max_stripes — kernel limit 1024, work tiles, temporary
storage headroom and buffer length — which stays a hard cap).
The work-tile bound changes the count inside a turn while the context is short.
On 40 cores, 8-head Gemma4 (E2B, 128 partitions) runs 32 partitions up to 1,008
tokens, 64 up to 2,032 and 128 beyond; 16-head Gemma4 switches once, at 1,009;
the 32-head D128 models (32 partitions) never switch beyond 512 tokens. Each
count is a pure function of the device, the query heads and the context, so a
turn still reproduces exactly. Any count at or above the page count gives the
same bits (the extra stripes are empty and add exact zeros), so dropping the
bound would hold one count per turn; it stays until the cost of those empty
threadgroups is measured on a quiet machine.
Muse-Glimmer plans single-row decode the same way in the scheduled lane and the
whole-turn lane (no cache owner, or MLX_SERVE_FORCE_SERIAL=1), so one prompt
gives one transcript in either lane. Multi-row scheduled waves keep generic V2:
the grouped kernels serve one sequence per dispatch, so a scheduled row's bits
depend on whether other rows shared its step (see
concurrent inference).
The native Gemma GGUF loader also records the nested text dtype so prefill metadata
can be prepared once, while Q4 decode consumes the original scale precision.
See the Gemma GGUF performance reference
for the numerical checks, memory tradeoff, and measured limits.
| Var | Effect |
|---|---|
MLX_CACHE_LIMIT_GB |
Hard Metal pool ceiling |
TTS_MLX_CACHE_LIMIT |
Optional TTS-lifetime ceiling on the shared MLX free-buffer pool, in GiB (fractional values allowed). Unset/0 adds no constraint; positive ceilings compose by minimum with the global policy. See TTS memory. |
MLX_GPU_HEADROOM_GB |
Headroom term in the auto-sizing formula |
MLX_RESIDENCY_SET_MAX_PCT |
Upstream MLX: one residency set's size, as % of the recommended working set (default 5). Splits wired memory across sets, never changes how much is wired; <= 0 or >= 100 = one set (mlx/utils.h:205-218) |
MLX_RESIDENCY_DEBUG=1 |
Upstream MLX: log each residency set as it is created |
MLX_NATIVE_GGUF_CACHE_MAX_GB |
Total-size cap on the native GGUF cache (packed weights, extracted tokenizer assets, companions), in GiB — fractional allowed, parsed like MLX_CACHE_LIMIT_GB. <=0, unset or unparseable = unlimited. When set, least-recently-used entries are evicted until under the cap; entries younger than 24 h, entries whose key lock is held, and the entry being loaded are never evicted, and .mlx-qwen4-assets-v2-* sidecars count toward the cap but are only removed through it. Publication itself is admitted atomically under a root-wide flock, so concurrent conversions on different keys cannot jointly exceed the cap — a conversion that still cannot fit after eligible eviction is not published and the load reads it from its staging directory instead — a reload then reuses that verified same-process staging dir rather than re-converting and leaking a copy per load. Refused loads keep a per-pid fallback dir (a completed staging dir, or .mlx-qwen4-assets-live-{identity}-{pid} for tokenizer sidecars) that never counts toward the cap and outlives every sweep while the process lives. Reclaimed bytes are logged |
MLX_QWEN4_PACKED_CACHE_MAX_GB |
Total-size cap across every Qwen4 prepared-SSD-cache identity directory, in GiB (fractional allowed). <=0 or unset keeps the per-checkpoint 96 GiB budget with no global cap. When set, least-recently-used non-current identity directories are evicted first, then the oldest .pack files of the current directory |
MLX_QWEN4_IGNORE_MEMORY_PRESSURE |
=1 skips the macOS memory-pressure gate only — the kern.memorystatus_vm_pressure_level poll (every 250 ms for up to ~3 s) inside available_memory, reached at Qwen4 bootstrap, admission and growth. Diagnostic opt-out; the free/file-backed headroom checks and the SSD-read safety limit (check_active_memory_budget, which never reads this flag) still apply |
Qwen3.5 / Qwen3.6 MTP (Multi-Token Prediction) speculative decoding adds
runtime knobs gating individual optimizations across the W6.5–Phase C perf
chain (the former "one unconditional warmup hook for verify prewarm" was
deleted with the compiled-C++ MTP path — nothing pre-warms the eager verify
at model load). The env vars are read at most once per process and cached;
the truthy/falsy vocabulary is uniform (1 / true / on and 0 /
false / off, case-insensitive, with trim()) everywhere except
MLX_NEMOTRON_MTP_DEFAULT, which matches the literal string 1 and
nothing else. The last row of the table is a per-family DEFAULT policy
switch rather than an optimization knob. The adaptive-depth knob
is a TypeScript ChatConfig field (not an env var) because it interacts
with
the user-set mtpDepth and needs per-session resolution.
| Knob | Default | Workstream | Direction | Notes |
|---|---|---|---|---|
| (eager verify prewarm) | W6.7 | REMOVED. The once-per-process 10-dummy-shape prewarm was deleted with the compiled-C++ MTP path (ee88b92b); the verify forward is now eager Rust, re-traced per cycle, and nothing warms it at model load. |
||
mtpAdaptiveDepth (TS field) |
ON* | W6.8 | per-session | TS ChatConfig field. * defaults ON when enableMtp=true and mtpDepth is unset; defaults OFF (pinned) when mtpDepth is set explicitly. |
MLX_MTP_CHAINED_CYCLES |
M5+: ON / M1–M4: OFF | W6.5 | config-gated | Each cycle's verify_hidden[K] slice seeds the next cycle's first MTP draft, skipping the Step-A target forward on full-accept cycles (slice batched into the next-cycle async_eval) — exactly MTPLX's 1-forward-per-cycle design. Default ON on M5+ (GPU arch gen ≥ 17), default OFF on M1–M4 (gen 13–16); MLX_MTP_CHAINED_CYCLES=0/false/off forces OFF even on M5+, =1/true/on forces ON even on M1–M4. CONFIG-DEPENDENT (see "Same-checkpoint engine head-to-head"): on M3 Max / bf16 / nvfp4 / hard prompt a controlled A/B found it neutral@depth1 / a regression@depth3 (Step-A only ~14% of cycle there + an unresolved lazy-slice acceptance regression); but on M5 Max / affine-int4 Step-A is ~37% of the cycle and removing it wins on BOTH a naive cold A/B (≈1.17×, 45.0→52.8 tok/s) and the controlled harness (+23.7pp ratio, 1.116 vs 0.879, beyond noise; acceptance also up 1.30 vs 0.91) — the biggest structural lever toward MTPLX parity. M5+ default ON is measured net-positive (affine +16% / nvfp4 byte-identical to AR); M1–M4 stays default OFF pending the lazy-slice eval-scheduling fix. |
MLX_MTP_BUCKETED_VERIFY |
W6.29 | REMOVED (dead). Per-bucket compiled verify graphs were deleted with the compiled-C++ MTP path (ee88b92b); nothing reads this var today. The measured long-decode win it documented (AR +12.0%, MTP +26.1% on qwen3.5-4b / M3 Max) is currently not implemented — the verify forward re-traces eagerly per cycle at whatever kv_len exists. Restoring per-kv-len compiled verify graphs is an open perf follow-up. |
||
MLX_MTP_ACCEPT_GATE |
ON | — | opt-OUT | MTP acceptance gate, depth-1-scoped, adaptive-exempt, and exact-binomial. When ON (default), the model aggregates the first-draft acceptance counts (accepted/attempted at draft slot 0) across completed fixed depth-1 turns and disables speculation for the NEXT fixed depth-1 turn only when an exact-binomial test at the 5% level shows the aggregate is inconsistent with the 0.6 break-even — plain AR — so a weak draft head does not pay the verify cost for zero speedup. The exact test preserves a uniform false-gate bound at every sample size: a 1-cycle turn, a 2-of-4 streak from a healthy 0.756 head (~25% of turns), and 0-of-2 / 0-of-3 / 1-of-5 aggregates (~5.95% / 1.45% / 1.43% false rates) never gate; only extreme aggregates like 0-of-4 (false rate ~0.35%) do. The aggregate is bounded (halved past ~512 samples) so a long healthy phase cannot drown out a later degradation. Depth > 1 turns and mtpAdaptiveDepth turns (which sweep depths 1-5) are never gated and do not publish history. State is per-model; the first turn of a model load probes, after 3 consecutive gated turns the gate re-probes, and a full session reset (reset_caches) clears the history. Measured first-draft acceptance on real checkpoints: 1.000 (qwen3.5-4b, counting), 0.756 (qwen3.5-4b, complex task), 0.000 (qwen3.5-0.8b — MTP is a ~0.6× net loss there). Set to 0 / false / off to always run MTP when requested. |
MLX_NEMOTRON_MTP_DEFAULT |
OFF | — | opt-IN | NemotronH only, and a default-policy switch rather than a perf knob. It gates NemotronHModel::mtp_auto_enabled() (models/nemotron_h/model.rs), the native getter ChatSession consults before auto-setting enableMtp for a caller who set nothing. Unset (the default) the getter returns false, so the family never auto-enables MTP even though hasMtpWeights() is true on every shipped checkpoint; =1 makes it return true, but only on a checkpoint that actually carries a head (the env read is ANDed with mtp_active). The one knob here that is off the shared truthy vocabulary: the read is std::env::var("MLX_NEMOTRON_MTP_DEFAULT").is_ok_and(|v| v == "1"), so true/on/TRUEdo NOT enable it. Cached in a process-wideOnceLockat the firstmtpAutoEnabled()call, so it is neither per-model nor changeable afterwards. It never blocks an explicit request —enableMtp: trueorgeneration_mode: "mtp" runs MTP with the var unset. Why the default is off (and why the throughput argument for it was retracted): docs/models.md. |
Phase C replaced the per-cycle MTP draft cache (zeroed every cycle —
heads saw only the in-cycle draft chain) with a persistent
committed-history cache: each cycle commits the full K+2 sequence
[last_committed, d_0..d_{K-1}, boundary] into a separate MTP K/V cache
so subsequent cycles' drafts attend the whole committed prefix.
Prompt-prefill seeds that cache from the prompt's hiddens before decode.
The committed-history cache is its own coordinate space, decoupled
from the main KV cache (the C++ begin_cycle anchors the draft RoPE
offset to g_mtp_committed_len, not the main offset). On turns that
skip prompt-prefill (cache-reuse / VLM / delta / prompt < 2 tokens) the
cache simply starts empty and fills
contiguously from decode tokens — internally consistent, and
speculative decoding stays verify-correct regardless of draft quality.
Measured on qwen3.6-27b-nvfp4-mtp / M3 Max / depth=3 / T=0, on a
moderately-predictable mixed prompt (see the prompt-dependence note
below — these absolute numbers are meaningless without the prompt):
| Path | mean accepted/cycle | per-position acceptance | MTP/AR decode |
|---|---|---|---|
| committed-history + prompt-prefill | 2.15 | [0.854, 0.715, 0.585] |
~1.31× |
| committed-history, no prompt-prefill | 1.75 | [0.812, 0.565, 0.381] |
~1.13× |
| (pre-Phase-C per-cycle cache) | 1.56 | — | < 1× |
Acceptance is dominated by prompt predictability, not by the checkpoint. The numbers above are NOT a fixed property of the checkpoint — on the SAME
qwen3.6-27b-nvfp4-mtp/ depth=3 / T=0, committed-history + prompt-prefill acceptance ranges from 1.44/cycle[0.735, 0.471, 0.235]on a novel three-paragraph prose essay (a deliberately HARD case) up to 2.76–3.00/cycle ([0.96–1.0, 0.92–1.0, 0.92–1.0], i.e. the depth-3 ceiling) on predictable text (counting, lists, recitation, repetition). MTP/AR decode speedup tracks acceptance directly: ~1.06× on the prose essay, 1.46× (AR 20.8 → MTP 30.3 tok/s) on a counting prompt. So when comparing against another engine's headline number (e.g. MTPLX's ~2.24×, measured on its own unpublished "recorded" prompt at T=0.6), the prompt and sampler dominate — a single acceptance figure is only meaningful alongside the exact prompt that produced it. Generation length does NOT move acceptance (essay is flat ~1.44→1.30 from 120→400 tokens). Verified 2026-05-30.
T=0 parity holds in distribution: every MTP-emitted token equals
argmax(verify_logits). MTP and AR outputs agree on a contiguous prefix
and then diverge at an isolated argmax near-tie — and that flip can land
early. The recurring offset-16 "Autumn is often regarded / described"
flip was diagnosed: at that token AR and the batched verify rank the
same top-2 tokens within one bf16 ulp (AR logits 21.500 / 21.375;
verify 21.375 / 21.375), so the verify forward merely tie-breaks to the
other token. One flip then decorrelates all downstream text. This is
benign lossless speculative decoding (vLLM / MTPLX / dflash-mlx all
document it) — not a verify-path bug. Because a near-tie can flip at
any offset, the MTP parity gate treats text divergence as
informational only; the blocking correctness gate is acceptance health.
The verify forward is one full 27B forward over T = depth+1 tokens and
is ~58–62 % of the MTP cycle; its cost grows with depth while later draft
slots accept progressively less. Measured on
qwen3.6-27b-nvfp4-mtp-oproj8 / T=0 / 256 tokens (depths 1–3 a
same-session sequential A/B; absolute ratios are thermal-sensitive
~±10 %, so the cross-depth ordering is the signal):
| Depth | M3 Max ratio | M5 Max ratio | M5 K̄ | M5 per-position acceptance |
|---|---|---|---|---|
| 1 | 1.14× | 1.15× | 0.87 | [0.865] |
| 2 | 1.12× | 1.15× | 1.42 | [0.811, 0.608] |
| 3 | 0.93× | 1.04× | 1.98 | [0.828, 0.656, 0.508] |
| adaptive | 1.07× | 1.12–1.13× | 1.00 | [0.86, 0.54, 0.40, …] |
Note (chained cycles is CONFIG-DEPENDENT — default ON on M5+ (gen ≥ 17), default OFF on M1–M4; see "Same-checkpoint engine head-to-head" below).
MLX_MTP_CHAINED_CYCLES=0/1overrides either direction. This table was measured with the Step-A target forward present every cycle (MLX_MTP_CHAINED_CYCLES=OFF). Chained cycles skips that forward on full-accept cycles, seeding the next draft fromverify_hidden[K]— exactly MTPLX's design (the engine that the same-checkpoint head-to-head below shows is 1.40× faster precisely because it never runs Step-A).The win/loss depends on how large Step-A's share of the cycle is and whether the chained-hidden's lazy slice stalls the pipeline:
- M3 Max / bf16 / nvfp4 / hard prose prompt (the original verdict): a controlled A/B (warmup + cooldown + alternating-order interleaved pairs, 6 repeats/cell, T=0; cross-config AR drift 1.0% ⇒ fair, self-normalized) found it not a win — depth 1 1.11× ON vs 1.08× OFF (inconclusive, within ±11pp noise); depth 3 0.90× ON vs 0.96× OFF (a regression), driven by an acceptance drop (meanAccepted 1.28 vs 1.43) plus a lazy-eval scheduling stall on the chained-hidden slice. Here Step-A is a smaller fraction of the cycle, so removing it doesn't pay.
- M5 Max / affine-int4 (2026-05-30): Step-A is ~37 % of the cycle. Both a naive cold count A/B (45.0 → 52.8 tok/s, ≈1.17×) AND the controlled harness on the affine checkpoint (depth 3, self-normalized ratio) agree it WINS: chained-ON ratio 1.116 vs OFF 0.879, +23.7pp beyond the ±19.8pp noise band, with acceptance higher under ON (1.30 vs 0.91) — no acceptance penalty here, the opposite of the M3/nvfp4 case.
So it is now config-gated by GPU arch gen: default ON on M5+ (gen ≥ 17), where it is measured net-positive (affine +16 % / nvfp4 byte-identical to AR); default OFF on M1–M4 (gen 13–16) until the chained-hidden lazy-slice eval-scheduling stall (the depth-3 acceptance regression) is fixed so it also wins on M3/bf16/nvfp4.
MLX_MTP_CHAINED_CYCLES=0/1overrides either direction. On M5-class + affine-quant it is the single biggest structural lever toward MTPLX parity. The earlier "Step-A is only ~14 % of cycle wall" claim is config-specific (M3/bf16); on M5/affine it is ~37 %. (A still-earlier naive same-binary A/B reported a spurious depth-1 "win" from a thermally-confounded run — the flag-invariant AR baseline had swung ~2×; the controlled harness exists to defeat exactly that confound.)
Depth 1 is optimal on both M3 Max and M5 Max — the 3rd draft slot's
~50 % acceptance still does not pay for the wider, slower verify forward.
On M5 Max depth 3 climbs from a net regression (0.93×) to marginally
positive (1.04×) — the Neural Accelerator does shave a little off the
wider verify — but it still loses to depth 1. The W6.8 adaptive policy
underperforms the depth-1 pin on both hosts. For this hardware/model
class, pin mtpDepth: 1.
M5 Neural Accelerator does NOT widen the MTP/AR gap. The headline
MTP/AR ratio at the optimal depth is essentially hardware-invariant: 1.14×
on M3 Max, 1.15× on M5 Max. The reason is symmetry — stock MLX qmv and
the AR forward path both benefit equally from the NA, so the MTP cycle's
verify forward gets faster in the same ratio as the AR baseline it is
measured against. See "Verify-kernel investigation" below for the M5 Max
microbench result that pins this down.
The verify forward is the bottleneck, and four attempts to make its
small-M (M = depth+1) quantized matmuls cheaper all failed:
| Attempt | Result on M3 Max | Result on M5 Max |
|---|---|---|
Batched GEMM (qmm / qmm_t_splitk) at M=4 |
~38 % slower than qmv; split-K low-precision accumulation also degrades acceptance → ratio collapses to ~0.75× |
(not re-tested — split-K accuracy regression is hardware-independent) |
multi3 multi-row qmv (W6.30 port) |
0.94× geomean vs stock qmv — slower even in its native affine-4-bit M=3 envelope |
0.94× geomean — identical to M3; the M5 Neural Accelerator helps stock qmv symmetrically (kernel + microbench removed as a proven dead-end, 2026-06-01) |
SPARSE_ACCEPT |
zero measured win (W6.19) | (not re-tested — sync-collapse knobs are CPU-side; hardware-invariant) |
VERIFY_ASYNC_EVAL (W6.9) |
zero measured win on M3 Max (overlap budget too small) | +5–6% ratio at depth 3 on qwen3.6-27b-nvfp4-mtp-oproj8 (~1.07× → ~1.13×) with byte-identical acceptance. Flipped to default ON in Phase 3 (2026-05-26). |
| End-to-end MTP/AR ratio at optimal depth | 1.14× (depth 1) | 1.15× (depth 1) — Neural Accelerator does not widen the gap |
Stock MLX's small-M qmv is already near-optimal on both M3 and M5 Max:
at these shapes the verify is not dominated by per-row weight re-reads,
so reading weights once (the qmm / multi3 approach) buys nothing and
the extra machinery costs more. Crucially, the M5 Neural Accelerator
benefits stock qmv and multi3 in the same proportion, so the
microbench ratio is invariant — and the end-to-end MTP/AR ratio is
invariant too, because the AR baseline gets the same NA uplift as the
MTP verify forward.
Why MTPLX reaches ~2.2× and this does not — it is NOT the hardware, and NOT a fork.
A direct audit of the MTPLX source (2026-05-30) settles this. MTPLX ships
stock PyPI mlx 0.31.2 — there is no required private fork. The
mlx-mtplx-…-qmm fork is optional (surfaced only as
optional_fast_mlx_fork_active), and its only content is the same small-M
qmv retune we already benchmarked at 0.94×. Its custom verify kernels
(multi-row qmv, fused-MLP, top-k lmhead) are default-OFF probes whose
own docstrings call them "slower than stock." MTPLX's headline 2.24× is an
acceptance multiplier, not a kernel one: depth-3 + T=0.6 + fan-pinned
M5 Max on a checkpoint hand-tuned for ~95% acceptance
(Qwen3.6-27B-MTPLX-Optimized-Speed, per-position [1.00, 0.98, 0.94]).
The monotonic depth ladder (AR 24.6 → D1 41.6 → D2 48.6 → D3 54.5 tok/s) is
driven by acceptance, not faster GEMMs. The one genuinely valuable custom
kernel they have — a fused GatedDeltaNet recurrence-from-conv-tape — is a
verify-dispatch micro-opt for recurrent layers, not a multiplier.
(dflash-mlx's 2.95–4.4× is a _different architecture — a separately-trained
block-diffusion drafter — out of scope for native MTP heads here.)
The "acceptance multiplier" story above is only half the MTPLX gap. To
isolate the engine from the checkpoint and the prompt, we ran MTPLX's own
engine on its own checkpoint (Youssofal/Qwen3.6-27B-MTPLX-Optimized-Speed,
affine-int4) against our engine on the identical checkpoint — same M5 Max,
same count prompt, T=0 greedy, depth 3, no fan-pinning on either side
(MTPLX thermal: not configured). Our loader reads their checkpoint fully:
the trunk via the standard quantization key, and the mtp.safetensors
sidecar (merged in as a model-*-of-*.safetensors shard) with the affine
group_size inferred per-tensor from the scale-column shape — so the
non-standard mtplx_mtp_quantization key being unread does not mis-dequant
the heads (verified: they hit the perfect [1,1,1] depth-3 ceiling on
counting).
| Metric (count prompt, T=0, depth 3, no fan-pin) | Ours | MTPLX | Winner |
|---|---|---|---|
| AR decode | 25.82 tok/s | 22.74 tok/s | us, 1.14× |
| MTP acceptance / cycle | 2.94 [1,1,1] |
3.0 [20/20,20/20,20/20] |
tie (both perfect) |
| MTP D3 decode | 40.18 tok/s | 56.43 tok/s | MTPLX, 1.40× |
So at matched (perfect) acceptance and with our faster AR forward, MTPLX's MTP cycle is still 1.40× faster. That residual is NOT acceptance, NOT quant format, NOT the verify kernel, NOT hardware — it is a structural difference in trunk forwards per cycle:
- MTPLX runs 1 full-trunk forward per full-accept cycle — the
D+1-wide verify (generation.py:4737). The next cycle's draft seed is the same forward's last-position hidden (hidden = verify_hidden[:, -1:, :],generation.py:5037) and the bonus token is sampled from its logits (pending_primary,generation.py:5051-5068). No separate target forward. - We run 2 — a single-token Step-A main-model forward, gated by
do_step_a(nowengine/mtp_turn.rs, driven by the stepper'sforward_with_hidden) — plus theD+1verify. The Step-A forward exists only to produce the next draft seed, which the verify forward already computes and we discard on the default path. (At the time of the original measurement this ran through the now-deletedforward_compiled_with_hidden; the eagerDenseMtpStepperpreserves the same 2-forward cycle structure.)
Timing arithmetic matches the measurement exactly: our single-token trunk
forward ≈ 38.7 ms (25.82 tok/s), 4-wide verify ≈ 58 ms, drafts ≈ 9 ms.
Ours = 38.7 + 9 + 58 ≈ 106 ms / 4 tok ≈ 40 tok/s (measured 40.18); MTPLX
= 9 + 58 ≈ 67 ms / 4 tok ≈ 56 tok/s (measured 56.43, verify_ms_per_call=57.97,
draft_time_s/20 ≈ 9.2 ms). The ~38.7 ms Step-A forward — ~37 % of our
cycle on this M5/affine config — is the entire gap.
The lever exists and is wired, now default ON on M5+ (gen ≥ 17):
MLX_MTP_CHAINED_CYCLES skips Step-A and seeds the next draft from
verify_hidden[K] — exactly MTPLX's trick, parity-exact at T=0. (Set =0 to
force it OFF on M5 for a bisect.) Two measurements on this checkpoint agree it
wins:
- Naive cold count A/B: 45.0 → 52.8 tok/s (≈1.17×), closing most of the gap to MTPLX's 56.4 (remaining ~7 % is second-order).
- Controlled harness (self-normalized MTP/AR ratio, interleaved, 4 repeats, depth 3): chained-ON ratio 1.116 vs OFF 0.879 — chained-ON WINS by +23.7pp, beyond the ±19.8pp noise band. Acceptance was higher with ON (meanAccepted 1.30 vs 0.91, per-position
[0.633, 0.433, 0.237]vs[0.638, 0.213, 0.064]), so there is no acceptance penalty here — unlike the M3/nvfp4 case. (The AR-stability gate flagged thermal instability, CV 15.2%, so absolute tok/s are discounted; the self-normalized ratio and deterministic acceptance are what the harness is designed to trust through exactly that.)
This overturns the earlier "not a win" verdict for this config (see the
chained-cycles note above): that verdict was measured on M3 Max / bf16 /
nvfp4 / a hard prose prompt, where Step-A is a smaller fraction of the cycle,
the chained-hidden's lazy-eval slice introduced a scheduling stall, AND chaining
lowered acceptance. On M5 Max / affine-int4, Step-A is ~37 % of the cycle,
removing it clearly wins, and acceptance does not regress. The flag is
config-dependent, not globally good-or-bad — so it is config-gated by GPU arch
gen: default ON on M5+ (gen ≥ 17), default OFF on M1–M4 (gen 13–16), with
MLX_MTP_CHAINED_CYCLES=0/1 overriding either way. Remaining before flipping
the default ON on M1–M4 too: (1) resolve the chained-hidden lazy-slice
eval-scheduling stall so it also wins on M3/bf16/nvfp4; (2) confirm the depth-3
acceptance asymmetry ([…,0.213,0.064] OFF vs […,0.433,0.237] ON — the deeper
draft slots accept far better under ON on the warm essay) is a genuine effect,
not a warm-cache Step-A-path artifact worth fixing on the OFF path too.
Bottom line: the full MTPLX advantage decomposes into (1) an acceptance component (the ~2.24× headline = T=0.6 + a prompt tuned for ~95 % acceptance) and (2) a structural per-cycle-forward component (~1.40× at matched acceptance, = the Step-A forward we run and they don't). (1) is a prompt/sampler artifact; (2) is a real, addressable engine win via Step-A elimination.
Native-heads MTP/AR ceiling ≈ 1.1–1.15× on both M3 Max and M5 Max
(depth 1, with Step-A present). The 1.40× same-checkpoint engine gap
above shows the ceiling is NOT intrinsic to native heads — eliminating the
redundant Step-A forward (already implemented behind
MLX_MTP_CHAINED_CYCLES) recovers most of it without any custom
verify-path kernel, private fork, or the dflash separately-trained-drafter
architecture. M5-class hardware does not unlock it; cycle structure does.
Long-context behaviour. The verify forward's attention cost scales with context length, eroding the speculative advantage on long prompts: on a ~1k-token prompt the MTP/AR ratio drops toward parity even though per-cycle acceptance stays healthy. A future adaptive context-length guard could fall back to plain AR decode once the prompt length crosses the break-even point.
Interactions:
- Setting
mtpDepthexplicitly disables adaptive depth by default; passmtpAdaptiveDepth: truealongside to keep adaptation enabled withmtpDepthas the initial seed.
Cross-references:
- TS field JSDoc:
enableMtp/mtpDepth/mtpAdaptiveDepthonChatSession.sendinpackages/lm/src/chat-session.ts. - Source of truth (env-var readers + inventory table):
crates/mlx-core/src/models/qwen3_5/mtp_decode.rs(mtp_chained_cycles_enabled,mtp_accept_gate_enabled). The W6.29 bucketed-verify dispatcher no longer exists — it was C++-side state in the deletedmlx_qwen35.cpp. The verify async-eval, deferred-hidden, sparse-accept, batched-target and prompt-prefill knobs are gone: each of those behaviors is now unconditional. - Phase C committed-history MTP cache: eager policy in
crates/mlx-core/src/models/qwen3_5/model.rs(DenseMtpStepper::begin_cycle/commit_mtp,committed_len); the old C++ policy filemlx_qwen35_mtp_compiled.cpp(g_mtp_committed_len) was deleted with the compiled path; prompt-prefill seed lives inbegin_mtp_decode's prompt-prefix seed block (viachunked_prefill_with_hidden). - W6.8 adaptive-depth policy:
crates/mlx-core/src/models/qwen3_5/adaptive_depth.rs.
token.eval()immediately after sampling — without it MLX builds an unbounded lazy graph.synchronize_and_clear_cache()every 256 steps — prevents memory accumulation during long generations.- Dtype-aware scalar ops — any
f32scalar in a binary op with bf16 promotes the entire result to f32. - Token-only eval — caches materialize through the dependency graph; no need to eval every cache tensor explicitly.
- For bf16 / f16 data extraction: use
to_uint16_native()instead of round-tripping through f32.
mlx_gpu_architecture_gen() (FFI in crates/mlx-sys/src/lib.rs) returns a generation number:
| Chip | Gen |
|---|---|
| M1 | 13 |
| M2 | 14 |
| M3 | 15 |
| M4 | 16 |
| M5 | 17 |
Qwen3.5 Metal GDN recurrence uses the per-step kernel. The native Metal
chunked-prefill experiment has been removed. Its former shader was scalar-FMA
plus simd_sum, without simdgroup_matrix or NAX matmul instructions.
Historical measurements on June 4, 2026, found it 2.8–3.5× slower in prefill
TTFT on M5 Max (24–31× per isolated call) across 580–5,384 prompt tokens, and
about 2× slower on M3. The two algorithms also differed by 1–2 BF16 ULP and
could change greedy continuations. These are historical measurements, not
validation of the final cleanup binary. See the
Splash/Qwen3.8 reference for the later removal
screen and its limitations.
The separate CUDA/non-Metal chunked_ops implementation remains. Eligible
long, unmasked prefill uses it by default; MLX_GDN_KERNEL=perstep selects the
serial ops reference, while chunked_ops retains the chunked route. Its decay
is computed directly in log space so strong decay does not produce log(0).
This path and its numerical regressions were not removed with the Metal shader.
| Scheme | How it's invoked |
|---|---|
| 4-bit affine | mlx_quantized_matmul (mode affine, configurable group size and bits) |
| MXFP4 | --q-mode mxfp4, or the low FFN class in --q-recipe unsloth --q-mxfp; 4-bit microscaling with group size 32 |
| MXFP8 | mlx_gather_qmm with mode="mxfp8" (used for MoE expert routing); returns [quantized, scales] |
| NVFP4 | --q-recipe unsloth --q-mode nvfp4; low-class FFNs use NVFP4 4/16 weight storage with standard MLX A16 quantized matmul in the fixed DGX class map |
| FP8 E4M3 import | Source .weight_scale_inv checkpoints are dequantized before expert stacking; no re-quantization after stacking |
| Plain FP8 weights | Unsloth DGX high class: fp8_e4m3 raw U8 weights + per-output BF16 scales; current runtime dequantizes once to BF16 and runs A16 matmul/gather-mm (not native W8A8) |
| FP8 KV cache | Paged-adapter only — KVCacheDType::Fp8 with per-layer scale management via KvScaleManager. FP8 KV is intentionally rejected by the flat-path attach. |
| ggml K-quants | --gguf-kquant on a GGUF import; modes q6k / q4k / q5k. Consume-only — the bridge's mlx_quantize throws; producing needs ggml's make_qkx2_quants |
| Symmetric ggml Q4_0 / Q8_0 | GGUF import, automatic. w = d*(q - Z) has no stored offset, so .biases is omitted and rebuilt from symmetric_zero_point in config.json |
K-quants are affine per sub-block, which is why they need no new kernel family: scale/bias are computed by a two-level decode instead of a scalar load, and qdot / qouter / dequantize are unchanged.
Q4_K/Q5_K y = d*sc[j]*q - dmin*m[j] -> scale = d*sc[j] bias = -dmin*m[j]
Q6_K y = d*sc[j]*(q-32) -> scale = d*sc[j] bias = -32*d*sc[j]
Sub-scales ride in the existing .scales / .biases sidecars as interleaved pairs, so FFI arity stays at two companion arrays. .biases holds ggml's d, a scale — not an additive bias, and its dtype must stay float16: resolve_kquant_group requires it exactly, and the converter's dtype passes skip it for that reason (kquant_biases_to_preserve).
Density, and where it differs from ggml:
| mode | mlx-node | ggml | why |
|---|---|---|---|
| q6k | 6.5625 bpw | 6.5625 | exact |
| q4k | 4.6250 bpw | 4.5000 | sub-scales stored unpacked, +0.125 |
| q5k | 5.6250 bpw | 5.5000 | same |
| Q4_0 | 4.5000 bpw | 4.5000 | exact, once the derived bias is dropped |
| Q8_0 | 8.5000 bpw | 8.5000 | exact, same reason |
ggml packs Q4_K/Q5_K sub-scales into 6 bits. Storing them unpacked costs 0.125 bpw but keeps the affine pointer-walk contract (scales += block_size/group_size) intact at ~20 call sites and keeps a divergent branch out of qmv's innermost loop.
MLX_PORTABLE_KQUANT controls a separate BF16/FP16 prefill path on Metal
without NAX. It keeps 16-bit matrix operands with FP32 accumulation and uses
a 64×64×32 SIMD tile (10 KiB threadgroup scratch). It does not dequantize or
retain a dense model copy. Automatic selection requires at least 128 input
rows, at least 1,024 output columns, contiguous operands, native q4k/q5k/q6k/iq4xs,
and enough tiles to avoid replacing stock split-K arithmetic. Pipeline limits
are checked on the actual device; unsupported compilation keeps the default
K-quant matmul (kquant::quantized_matmul).
Set 0 to roll back, or 1 to force the eligible path for qualification on
an M5. NAX, decode, short verification, training, and unsupported formats keep
their existing routes by default. See the portable investigation.
K-quant prefill takes mlx-node's own NAX kernel (kquant_qmm_t_nax from metal/kquant/kquant_nax.h, prebuilt into paged_attn.metallib) under the same gate as MLX's affine qmm: is_nax_available(), transposed weight, K % 64 == 0, and TF32 enabled or a non-F32 input (crates/mlx-sys/src/mlx_kquant_metal.cpp:596-598). Without the NAX kernel, K-quant qmm_t runs 3.40–3.49× the affine control; with it, it is level with affine (M=512 N=8192 K=8192, bf16, M5 Max; measured on the fork's kernel, which ours matches bit for bit). That is removal of a penalty, not a win over affine — an A/A affine control on the same harness lands 0.53% off unity, which is the noise floor those deltas sit near.
kquant_nax_bench only means anything on applegpu_g17s / macOS ≥ 26.2; elsewhere both arms fall back to the same simdgroup kernel. KQ_NAX_REQUIRE=1 turns its skip — and the NAX band inside kquant_mode_guards — into a hard failure.
crates/mlx-core/src/convert.rs supports:
- mlx-lm-style mixed-bit:
mixed_2_6,mixed_3_4,mixed_3_6,mixed_4_6 qwen3_5— Qwen3.5-tuned recipeunsloth— legacy affine requires imatrix calibration.--q-mxfpselects the fixed tensor-class map translated from NVFP4/FP8 to MXFP4/MXFP8;--q-mode nvfp4selects the DGX weight map with NVFP4/plain per-output E4M3 FP8. These fixed maps may omit the imatrix; AWQ pre-scaling is then skipped and quality may be lower, while the class map remains unchanged. Qwen uses its final-eight FFN split; exact Gemma-4-26B-A4B MoE checkpoints keep all dense/expert FFNs low, attention q/k/v/o high, and all other tensors BF16, with no final-eight or head exception. Flat GGUF imatrix statistics are not promoted into an invented expert-axis calibration rule. DGX runtime remains A16 for both weight classes and does not preserve Unsloth's calibrated global scales, W4A4/W8A8 activation execution, or calibrated FP8 KV-cache scales; no upstream numerical or performance parity is claimed. Plain affine alone keeps the legacy Dynamic 2.0 map.
AWQ-style imatrix pre-scaling is supported for improved low-bit quality.
quant_predicate defaults: router gates → 8-bit; everything else → 4-bit.