Skip to content

Add a fused full-attention path for head_dim 256 on NAX devices - #3842

Merged
zcbenz merged 6 commits into
ml-explore:mainfrom
wyanzhao:nax-256-dsplit
Aug 19, 2026
Merged

zcbenz merged 6 commits into
ml-explore:mainfrom
wyanzhao:nax-256-dsplit

Conversation

@wyanzhao

@wyanzhao wyanzhao commented Jul 13, 2026 •

Copy link
Copy Markdown
Contributor

Proposed changes

Add a fused NAX full-attention path for head_dim == 256 without
materializing the [n_heads, qL, kL] score tensor.

Area Change
Kernel configuration bq=64, bk=32, wm=4, wn=2
Parallel decomposition Separate attention_nax_dsplit kernel: split D and Dv across the WN = 2 simdgroups of the second warp dimension
Routing Causal, no array mask, qL >= 1024, NAX-compatible dtype/TF32 mode
Ragged inputs No padding copies. Ragged qL runs the kernel's unaligned-Q path (Q is loaded once, outside the KV loop, so it is free). Ragged kL from a KV-cache slice is read past its end up to the next bk multiple (the extra columns are causally masked); ragged kL from a plain array runs the unaligned-K path (~13% slower over the KV loop, still ~2x the unfused graph)
Fallback safety Preserve fallback for unsupported shapes and float32 with TF32 disabled

Motivation

use_fallback currently allows fused full attention only for
head_dim in {64, 80, 128}. Full-attention layers in the Qwen3.5/3.6
family use head_dim == 256, so they fall back to an unfused graph that:

  • materializes the full score tensor;
  • computes the full rectangle before applying the causal mask; and
  • becomes increasingly expensive at long context lengths.

A direct bd=256 instantiation of steel_attention_nax reaches only about
half the efficiency of bd=128 on M5 Max (22 vs 48 TF). The limiting
resource is the per-simdgroup accumulator byte budget: the bd=256 output
and score accumulator set is twice as large. A half-precision-accumulator
probe recovered +51%, while register pressure, scheduling fences,
accumulation-chain order, instruction footprint, and Q reload traffic were
ruled out by measurement.

Design

Head-dimension split

attention_nax_dsplit, a separate kernel from attention_nax dispatched
for bd = 256 with wn = 2, assigns each simdgroup one half of D and
Dv:

  1. Each half computes its partial Q @ K.T scores.
  2. The two halves exchange score fragments pair-by-pair through a 16 KB
    threadgroup buffer.
  3. Both halves run softmax on the complete score tile.
  4. Each half accumulates P @ V for its owned half of Dv.

Both halves use the same fragment-to-lane mapping, so the exchange can use a
per-lane-linear layout without coordinate conversion. Q fragments remain
register-resident across the KV loop.

This reduces the per-simdgroup accumulator set to 2 S + 8 O fragments
(320 B/lane), matching the working set of the existing bd=128 kernel.

Ragged causal inputs

There is no padding. Ragged query lengths take the kernel's unaligned-Q
path (align_Q = false), which measures within noise of the aligned one:
Q is loaded into registers once, before the KV loop, so raggedness never
touches the hot loop.

Ragged key/value lengths matter more: the unaligned-K pipeline
(align_K = false) carries bounds-checked K/V loads and column masking
inside the KV loop and measures ~13% slower over the whole loop on M5 Max.
The dispatch avoids it for the common case without copying anything: K/V
that come out of a preallocated KV cache (mlx-lm's KVCache grows its
buffer 256 positions at a time) are slices of a longer buffer,
keys[..., :kL, :]. When such a slice can be read past kL up to the next
bk multiple — the head's row pitch (strides[1] / strides[2]) reaches
that far and the buffer bound allows it — the kernel is simply told the
widened kL. This is safe under causal masking with qL_off computed from
the true lengths: every extra key column lies strictly beyond the
diagonal of every real query row, so it receives exactly zero probability
and its (finite) V row never reaches the output.

The check deliberately differs from the CUDA use_cudnn_for_decoding
heuristic in two ways: Metal buffers are page-rounded, so it tests
offset + extent <= buffer_size instead of equality, and it does not
require the cache capacity to be a multiple of 256 — after a chunked
prefill mlx-lm's cache is e.g. 2049 + 1280 = 3329 positions long, and
only the row pitch reaching the widened length matters. K/V that are not
such slices (a plain array with a ragged kL) run the unaligned-K path.

Routing and fallback behavior

Condition Behavior
NAX, causal, no array mask, head_dim == value_head_dim == 256, qL >= 1024 Fused wn=2 path
qL < 1024 Existing fallback
Array mask present Existing fallback
float32 with MLX_ENABLE_TF32=0 Existing fallback
Non-NAX device Existing fallback

The qL >= 1024 threshold is based on measured M5 Max behavior. Below that
point, long-cache rectangles do not provide enough query blocks to fill the
machine (512 x 8192 measured approximately 1.08x the unfused time on an
earlier variant).

Performance

Unless noted otherwise: M5 Max, 16 Q / 2 KV heads, bf16, causal.

Kernel benchmarks

Shape Before After Improvement
8192 x 8192 24.92 ms / 22.1 TF 12.84 ms / 42.8 TF 1.94x throughput
4096 x 4096 8.87 ms 3.45 ms 2.57x faster
2048 x 8192 8.92 ms 5.77 ms 1.55x faster
1024 x 8192 4.65 ms 3.58 ms 1.30x faster
2048 x 2048 2.39 ms 1.19 ms 2.01x faster

At 8192 x 8192, the new kernel reaches approximately 90% of the
bd=128 kernel's throughput.

Ragged lengths

M5 Max, 16 Q / 2 KV heads, bf16, causal, min of alternating A/B runs:

