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.
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.