开源llama.cpp Releases·原文 2026年9月16日

llama.cpp b10994修复Metal后端MUL_MAT_ID在激活值超过f16范围时产生NaN

该版本修复了Metal路径上mul_mm_id将激活值收窄为half导致溢出为inf、进而使 8x8累加器整块变成NaN的问题;Mistral Small 4在Metal上预填充超过等于 32 个token时会返回全NaN词表。

AI解读:llama.cpp发布b10994,修的是Metal后端一个会让整段推理直接崩掉的问题:mul_mm_id在走simdgroup MMA时把激活值收窄成half(f16),而f16最大只能到 65504,一旦模型的激活值超过这个数就会变成inf,随后simdgroup_multiply_accumulate把整个 8x8累加器tile变成NaN。

这不是理论上的边角情况。发布说明里点名的真实模型是Mistral Small 4(arch mistral4,128 个专家、4 个激活),它某一层的激活值能达到约 1e5;在Metal上,凡是预填充达到或超过 32 个token,返回的就是一整份NaN词表。用Apple M2 Max跑b10156的复现日志显示,MUL_MAT_ID在index 0就出现NaN(MTL0=nan,CPU=583442.375000)。

切换点很清楚:ne21_mm_id_min(32)以下的mul_mv_id路径用f32承载同样的数值,结果是正确的,所有CPU路径也正确。也就是说 32 是那条分界线,32 以下走对的路,32 及以上走错的路。

修复同时补上了测试。原有的init_mul_mat_id_tensors用均匀分布 [-1, 1] 初始化,根本没法把操作数推出f16范围,所以这个问题此前测不出来。test_mul_mat_id新增amax参数(默认 1.0f,完全保留历史初始化),只缩放f32激活值,量化权重保持正常范围。

新增六个用例覆盖两种形状:q4_K、128 专家、4 激活、4096x2048对应真实模型;q8_0、8 专家、2 激活、512x256用最小规模复现同一失败。n=16 在切换点以下,是必须保持绿色的对照组;n=32 和n=64 在切换点以上,目前在Metal上失败。

对在Apple芯片上跑MoE、尤其是专家数多且激活值偏大的模型的用户,升级到b10994是直接的绕开方式;只跑CPU或只在 32 token以下做短提示的场景不受这个bug影响。发布页照例提供macOS、iOS、Linux、Windows、Android的预编译包,含CUDA 12.8/13.3、ROCm 10.0、Vulkan、SYCL、OpenVINO等后端变体,另有UI包。

llama.cpp发布b10994,修复Metal后端在mul_mm_id中当激活值超出f16范围时产生NaN的问题(PR #26223)。

发布说明给出的根因是:Metal的mul_mm_id路径为simdgroup MMA把src1收窄为half(每个实例化中的S1 = half;ggml-metal.metal:10582、:10595,tensor-ops路径镜像在 :10643/:10654)。f16在 65504 饱和,激活值超过该值即产生inf,simdgroup_multiply_accumulate随即把整个 8x8累加器tile变成NaN。

ne21_mm_id_min(32)以下的mul_mv_id路径以f32承载同样的值,结果正确;所有CPU路径也正确。

发布说明点名的真实模型是Mistral Small 4(arch mistral4,128 个专家、4 个激活),其中某一层的激活值达到约 1e5;在Metal上,每次预填充达到或超过 32 个token都会返回完全为NaN的词表。

复现记录(Apple M2 Max、macOS、llama.cpp b10156):MUL_MAT_ID(type_a=q8_0,...,n=32,k=256,amax=100000.000000): [MUL_MAT_ID] NaN at index 0 (MTL0=nan CPU=583442.375000) FAIL。

测试侧如何让这个bug可被复现

发布说明指出,该问题此前无法被测试:init_mul_mat_id_tensors以均匀分布 [-1, 1] 初始化,现有用例无法把操作数推出f16范围。

test_mul_mat_id新增amax参数(默认 1.0f,精确保留历史初始化),只缩放f32激活值,量化权重保持正常范围。

  • 新增六个用例。n=16 位于mul_mv_id到mul_mm_id的切换点以下,是必须保持绿色的对照;n=32 与n=64 位于切换点以上,目前在Metal上失败。
  • 覆盖两种形状,因为该问题不针对特定模型或规模:q4_K、128 专家、4 激活、4096x2048对应真实模型;q8_0、8 专家、2 激活、512x256以最小规模显示同样的失败。

提交与发布信息

提交信息显示改动包含修复var scope,并一度将测试注释掉以免暂时打断CI。发布说明署名的协助者为Claude Fable 5.1,共同署名者包括Claude Fable 5与Georgi Gerganov。

  • b10994提供macOS Apple Silicon(arm64)、macOS Intel(x64)、iOS XCFramework等构建。
  • Linux侧提供Ubuntu x64/arm64/s390x CPU、Vulkan、CUDA 12.8与CUDA 13.3(含CUDA库)、ROCm 10.0、OpenVINO 2026.3.1、SYCL FP32/FP16。
  • Windows侧提供x64/arm64 CPU、OpenCL Adreno、CUDA 12.4与CUDA 13.4(含库)、Vulkan、OpenVINO、SYCL、ROCm 10.0。
  • 另有Android arm64 CPU构建与UI包;macOS Apple Silicon的KleidiAI构建和openEuler构建标记为DISABLED。

信息来源

llama.cpp Releases原始来源