Brief
llama.cpp fixes Metal NaN for activations beyond f16 range
A Metal bug in llama.cpp turned entire vocabularies into NaN for models whose activations exceed f16 range. The fix targets the mul_mm_id path, and a new test parameter reproduces the failure.
The Metal mul_mm_id path in llama.cpp narrowed src1 to half for the simdgroup MMA. Because f16 saturates at 65504, activations above that became inf, and simdgroup_multiply_accumulate then turned the whole 8x8 accumulator tile into NaN. The mul_mv_id path below ne21_mm_id_min (32) carries the same values in f32 and was already correct, as was every CPU path.
The failure was untestable before because init_mul_mat_id_tensors initializes uniform [-1, 1]. test_mul_mat_id now gains an amax parameter (default 1.0f) that scales only the f32 activations. Six cases cover n=16 as a control and n=32 and n=64 as failing on Metal. The real model behind this is Mistral Small 4, one of whose layers reaches ~1e5 activations: on Metal every prefill of >=32 tokens returned an entirely NaN vocabulary.
Our reading
Our reading is that this fix matters mainly for local Metal users running large mixture-of-experts models, and the uncovered dense kernel suggests the same class of failure may still lurk elsewhere.
What to do or watch
Watch the dense kernel_mul_mm path: source text says it has the identical f16-to-half conversion of src1 and is expected to fail the same way, but the new amax test cases only cover mul_mm_id — so whether that dense path actually produces NaN on Metal remains the open question. If you run large mixture-of-experts models on Metal, update past this fix and re-check prefill output, since a single layer reaching activations beyond f16 range was enough to return an entirely NaN vocabulary.
Source details and supporting facts
Each line is stated by the page named above it.
Stated by github.com
- The Metal mul_mm_id path narrows src1 to half for the simdgroup MMA.
- f16 saturates at 65504, so a model whose activations exceed that produces inf, and simdgroup_multiply_accumulate then turns the whole 8x8 accumulator tile into NaN.
- The real model behind this is Mistral Small 4 (arch mistral4, 128 experts / 4 active), one of whose layers reaches ~1e5 activations: on Metal every prefill of >=32 tokens returns an entirely NaN vocabulary.
- kernel_mul_mm (dense) has the identical conversion and is expected to fail the same way; it is not covered here.
Sources
- llama.cpp releases (GitHub)Text stored 16 September 2026
How this story was checked. Written from the 1 page listed above, stored 16 September 2026; claims checked against that stored text on 16 September 2026.
What that means
- 4 of 4 reported statements were confirmed against the page that carries them; the rest were removed rather than published.
- Figures in the text were required to appear in the stored source text: yes. Identifiers: yes.
- The check reads stored text only: no claim rests on a fresh look that did not happen.
- Where the reporting was silent, the text says so instead of filling the gap.