Shape (qL x kL) Path Time
4096 x 4096 aligned reference 3.46 ms
4064 x 4096 ragged Q, unaligned-Q path 3.61 ms
4096 x 4112 ragged K sliced from a longer cache (widened, aligned pipeline) 3.53 ms
4096 x 4112 ragged K from a plain array (unaligned-K path) 4.02 ms
8192 x 8200 ragged K, aligned vs unaligned-K path 15.2 ms vs 17.2 ms

Padding Q instead would have cost more than it saved (3.84 ms for
4064 x 4096), which is why the earlier padding copies were dropped.

End-to-end prefill

Qwen3.6-35B-A3B 4-bit, ragged prompts:

Prompt / chunking Before After Change Peak memory
8k, default 2048-token chunks 4039 tok/s 4193 tok/s +3.8% —
8k, single chunk 4209 tok/s 4822 tok/s +14.6% —
16k, 8k-token chunks 3457 tok/s 4180 tok/s +20.9% 29.9 -> 26.1 GB
32k, 8k-token chunks 2833 tok/s 3598 tok/s +27.0% 34.6 -> 26.5 GB
32k, single chunk 1771 tok/s 2873 tok/s +62.2% 72.9 -> 41.8 GB

Transient memory for the 4096 x 4096 kernel case drops from 663 MB to
75 MB because the fused path no longer materializes the score tensor.

The end-to-end numbers were taken with an earlier revision that padded
ragged inputs with copies; mlx-lm hands over cache slices, which now run
the same aligned kernel without the pad/copy work.

Numerical behavior

Against an f32 reference, the fused bf16 output is more accurate than the
previous unfused graph:

Implementation Relative L2 error
Unfused bf16 graph 2.3e-3
Fused NAX kernel 1.6e-3

Validation

  • Fallback and fused shapes checked against the primitives reference in
    float32 and bfloat16.
  • Committed coverage includes aligned/ragged squares and rectangles, plus
    ragged K/V sliced out of a longer preallocated cache with large, finite
    stale rows behind the slice (any leak into the output would be loud).
  • Verified directly against mlx_lm.models.cache.KVCache slices from a
    chunked prefill (2049 then 1031 tokens) in bf16 and fp32/TF32.
  • Additional local boundary checks cover odd lengths 1031 and 2049, plus
    the exact qL == 1024 routing boundary.
  • A subprocess test with MLX_ENABLE_TF32=0 verifies that
    float32 + head_dim 256 remains on the fallback path.
  • Full local test_fast_sdpa.py result: 20 tests, 1 platform-dependent
    skip, otherwise passing
    .

Checklist

  • I have read the CONTRIBUTING document.
  • I have run pre-commit run --all-files.
  • I have added tests that cover the new fused and fallback behavior.
  • No documentation changes are required for this internal kernel path.

@wyanzhao
wyanzhao marked this pull request as ready for review July 13, 2026 20:54
@pierre427

Copy link
Copy Markdown
Contributor

Validated on M5 Max (mlx 0.32.1, NAX active). Built the merge-base
(4367c73b) and PR head (1aee415d) from source with the
MACOSX_DEPLOYMENT_TARGET=26.2 NAX gate; bf16 4096³ GEMM canary confirms NAX in
both. D=256 causal SDPA A/B, GQA 32/8 and 64/8, S∈{16384,32768}, L∈{1..2048},
warmup + 40-iter median, run in both orderings (base-first and head-first) to
control for ambient throttling.

  • Win on the large-L prefill shapes the kernel targets — L≥1024 at both GQA
    ratios and both S: head faster in both orderings (≈1.2–1.5×; the fused path
    avoids materializing the [n_heads, qL, kL] score tensor, so the gain grows
    with qL). These are exactly the Qwen3.5/3.6 full-attention (head_dim=256) shapes.
  • No regressions on any shape in either ordering; decode (L=1) and mid-L
    (256/512) are within run-to-run noise.
  • Correctness fp16-clean vs an fp32 reference (max_abs_err ~1e-4).

Caveat: this run was on a machine with intermittent GPU contention (canary drifted
2.5→4.4 ms across passes), so treat the magnitude as a direction-robust ~1.2–1.5×
rather than a precise figure — a clean-machine re-run would tighten it.

@zcbenz zcbenz added the await verification This pull request is non-trivial and requires a human expert to verify its correctness. label Aug 5, 2026
Comment thread mlx/backend/metal/kernels/steel/attn/kernels/steel_attention_nax.h Outdated
@zcbenz

zcbenz commented Aug 15, 2026

Copy link
Copy Markdown
Member

@jessegross Would you mind testing how much this is helping in production?

Comment thread mlx/backend/metal/kernels/steel/attn/kernels/steel_attention_nax.h
enc.add_temporary(*o_pad);
qL = qLp;
}
if (kL % bk) {

@zcbenz zcbenz Aug 16, 2026 •

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The KV cache written for MLX is usually guaranteed to be padded to 256 otherwise it would be too slow:
https://github.com/ml-explore/mlx/blob/main/docs/src/usage/kv_cache.rst

And we make use of this fact when using cuDNN SDPA for decoding:

// Check if k/v are slices from fixed-size kv cache.
auto is_slice = [](const array& kv) {
// Get pre-sliced sequence length from strides, and check if the buffer
// belongs to a contiguous kv cache.
int64_t T_kv = kv.strides(1) / kv.strides(2);
if (kv.size() / kv.shape(2) * T_kv != kv.buffer_size() / kv.itemsize()) {
return false;
}
// It is possible to use heuristic to check slices, but for now just make
// mlx-lm work.
return T_kv % kv_cache_step == 0;
};

So I think there is no need to pad K/V here, we just need to add heuristics to know it is already padded. The original kernel can also make use of this fact and we can probably do it in another PR.

There is probably no longer need to pad queries as well since unaligned KV might be responsible for most of the slowdown.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done — the pad copies for Q, K and V are gone.

K/V: when they are slices of a longer buffer (keys[..., :kL, :] from a preallocated cache), the dispatch just tells the kernel a kL widened to the next bk multiple. With qL_off computed from the true lengths, every extra key column lies strictly beyond the causal diagonal of every real query row, so it gets exactly zero probability and its (finite) V row never reaches the output — same assumption the cuDNN decoding path makes about the rows behind a cache slice. Two deliberate differences from use_cudnn_for_decoding: Metal buffers are page-rounded, so the check is offset + extent <= buffer_size rather than equality; and I don't require the cache capacity to be a multiple of 256, only that the head's row pitch (strides[1] / strides[2]) reaches the widened length — after a chunked prefill mlx-lm's KVCache is e.g. 2049 + 1280 = 3329 positions long. Verified directly against mlx_lm.models.cache.KVCache slices (bf16 and fp32/TF32); the tests now include cache slices with large stale rows behind them.
Q: you were right — the unaligned-Q path is free (Q is loaded once, outside the KV loop; 4064x4096 measures 3.61 ms vs 3.46 ms aligned, while padding Q cost 3.84 ms), so Q is left ragged.

Kept this to the split kernel; the base kernel can pick it up in a follow-up as you suggested.

Comment thread mlx/backend/metal/kernels/steel/attn/kernels/steel_attention_nax.h Outdated
@zcbenz zcbenz removed the await verification This pull request is non-trivial and requires a human expert to verify its correctness. label Aug 16, 2026
@jessegross

Copy link
Copy Markdown
Contributor

@zcbenz I can take a look tomorrow but in general I think it is worth fusing head dim 256. It would be nice to have a version for pre M5 devices as well though :)

