开源llama.cpp Releases·原文 2026年10月29日

llama.cpp合并SYCL重排反量化向量化改动,MXFP4 MoE预填充提速约 57%–61%

llama.cpp发布b11979版本,合入PR #30976:在SYCL后端把重排布局下多种量化格式的反量化存储改为向量化写,覆盖Q4_0、Q4_K、Q5_K、Q6_K、Q3_K、IQ3、Q2_K和MXFP4。提交说明给出的测速均来自Arc Pro B70单卡,提速集中在“已安装重排布局张量后再做prefill”的场景。

AI解读:这次改动只针对SYCL后端重排布局(reordered layout)下的反量化转换内核:把逐元素标量存储改为vec存储,或重新划分work-group/线程映射,使每个线程一次写出 4 个或 8 个连续输出。提交说明称逐元素数学与舍入不变,输出与标量内核逐位相同。

提速数据有条件:几乎全部为Arc Pro B70上的“primed prefill”,即张量装好重排布局、先解码再做prefill的场景。多数格式的fresh-layout prefill、decode和困惑度(PPL)没有变化,PPL在运行误差内。因此不能把数字理解为所有推理场景普遍提速。

按提交说明,Q4_0在该场景下pp512/pp2048提升 6.4%–7.3%(测试到 27B稠密模型);Q4_K提升 14.7%–15.3%;Q6_K在Q6_K 27B模型上提升 28.7%;Q3_K提升 16.8%;IQ3_XXS/IQ3_S提升 20.3%–21.1%;Q2_K提升 13.1%;Q5_K在q5_K主导的 27B模型上提升 3.6%。

MXFP4的收益最大:在MXFP4 MoE 20B模型上,pp2048提升 57%、pp512提升 61%,primed与fresh两种预填充均如此,decode不变。提交说明指出,MXFP4的重排布局只安装在MoE专家路径(reorder_qw_mxfp4_moe)上,稠密MXFP4张量仍用普通布局,所以这一数字不适用于稠密MXFP4模型。

实现上有两处删除:Q4_0和MXFP4的旧标量重排内核被移除,Q4_0采用与 #30775 中Q8_0内核相同的扁平sycl::range启动方式,每work-item处理 4 个qs字节(8 个输出)。Q5_K的GET_ROWS重排路径也改为同样的 32 线程分解。

b11979同时提供macOS、Linux、Windows、Android等平台的预编译包,含SYCL FP32/FP16、CUDA 12/13、Vulkan、ROCm等后端版本;上述性能数字仅对应SYCL后端改动。

llama.cpp发布b11979版本,合入PR #30976“sycl : vectorize reordered-layout dequantization stores (Q4_0, K-quants, IQ3, MXFP4)”。改动对象是SYCL后端在重排布局下把量化权重转换为浮点的内核,做法是扩大每次写出的向量宽度,或重新划分线程与输出之间的映射。

提交说明称,各内核的逐元素数学和舍入保持不变,输出与改动前的标量内核逐位相同。所有测速都在Arc Pro B70上完成;除MXFP4外,提速均出现在“primed prefill”(张量已安装重排布局、先解码再prefill)场景,fresh-layout prefill、decode与PPL多数没有变化。

各量化格式的改动与测速

Q4_0:原内核每个work-item处理 1 个qs字节、做两次相隔 16 个元素的标量存储;改为与 #30775 中Q8_0相同的扁平sycl::range启动,每个work-item处理 4 个qs字节(8 个输出)、做两次vec存储。原有的k % QK4_0 == 0 断言保证按整 32 元素块覆盖,k/8启动恰好覆盖每一行。旧标量重排内核被删除。Arc Pro B70、primed prefill下Q4_0模型pp512/pp2048提升 6.4%–7.3%(测试到 27B稠密)。

Q4_K:原内核已用sycl::vec加载量化输入,但用标量存储写出反量化结果;改为在vec寄存器中构建结果并发出两次向量存储。Arc Pro B70、primed prefill下Q4_K主导模型提升 14.7%–15.3%(测试到 27B稠密)。

Q5_K:原内核每个线程只有 4 个输出,瓶颈不在存储;单独加宽存储在混合量化模型上处于噪声范围,在q5_K主导模型上为 +0.5%。改为每块 32 线程、每线程消费 4 个连续ql和qh字节并发出两次对齐vec存储;4 元素一组共享一对scale,逐元素数学与输出不变。q5_K主导的 27B模型primed prefill提升 3.6%。GET_ROWS重排路径改为同样的 32 线程分解。

Q6_K:原内核每线程写 4 个相隔 32 元素的标量;重新划分同一 64 线程工作组,使每线程发出一个对齐vec。4 元素一组共享一个 16 元素scale组、一个ql nibble来源和一对qh位,逐元素数学与输出不变。Q6_K 27B模型primed prefill提升 28.7%。

Q3_K:每线程已计算 4 个连续输出,改为在vec寄存器中构建并发出一次向量存储。Q3_K主导的 27B模型primed prefill提升 16.8%。

IQ3:iq3_xxs与iq3_s两个内核每线程写 8 个连续输出(来自两次grid查找各 4 个),改为构建两个vec寄存器并发出两次向量存储。IQ3_XXS与IQ3_S主导的 27B模型primed prefill提升 20.3%–21.1%。

Q2_K:按Q6_K同样方式重新划分,每线程发出一个对齐vec,而不是 4 个相隔 32 的标量;4 元素一组共享一个scale字节和一个 2-bit lane。Q2_K主导的 27B模型primed prefill提升 13.1%。

MXFP4:采用与Q4_0相同的扁平sycl::range启动,每个work-item处理 4 个qs字节(8 个输出)、两次vec存储;相同的LUT查找与逐块scale,输出逐位相同;k % QK_MXFP4 == 0 断言保证按整块覆盖,k/8启动恰好覆盖每一行,旧标量内核被删除。提交说明称,MXFP4重排布局只安装在MoE专家路径reorder_qw_mxfp4_moe上,稠密MXFP4张量保持普通布局。MXFP4 MoE 20B模型上,primed和fresh prefill均约提升 57%(pp2048)/61%(pp512),decode不变,PPL在该模型运行噪声内。

这些数字的适用范围

提交说明的测速全部来自单张Arc Pro B70;除MXFP4外,提速场景限定为已安装重排布局张量后的prefill,且模型规模测试到 27B稠密。fresh-layout prefill、decode与PPL未显示变化,因此这些百分比不应外推到首次加载布局、解码吞吐或其他显卡。

MXFP4的 57%–61% 对应MoE 20B模型的专家路径;稠密MXFP4不走该重排布局,提交说明未给出其对应数字。

版本产物

b11979发布页提供面向macOS(Apple Silicon arm64、Intel x64)、iOS、Linux、Android、Windows等平台的预编译包,Linux与Windows包包含CPU、Vulkan、CUDA 12/13、ROCm、OpenVINO以及SYCL FP32/FP16等后端。上述性能改动作用于SYCL后端,其他后端的构建产物不因此获得相同测速结果。

信息来源

llama.cpp Releases原始来源