wyanzhao and others added 5 commits August 18, 2026 17:10
use_fallback currently allows the fused full-attention path only for
head_dim in {64, 80, 128}. Models with head_dim 256 (e.g. the Qwen3.5/3.6
family's full-attention layers) fall back to the unfused graph, which
materializes the [n_heads, qL, kL] score tensor and computes the full
rectangle before masking.

A naive bd=256 instantiation of steel_attention_nax runs at half the
efficiency of the bd=128 kernel on every shape (22 vs 48 TF on M5 Max):
the per-simdgroup accumulator working set doubles (16 + 4 f32 output and
score fragments), and accumulator bytes are the resource that gates
tensor-unit throughput at this width (a probe with fp16 accumulators
recovers +51%; static register pressure, scheduling fences, accumulation
chain order, instruction footprint and Q reload traffic were each ruled
out by measurement).

This adds a head-dim-split configuration instead (bq=64, bk=32, wm=4,
wn=2): the second warp dimension splits D, so each simdgroup computes
Q@K.T partials over one half of D (exchanged pair-by-pair through a 16KB
threadgroup buffer in per-lane-linear layout, which needs no coordinate
math because both halves share the fragment-to-lane map), runs softmax
redundantly on the full score tile, and accumulates P@V for its half of
Dv. That puts the accumulator set (2 S + 8 O fragments, 320B/lane) at
exact bd=128 parity, and throughput follows: 42.8 TF at 8192^2 causal,
90% of the bd=128 kernel. Q fragments stay register-resident across the
KV loop, and mma chains keep a single accumulator (alternating
accumulators measures -25%).

Routing and raggedness (all numbers M5 Max, 16 Q / 2 KV heads, bf16,
causal):

- The fused path is taken for causal shapes with qL >= 1024 and no array
  mask, where it beats the unfused graph everywhere measured, offset
  rectangles included: 2048x8192 5.77 vs 8.92 ms, 4096^2 3.45 vs
  8.87 ms, 1024x8192 3.58 vs 4.65 ms, 2048^2 1.19 vs 2.39 ms. Below
  ~1k queries against a long cache there are too few query blocks to
  fill the machine (512x8192 measured 1.08x on an earlier variant), so
  those keep the fallback.
- Ragged lengths are padded up to the tile sizes at dispatch and the
  padded query rows sliced off on the way out: the unaligned
  (align_Q/align_K = false) pipelines slow every block down ~3x at this
  width, and real decoder prefills are essentially never tile-aligned.
  With qL_off computed from the true lengths the causal bound
  c <= r + qL_off never reaches a padded key column; pads are
  zero-filled so padded value rows cannot poison P @ V; cost is ~0.6 ms
  against a ~49 ms cliff.
- float32 with tf32 disabled keeps the fallback, since only the NAX
  kernel is instantiated at this width.

Transient memory at 4096^2 drops 663 MB -> 75 MB (no score
materialization), and the fused output is closer to an f32 reference
than the unfused bf16 graph (relative L2 1.6e-3 vs 2.3e-3).

End to end (Qwen3.6-35B-A3B 4-bit, 8k ragged prompt):
  default 2048-token chunked prefill: 4039 -> 4193 tok/s (+3.8%)
  single-chunk prefill:               4209 -> 4822 tok/s (+14.6%)

Tests: fallback shapes and fused shapes (aligned and ragged squares, a
boundary rectangle and a ragged rectangle) against the primitives
reference in float32 and bfloat16, plus a subprocess case with
MLX_ENABLE_TF32=0 covering the float32 fallback predicate (a mismatch
there fails as a kernel-load error).
… end

- sdpa_full_self_attention_nax no longer pads Q/K/V. Ragged Q takes the
  kernel's unaligned-Q path (measured free: Q is loaded once, outside the
  KV loop). K/V that are slices of a longer buffer (mlx-lm KVCache) are read
  past kL up to the next bk multiple; with qL_off from the true lengths the
  extra columns are causally masked to zero probability. Other ragged K/V
  take the unaligned-K path (~13% over the KV loop on M5 Max).
- attention_nax_dsplit: express the head-dim split constants via WN.
- Tests: ragged K/V sliced out of a longer cache with stale rows behind it.

@zcbenz zcbenz left a comment

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Did some simple inference with Qwen3.8 and I can observe consistent improvement on long prompt prefill, nice improvement!

WindChimeRan added a commit to vllm-project/vllm-metal that referenced this pull request Aug 19, 2026
<img width="2496" height="884" alt="bench_nax_prefill"
src="https://github.com/user-attachments/assets/d482a22b-2e37-435c-8fbc-a7097deb2a6d"
/>

## Summary

Adds an automatic, opt-out paged-attention prefill kernel for the NAX
tensor units on M5 GPUs.

- Supports fp16/bf16 MHA/GQA/MQA with head sizes 64, 96, 128, 256, or
512 and block sizes 8, 16, or 32.
- Pure decode, speculative verification windows, TurboQuant, unsupported
hardware/shapes, and missing NAX libraries retain their existing paths.
GDN linear attention is unchanged.
- Mixed batches use NAX for the whole prefill dispatch, including
`q_len=1` decode passengers.
- Set `VLLM_METAL_DISABLE_NAX=1` to force the existing tiled kernel.

The pure-decode kernel is unchanged. End-to-end TPOT can still improve
during chunked prefill because prefill and decode share engine steps.

## Future work

* maybe GDN support
* although this PR's head_dim 256 512 is better than main, it's still
under-utilize the NAX. We could refer to upstream PR for reference on
future work ml-explore/mlx#3842

## Performance

Measured on an M5 Pro with macOS 26.6 and MLX 0.32.0:

| Workload | Metric | This PR vs main |
|---|---|---:|
| Qwen3-8B-bf16, 8 x 4096-token prompts | Mean TTFT | **-21%** |
| Qwen3-8B-bf16, 6 x 7900-token prompts | Mean TTFT | **-23%** |
| Qwen3-0.6B, sonnet 2048 in / 32 out, concurrency 32 | Total tok/s |
**+33%** |
| Same | Mean TTFT | **-41%** |
| Gemma 4 E4B bf16, sonnet 2048 / 32, concurrency 32 | Total tok/s |
**+27%** |
| Same | Mean TTFT | **-24%** |
| Qwen3.8-27B 8-bit | Serving throughput / TTFT | no measurable change |

Qwen3.8's head-size-256 attention kernel is about 2.1x faster, but only
16 of 64 layers use SDPA and 8-bit weight dequantization dominates the
serving step.

## Validation

All 60 M5 specialization pipelines pass parity. Targeted MHA/MQA, fully
masked rows, dual-arm golden generation, AOT packaging, affected tests,
and GitHub CI also pass.

<details>
<summary>Kernel and serving measurements</summary>

### Paged bf16 attention speedup

| Head size | 2K | 4K | 8K |
|---:|---:|---:|---:|
| 96 | 2.77x | 3.14x | |
| 128 | 2.98x | 3.63x | 3.36x |
| 256 | 2.16x | 2.02x | 1.85x |
| 512 | 4.36x | 5.69x | |

At actual head-size-256 GQA geometry, the 4K kernel speedup is 2.09x for
Qwen3.5 (8 query / 2 KV heads) and 2.06x for Qwen3.8 (24 / 4).

Additional serving results:

- Qwen3-8B, 24 x 4096-token prompts at concurrency 4: +57% total tok/s
and -38% mean TTFT.
- Gemma 4 E4B: -22% duration and -22% mean TPOT.
- Qwen3-0.6B 2048/32: -25% mean TPOT. Decode itself does not use NAX;
faster chunked prefill reduces interference with decodes in the same
engine step.

Serving A/B runs use one build and toggle only `VLLM_METAL_DISABLE_NAX`.

</details>

<details>
<summary>Implementation, MLX provenance, and dataflow</summary>

The 449-line shader adapts MLX 0.32's `steel_attention_nax`: BQ=64,
BK=32, four simdgroups, 16x32x16 MPP `matmul2d`, register-fragment Q/K/V
loads, online softmax, relaxed precision, and the head-size-128
scheduling barrier. vLLM extends the schedule to head sizes 96, 256, and
512, using a midpoint PV-loop split for 256/512.

| Portion | Size | Implementation |
|---|---:|---|
| NAX fragment helpers | 55 non-comment lines | MLX lane mapping,
cooperative-tensor packing, and row reduction |
| Paged/varlen attention body | 259 non-comment lines | vLLM addressing,
masking, and online-softmax integration |
| Specializations | 43 non-comment lines | 30 kernels: two dtypes x five
head sizes x three block sizes |

The vLLM body adds block-table gathers, packed varlen queries,
chunked-prefix causal positions, partial-row/page guards, sliding
windows, softcap, sinks, and GQA head mapping.

The dataflow deliberately differs from the classic tiled kernel. Tiled
stages Q/K/V in threadgroup memory and reuses K/V across four
simdgroups; NAX follows MLX and loads register fragments directly,
relying on GPU caches. Controlled ablations found that staging NAX K/V
was 34% to 58% slower, while loading tiled K/V directly was 32% to 63%
slower. Direct loading is therefore specific to the NAX fragment
geometry, not a universal optimization.

The relaxed MPP descriptor changes accumulation order and truncates the
fp32 P operand of PV, so validation uses bounded parity rather than
bitwise identity.

</details>

<details>
<summary>Compatibility, correctness, and release packaging</summary>

NAX requires M5 hardware, macOS 26.2 or later, a successfully loaded NAX
metallib, fp16/bf16, a supported head/block size, and a
prefill-containing batch outside verification-window and TurboQuant
modes. Initialization failures warn and fall back.
`VLLM_METAL_DISABLE_NAX=1` is the emergency override.

Correctness coverage:

- 60 GQA pipelines: two dtypes x five head sizes x three block sizes x
fast/masked regimes. Worst observed divergence remains below 0.9 ULP
max.
- Targeted MHA and MQA fast/masked cases pass.
- Fully masked rows remain finite and match tiled attention.
- Qwen3-0.6B greedy generation produces identical golden tokens with NAX
enabled and disabled.
- Prebuilt and source/JIT initialization paths pass.

Official macOS 15 CI and release jobs select Xcode 26.3. Required native
artifacts target macOS 15.0; the separate NAX metallib targets macOS
26.2. Wheel verification requires the NAX metallib and checks both
native extensions' deployment target.

GitHub's M1 runner compiles and packages NAX. The 60-pipeline execution
matrix remains a manual M5 release check.

</details>

---------

Signed-off-by: Ranran <ranranhaoranzhang@gmail.com>
@pierre427

Copy link
Copy Markdown
Contributor

Re-validated on M5 Max at the current head. My earlier numbers (2026-07-16) were
against 1aee415d and are superseded by the kernel split and the padding-copy
removal since.

Built PR base db935c094 and head b86d328d6 from source with the
MACOSX_DEPLOYMENT_TARGET=26.2 NAX gate. D=256 causal SDPA, fp16, GQA 32/8 and
64/8, S ∈ {16384, 32768}, L ∈ {1, 8, 256, 512, 1024, 2048}, warmup + 40-iter
median, run in both orderings (base-first and head-first). bf16 4096³ GEMM canary
2.57–2.83 ms across the four passes confirms NAX active.

Noise floor first: same-arm drift between the two passes reaches 28.2% (base) and
23.6% (head) relative to the faster pass. Ratios below roughly 1.1x are not
separable from run-to-run variance on this machine.

Faster in both orderings at ≥1.03x — 8 of 24 shapes:

GQA S L speedup (min of the two orderings)
32/8 16384 1024 1.03x (marginal — inside the drift band)
32/8 16384 2048 1.18x
32/8 32768 1024 1.20x
32/8 32768 2048 1.33x
64/8 16384 1024 1.21x
64/8 16384 2048 1.46x
64/8 32768 1024 1.36x
64/8 32768 2048 1.55x

Under a strict >1.0-in-both-orderings criterion the count is 11 of 24; the extra
three (64/8·16k·L=512 = 1.006, 64/8·32k·L=8 = 1.005, 64/8·32k·L=256 = 1.005) are
well inside the noise floor and I would not claim them.

The robust result is L ≥ 2048, where the four wins (1.18x–1.55x) sit outside the
drift band and grow with both qL and S — consistent with avoiding the
[n_heads, qL, kL] score materialization. The L=1024 wins are directionally
consistent, but only the 64/8 rows clear the noise comfortably.

Slower in both orderings — two shapes:

GQA S L ordering A ordering B
32/8 16384 256 −7% −0.4%
32/8 32768 1 −2% −1.9%

Both are inside the drift band and the −7% does not reproduce in the second
ordering, so I read them as noise rather than real regressions — but they are not
zero.

Elsewhere: mid-L (256/512) spans 0.86–1.13x across orderings with no consistent
direction; decode-ish (L=1/8) spans 0.85–1.24x. Both bands are dominated by drift.

Correctness: fp16 vs an fp32 reference, max_abs_err 2.4e-4 – 4.9e-4 across
GQA 32/8 and 64/8 at L=64 and L=256.

Happy to post the raw per-shape logs for all four passes if that is useful for
recomputing the table independently.

@jessegross

Copy link
Copy Markdown
Contributor

This is what I see end-to-end on an M5 Max, which are consistent with @wyanzhao's numbers. The longer prompts are are chunked into batches of 2048.

That looks like a pretty worthwhile improvement to me.

qwen3.6:35b (MoE)

prompt tokens prefill (base) prefill (PR) speedup peak memory
2k 4582 tok/s 4703 tok/s +2.6% —
4k 4604 tok/s 4791 tok/s +4.1% −0.1 GiB
8k 3896 tok/s 4045 tok/s +3.8% −0.3 GiB
16k 3209 tok/s 3491 tok/s +8.8% −0.8 GiB

qwen3.8:27b (dense)

prompt tokens prefill (base) prefill (PR) speedup peak memory
2k 621 tok/s 633 tok/s +1.9% —
4k 621 tok/s 637 tok/s +2.5% −0.2 GiB
8k 613 tok/s 629 tok/s +2.5% −0.5 GiB
16k 593 tok/s 617 tok/s +4.1% −1.2 GiB

@zcbenz

zcbenz commented Aug 19, 2026

Copy link
Copy Markdown
Member

Thank you all for helping testing this PR! 🙏

@zcbenz
zcbenz merged commit 714a7ef into ml-explore:main Aug 19, 2026
davidtai added a commit to Layr-Labs/mlx that referenced this pull request Aug 25, 2026
* Return tuple in meshgrid (ml-explore#4229)

* Add endpoint parameter to linspace (ml-explore#4184)

Co-authored-by: Cheng <git@zcbenz.com>

* Fix vmap of partition/argpartition dropping the kth argument (ml-explore#4116)

* Fix nan_to_num replacing inf with 0 for float16 and bfloat16 (ml-explore#4222)

Co-authored-by: codeAnqiang-ma <273298913+codeAnqiang-ma@users.noreply.github.com>
Co-authored-by: Cheng <git@zcbenz.com>

* Fix einsum not broadcasting batch dimensions in batched tensordot (ml-explore#4125)

Co-authored-by: Cheng <git@zcbenz.com>

* Dequantize in float32 (ml-explore#4241)

* chore: Reject complex in erf and erfinv (ml-explore#4243)

* Fix cpu compilation failure of abs with uint (ml-explore#4240)

Co-authored-by: Cheng <git@zcbenz.com>

* Fix quantize matrix multiplication floor issue (ml-explore#4251)

* Only use MPI backend for world size > 1 (ml-explore#4210)

* chore: Reject complex in expm1, sigmoid and arctan2 (ml-explore#4257)

* Decompose small kernel-depth 3D convs into 2D convs (ml-explore#3785)

Co-authored-by: katlun-lgtm <264247399+katlun-lgtm@users.noreply.github.com>
Co-authored-by: Cheng <git@zcbenz.com>

* Fix Metal sort of a view with a negative stride (ml-explore#4252)

* Mirror the depth axis in the decomposed 3D conv when flipped (ml-explore#4277)

* Fix Metal row reductions on negative-stride views (ml-explore#4267)

Co-authored-by: Fu Xiaonan <214359569+FU-max-boop@users.noreply.github.com>

* [CUDA] Fix custom kernel cache collision for same name, different source (ml-explore#4273)

Co-authored-by: Cheng <git@zcbenz.com>

* Fix ops rejecting integers larger than INT32_MAX (ml-explore#4255)

Co-authored-by: Feli <feli@hnu.edu.cn>
Co-authored-by: Cheng <git@zcbenz.com>

* Fix var/std for complex numbers (ml-explore#4260)

* Fix int32 overflow in conv padded input and pad shapes (ml-explore#4258)

Co-authored-by: Cheng <git@zcbenz.com>

* chore: Reject complex in remainder (ml-explore#4270)

* chore: Compare the macOS SDK version as a version when gating JACCL (ml-explore#4286)

* Clamp ring socket transfers so a payload of 2 GiB or more can be sent (ml-explore#4281)

Co-authored-by: Cheng <git@zcbenz.com>

* chore: Use normalize_axis_index in split/unstack/partition/topk (ml-explore#4288)

* Remove grouped output in CI (ml-explore#4195)

* [CUDA] Fix finding cuda 13 headers in JIT compilation (ml-explore#3995)

* Refactor wheel building script (ml-explore#3818)

* Make mx.compile cache erasing thread safe (ml-explore#4248)

Co-authored-by: yentur <mr.yentur@gmail.com>

* Add builds for free-threaded python (ml-explore#3812)

* Fix int32 overflow in concatenate/repeat/kron (ml-explore#4303)

* python: Widen list elements that do not fit in int32 to int64 (ml-explore#4305)

* Propagate CPU errors to events (ml-explore#3742)

Co-authored-by: Alessio Pollero <alessio.pollero@gmail.com>

* Fix mx.arange dtype inference overflow regression (ml-explore#4324)

* Add workflow to update pull request limit bypass list (ml-explore#4320)

* Support head dimension 72 in Metal full attention (ml-explore#4330)

* Patch bump to 0.32.2 (ml-explore#4333)

* Preserve subnormal float values when casting to bool (ml-explore#4224)

* python: Support assigning through a bare Ellipsis index (ml-explore#4314)

* Fix divmod truncating the quotient for floats (ml-explore#4108)

Co-authored-by: Cheng <git@zcbenz.com>

* Add force_fused option to scaled_dot_product_attention (ml-explore#4185)

* chore: Reject negative eps in the normalization layers (ml-explore#4312)

* Bound GGUF metadata string/array values against the file mapping (ml-explore#4212)

Co-authored-by: x14ngch3n <x14ngch3n@users.noreply.github.com>
Co-authored-by: Cheng <git@zcbenz.com>

* Read each K/V byte once in gqa-8 decode attention (ml-explore#4077)

* Fix fft vmap and jvp for transforms over a subset of axes (ml-explore#4138)

* Fix median dropping NaN (ml-explore#4146)

* Fix the CPU scan over a size one axis with a padded stride (ml-explore#4139)

Co-authored-by: Cheng <git@zcbenz.com>

* chore: Validate the optimizer betas at construction (ml-explore#4310)

Co-authored-by: Cheng <git@zcbenz.com>

* `RMSNormVJP` backward writes a full `{n_rows, D}` `gw_temp` intermediate (ml-explore#4293)

* [Bug]: add default none value to axis parameter of the take_along_axis (ml-explore#4357)

Co-authored-by: Anastasiia Filippova <a_filippova@apple.com>

* Add a fused full-attention path for head_dim 256 on NAX devices (ml-explore#3842)

Co-authored-by: Cheng <git@zcbenz.com>

* Update nanobind to 2.15.0 (ml-explore#4337)

* Skip unnecessary simdgroup computations for quantised MOE matmuls on NAX (ml-explore#4352)

* Add AI usage policy (ml-explore#4331)

Co-authored-by: Jake Bowhay <60778417+j-bowhay@users.noreply.github.com>

* Raise cpu stream errors from synchronize (ml-explore#4338)

Co-authored-by: Cheng <git@zcbenz.com>

* chore: Validate eps in Adam at construction (ml-explore#4361)

Co-authored-by: Anastasiia Filippova <a_filippova@apple.com>

* Bound winograd conv2d working set by tiling the batch (ml-explore#4102)

Co-authored-by: Cheng <git@zcbenz.com>

* Use a 32-row block in qmm_t_nax when one block covers all of M (ml-explore#4171)

* chore: Deduplicate fftshift and ifftshift (ml-explore#4318)

* Fix Log and Equal is_equivalent ignoring primitive state (ml-explore#4266)

Co-authored-by: Cheng <git@zcbenz.com>

* Stabilize reduced-precision InstanceNorm (ml-explore#4230)

* chore: Normalize negative axes in sort and argsort (ml-explore#4332)

* Clean up main thread compile cache before python interpreter shuts down (ml-explore#4373)

* chore: Check malformed jaccl hostfile that miss rdma in pairs (ml-explore#4284)

Co-authored-by: Cheng <git@zcbenz.com>

* Round mxfp8 block scales up to avoid saturation (ml-explore#4353)

Co-authored-by: Daniel Hiltgen <daniel.hiltgen@ollama.com>
Co-authored-by: Cheng <git@zcbenz.com>

* Add support for the __array_namespace_info__  (ml-explore#4334)

* Stop a failed CUDA graph commit from poisoning the encoder (ml-explore#4356)

Co-authored-by: Cheng <git@zcbenz.com>

* Fix quantized kernels in JIT build (ml-explore#4372)

Co-authored-by: Cheng <git@zcbenz.com>

* Avoid zero work in stride-2 ConvTranspose3d (ml-explore#4343)

* [CUDA] Ce fused kernel (ml-explore#3947)

* Fix cpu exclusive scan for complex numbers (ml-explore#4272)

Co-authored-by: Cheng <git@zcbenz.com>

* Support Relocatable CUDA DLLs on Windows (ml-explore#4382)

* Use cast_to for fused AsType in compiled Metal kernels (ml-explore#4351)

Co-authored-by: katlun-lgtm <katlun@windyviews.com>
Co-authored-by: Cheng <zcbenz@gmail.com>

* python: Declare DLPackCompatible protocol members as methods (ml-explore#4384)

* Fix quantizing sliced arrays (ml-explore#4381)

* Fix einsum dropping a trailing empty subscript (ml-explore#4299)

Co-authored-by: Cheng <git@zcbenz.com>

* Add script to run python tests (ml-explore#4393)

* Hold GIL in AttachedData destructor (ml-explore#4391)

* Bound Metal buffer COUNT, not just bytes, in MetalAllocator

The Metal allocator throws `[metal::malloc] Resource limit (N) exceeded`
when num_resources_ (the live+cached Metal buffer COUNT) reaches
resource_limit_ (the iogpu.rsrc_limit sysctl, default ~499000). Freed
buffers are recycled into a size-keyed cache whose only trim is by BYTES
(release_cached_buffers takes a bytes-to-free target, max_pool_size_ ~=
physical RAM). Under churn with many distinct buffer shapes (varied prompt
lengths, growing KV caches, multiple co-resident models) the cache fills
with entries never reused at that exact size, so the COUNT climbs to the
limit while byte usage stays modest and the byte trim never fires — the
process crashes mid-inference on a machine with most of its RAM free.

malloc() now also reclaims by count: when num_resources_ crosses a 90%
high-water mark of resource_limit_, it clears the (pure-reuse) buffer
cache so the count drops back to the live working set. Clearing the cache
only costs re-allocation, never correctness, so the count limit becomes
unreachable by any request mix or batching method while the existing byte
limits keep total memory bounded.

Adds get_num_resources()/get_resource_limit() to the public memory API
(metal + no_gpu + cuda backends) so the count and its ceiling are
observable from callers. Adds an MLX_RESOURCE_LIMIT env override that can
only LOWER the ceiling (clamped to the OS limit, strictly validated) to
exercise the trim deterministically and as an operator safety valve.

* perf(mlx): opt-in Gemma 4 expert-QMM tile kernel with parallel descriptor builder (#4)

* perf(mlx): add opt-in Gemma 4 expert-QMM tile kernel with parallel descriptor builder

Adds a distinctly-named expert QMM implementation for the Gemma 4
26B-A4B MoE production shapes, gated by MLX_GATHER_QMM_EXPERT_SLICES:

- qmm_t_expert_impl: BM32 expert tile body (BM16 fallback rows) taking a
  private/by-value row count; the shared qmm_t_impl constant-address ABI
  and all ordinary gathered/batched/dense QMM routes are unchanged.
- build_gemma4_sorted_expert_tiles_bm32: one 128-thread threadgroup
  replaces the reference design's single-GPU-thread serial builder;
  parallel expert-range binary search, Hillis-Steele scan, and strided
  upper-bound descriptor emission.
- Selector runs after the NAX-first route and requires affine BF16
  transposed inputs, 4-bit gs=64 weights, 128 experts, assignment counts
  of exactly 4096/8192/16384, and the exact gate/up or down rank-3
  shapes; every miss keeps the legacy route. NAX engagement is
  non-engagement, never bypassed.
- device.{h,cpp}: one-shot request resolution, nonthrowing dual-symbol
  AOT probe/prewarm, relaxed-atomic diagnostics (requested, aotAvailable,
  naxAvailable, hits, per-class fallbacks).
- gpu_tests: exact-shape arithmetic parity, fallback, and counter
  invariant probes.

Retention standing (2026-08-09 production matrix): opt-in experiment.
Standalone profile dropped (prefill -10.2% vs bracket); paired
weighted-unsort+R1 profile retained-final (prefill +1.8%, TTFT -7.5%,
decode +3.3%, arrival E2E +12.0%). NOTE: this source post-dates the
benchmarked binaries/metallib (post-measurement kernel-body edit);
rebuild and re-verify before any performance claim.

* fix(mlx): fail-safe sortedness check in gemma expert tile builder; counter/atomic hygiene

Review-wave fixes for the R1 expert-QMM path:

- N1 (sortedness trust): build_gemma4_sorted_expert_tiles_bm32 now
  verifies each thread's post-binary-search segment boundary against the
  generalized invariant indices[start - 1] < lid <= indices[start]
  (edge threads check their single neighbor), votes per simdgroup via
  simd_or, folds the votes through threadgroup memory, and on any
  violation retracts count[0] to 0 (tile kernel then early-returns) and
  records the violation in count[1]; the buffer ABI is unchanged
  (count index 1 was previously unused). try_gemma4_expert_qmm allocates
  the second count element, drains the encoder after the builder, and
  re-routes a retracted call to the order-agnostic legacy path instead of
  dispatching the tile kernel (zero count is unambiguous: the selector's
  assignment gate guarantees M is 4096/8192/16384).
- N2 (route-condition duplication): the sorted-RHS gate literal that
  appeared (negated) in the diagnostics record and in the dispatch
  decision is now the shared static constexpr predicate
  takes_sorted_rhs_route, so future tuning of the 16/4 thresholds cannot
  desynchronize counter vs route.
- N3 (per-call bias normalization): gather_qmm_rhs no longer spends
  ensure_row_contiguous on biases before classification reads the raw
  tensor's fields; normalization runs only inside the winning-route
  branch (hit semantics unchanged; the legacy block keeps its own
  normalization point and ordering).
- N4 (armed_ data race): Gemma4ExpertQMMCounters::armed_ is now
  std::atomic<bool> with relaxed loads/stores in armed(), snapshot(),
  snapshot_and_disarm() (read-then-write order preserved) and
  clear_and_arm(); the class remains non-copyable, now enforced.

* fix(mlx): make the R1 sortedness fail-safe sound; proper retract attribution

F1: the per-expert boundary vote was a partial detector -- an inversion
inside a segment used by no other expert's boundary could escape, so
"re-route on any violation" overclaimed. build_gemma4_sorted_expert_tiles_bm32
now also runs a strided adjacent-pair scan: thread lid checks
indices[i-1] <= indices[i] for i = lid+1; i < M; i += 128, covering every
adjacent pair in [1, M) exactly once (1..128 iterations at the reachable
M in {4096,8192,16384}). Adjacent-pair monotonicity is transitive, so a
clean scan is a sound and complete sortedness oracle; it folds into the
same simd_or/threadgroup vote and the same retract (count[0]=0, count[1]=1).
The boundary checks stay as cheap, precise diagnostics.

F2: retracts were write-only in count[1] and surfaced as
fallback_metallib_unavailable -- misattribution in the only observable
surface. A dedicated fallback_sortedness_retracted counter now rides the
GemmA4 route counters and the C diagnostics ABI
(sizeof 80 -> 88, new uint64 at offset 80; existing offsets unchanged).
try_gemma4_expert_qmm returns the route class: count[0]==0 with count[1]==1
records fallback_sortedness_retracted, any other unusable build keeps
fallback_metallib_unavailable, then re-routes to the legacy path as before.

F4: new doctest drives the full armed() -> clear_and_arm() ->
snapshot_and_disarm() cycle and the attempts == hits + fallbacks invariant
including the new class; the route-table and counter-invariant tests now
cover fallback_sortedness_retracted.

Verified: cmake tests 262/262 + 3550 assertions pass; metal -Wall -Wextra
-fno-fast-math compile of kernels/quantized.metal is warning-free.

* perf(metal): E=256 expert-tile route + trust + gpu::eval UAF fix — darkbloom-base mirror (#7)

* perf(metal): instantiate E=256 expert-tile route for Qwen 3.5/3.6 MoE prefill (mirror of Cmlx/mlx 58fab46)

* fix(metal): use-after-free in gpu::eval for primitives that synchronize mid-eval (mirror)

* perf(metal): trust mode skips retract readback (mirror)

* fix(compile): preserve all-cache binding cleanup

---------

Co-authored-by: JasonHonKL <148705846+JasonHonKL@users.noreply.github.com>
Co-authored-by: AK <144495202+AKnassa@users.noreply.github.com>
Co-authored-by: Cheng <git@zcbenz.com>
Co-authored-by: Adityaj0 <93090622+Adityaj0@users.noreply.github.com>
Co-authored-by: anchor <codeanqiang@gmail.com>
Co-authored-by: codeAnqiang-ma <273298913+codeAnqiang-ma@users.noreply.github.com>
Co-authored-by: Rohan Gautam <rohan1gautam@gmail.com>
Co-authored-by: Ayaan Gazali <ayaangazali.work@gmail.com>
Co-authored-by: Erwin Zhang <59893706+erwinzhang7@users.noreply.github.com>
Co-authored-by: katlun-lgtm <katlun@gmail.com>
Co-authored-by: katlun-lgtm <264247399+katlun-lgtm@users.noreply.github.com>
Co-authored-by: robertomeroni <150194833+robertomeroni@users.noreply.github.com>
Co-authored-by: Fu Xiaonan <ht3fudatou@163.com>
Co-authored-by: Fu Xiaonan <214359569+FU-max-boop@users.noreply.github.com>
Co-authored-by: Hao Xu <hxu44@apple.com>
Co-authored-by: Feli <89400571+FeliGame@users.noreply.github.com>
Co-authored-by: Feli <feli@hnu.edu.cn>
Co-authored-by: Eyüp Can Akman <eyupcanakman@gmail.com>
Co-authored-by: Cheng <zcbenz@gmail.com>
Co-authored-by: yentur <mr.yentur@gmail.com>
Co-authored-by: Alessio Pollero <alessio.pollero@gmail.com>
Co-authored-by: Zhiqi Zhang <zhiqizhangg@gmail.com>
Co-authored-by: Daniel Hiltgen <dhiltgen@users.noreply.github.com>
Co-authored-by: Tanish Jain <recklurker@gmail.com>
Co-authored-by: hojin12312 <hojin12312@gmail.com>
Co-authored-by: Xiang Chen <46052474+x14ngch3n@users.noreply.github.com>
Co-authored-by: x14ngch3n <x14ngch3n@users.noreply.github.com>
Co-authored-by: Duhyeon, Kim <49020301+dudududukim@users.noreply.github.com>
Co-authored-by: rohith <kapellirohith@gmail.com>
Co-authored-by: Ishaan Samantray <devteam.aegis@gmail.com>
Co-authored-by: Aaishwarya Mishra <aaishwarymishra@gmail.com>
Co-authored-by: Anastasiia Filippova <a_filippova@apple.com>
Co-authored-by: Yanzhao Wang <19340816+wyanzhao@users.noreply.github.com>
Co-authored-by: XXXXRT666 <157766680+XXXXRT666@users.noreply.github.com>
Co-authored-by: Jake Bowhay <60778417+j-bowhay@users.noreply.github.com>
Co-authored-by: vraj patel <87225460+vraj00222@users.noreply.github.com>
Co-authored-by: Gusanidas <33495733+Gusanidas@users.noreply.github.com>
Co-authored-by: Dwijen Patel <dwijen@gmail.com>
Co-authored-by: Vladimir Iglovikov <ternaus@users.noreply.github.com>
Co-authored-by: Brian C. <94733710+deBrian07@users.noreply.github.com>
Co-authored-by: Daniel Hiltgen <daniel.hiltgen@ollama.com>
Co-authored-by: YH Yan <strayberry0w0@gmail.com>
Co-authored-by: katlun-lgtm <katlun@windyviews.com>
Co-authored-by: anupsv <6407789+anupsv@users.noreply.github.com>
Co-authored-by: Gajesh Naik <26431906+Gajesh2007@users.noreply.github.com>
Co-authored-by: David Tai <davidtai@Davids-MBP.lan>
tpegolotti pushed a commit that referenced this pull request Aug 26, 2026
tpegolotti pushed a commit that referenced this pull request Aug 26, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants