人工智能大模型推理引擎本地部署模型量化模型优化【免费下载链接】ik_llama.cppllama.cpp fork with additional SOTA quants and improved performance项目地址https://gitcode.com/GitHub_Trending/ik/ik_llama.cpp点击查看免费下载本文基于 ik_llama.cpp 仓库的 PR #446Fix bug in MMVQ kernel及其关闭的三个 issue#389、#398、#425完整复盘一次发生在 CUDA 量化矩阵-向量乘内核MMVQ中的隐蔽 bug 定位、根因分析与修复过程。读者将理解 MMVQ 内核在 MoE 模型推理中的关键作用、为什么一次只处理 2~3 个 token会成为触发崩溃的罕见条件以及当前仓库源码中该内核的实现结构。背景三个症状各异、根因相同的崩溃报告2025 年 5 月ik_llama.cpp 仓库陆续收到三份崩溃报告表面症状各不相同最终却都指向同一个根因#389 - llama-batched-bench 在 batch size 2 时崩溃用户QuPengfei在 Qwen3-235B-A22B128 个专家、8 个激活专家的 MoE 模型上以-ser 7,1等参数运行llama-batched-bench时程序在 PPprompt processing阶段刷出大量GGML_ASSERT(fms.S[j] 0) failed断言失败后 abortcore dumped。日志中反复出现的调用栈指向iqk_flash_attn_implissue 记录。#398 --fmoe引发 illegal memory access用户pt13762104在双 Tesla T4 上运行 Qwen3-30B-A3B 时发现只要开启-fmoe融合 MoE服务器运行片刻后必然报CUDA error: an illegal memory access was encountered错误定位在ggml_cuda_up_gate_unaryggml-cuda.cu中一次cudaMemcpyAsync且总是发生在 device 1issue 记录。#425 - 加载 DeepSeek-V3-0324 量化模型即崩溃用户nux用DeepSeek-V3-0324-IQ4_K_R4ik_llama.cpp 特有的 R4 量化运行llama-server时prompt processing 一开始就触发CUDA error: an illegal memory access was encountered同时内核日志出现NVRM: Xid 31 ... MMU Fault即使去掉-fmoe也无法避免issue 记录。这三份报告有个共同点问题都发生在多 GPU 环境双 T4、三块 3090 等而项目作者ikawrakow只有单 GPU 机器无法复现导致排查周期很长——直到社区成员ciprianveg在定位过程中提供了决定性帮助。根因定位MMVQ 内核在2~3 个 token路径上的缺陷PR #446 的描述给出了最终结论PR 记录The bug was in the CUDA matrix-vector multiplication kernel (a.k.a., MMVQ). It only shows up when the kernel processes 2 or 3 tokens. Hence, it was not observed during TG, and only showed up during PP when an expert in a MoE model ended up with having to process just 2 or 3 tokens from the batch (which is rare).关键信息有三点Bug 位于 MMVQmatrix-vector multiplication for quantized types内核即 CUDA 上的量化矩阵-向量乘实现只有内核一次处理 2 或 3 个 token 时才触发。因此在 TGtoken generation单 token 逐次解码阶段永远观察不到只有 PPprompt processing批量处理阶段才会暴露MoE 模型是放大器PP 阶段一个 batch 通常有大量 token但 MoE 的专家路由expert routing会把 token 分发到不同专家某个专家最终可能只分到 batch 中 2~3 个 token——此时 MMVQ 内核恰好处在最容易出错的路径上于是本应罕见的崩溃被 MoE 架构常态化了。这也解释了 #389 中-ser 7,1-ser即--smart-expert-reduction智能专家缩减参数为i,f组合见 llama-bench.cpp为什么会触发崩溃专家缩减改变了每个专家实际处理的 token 分布使单个专家只处理极少量 token的情况频繁出现从而撞上 MMVQ 的缺陷路径。PR #446 作者明确指出该 PR 修复的是真实存在的 bug无论能否复现这三个 issue都应该合并并认为此前 #442 中的其他改动可能并非必要I believe all other changes I made in #442 are not necessary。PR 于 2025-05-23 提交、次日确认合并仅修改了ggml/src/ggml-cuda/mmvq.cu一个文件5 行变更commit193a15b。源码级剖析MMVQ 内核的结构与可疑路径当前仓库中的 MMVQ 实现分散在ggml/src/ggml-cuda/下的一组文件中正好可以用来对照理解 bug 所在的代码区域。1. 参数聚合层mmvq.cu入口函数ggml_cuda_op_mul_mat_vec_q_implmmvq.cu把矩阵的维度信息、行区间row_low/row_high、batch 的 token 数src1_ncols等打包进mmvq_args再按权重张量的量化类型分发到对应的模板 kernel。值得注意的一处多 GPU 相关逻辑// the main device has a larger memory buffer to hold the results from all GPUs // nrows_dst nrows of the matrix that the kernel writes into const int64_t nrows_dst id ctx.device ? ne0 : row_diff;mmvq.cu从源码结构可以推断多 GPU 场景下非主设备offloaded 层所在设备只写入自己负责的行区间row_diff而主设备需要为整块ne0行准备结果缓冲区。这种每个设备写不同行数的布局一旦 kernel 对行号或结果指针的计算出现偏差就会造成越界写最终表现为illegal memory access——与 #398、#425 报错总是发生在非主设备device 1/device 2的现象吻合。ggml_cuda_op_mul_mat_vec_q_id中还有一个与 batch 大小直接相关的断言mmvq.cuGGML_ASSERT(src1-ne[1] MMVQ_MAX_BATCH_SIZE src1-ne[2] 1);其中src1-ne[1]就是一次送入 MMVQ 内核的 token 数batch 维度。2. batch 上限与内核特化mmvq.cuh与mmvq-templates.cuhMMVQ 并非一个万能内核而是针对不同 batch 大小做了模板特化。上限定义在 mmvq.cuh#define MMVQ_MAX_BATCH_SIZE 8 // Max. batch size for which to use MMVQ kernels.即当 token 数 ≤ 8 时矩阵-向量乘才会走 MMVQ 路径更大的 batch 会落入其他矩阵乘实现如 cuBLAS。在mul_mat_vec_q_cuda_Tmmvq-templates.cuh中ncols_y即 token 数从 1 到 8 分别实例化出独立的 kernelcase 1mul_mat_vec_qtype, 1, nwarps第 403-406 行case 2mul_mat_vec_qtype, 2, nwarps第 407-410 行case 3mul_mat_vec_qtype, 3, nwarps第 411-414 行case 4~8依此类推直到MMVQ_MAX_BATCH_SIZEncols_y 2和ncols_y 3正是 PR #446 描述的处理 2 或 3 个 token 时触发的模板实例。内核内还有一个随ncols_y变化的分支#if defined(GGML_USE_HIPBLAS) defined(__HIP_PLATFORM_AMD__) (defined(RDNA2) || defined(RDNA3)) constexpr int rows_per_cuda_block 1; #else constexpr int rows_per_cuda_block ncols_y 4 ? 1 : 2; #endifmmvq-templates.cuhrows_per_cuda_block表示一个 CUDA block 处理权重矩阵的几行token 数小于 4 时每 block 只处理 1 行否则处理 2 行。从源码结构可以推断ncols_y 2/ncols_y 3恰好落在batch 很小 每 block 单行这条特殊路径上任何针对行数、nrows_dst或结果写回索引dst[j*nrows_dst row0 threadIdx.x]见 mmvq-templates.cuh的边界计算错误都会在此处集中爆发——这与TG 不崩、PP 罕见崩的现象完全自洽。3. 量化类型全覆盖iqk_mmvq.cuik_llama.cpp 的另一特色是大量自研量化类型IQ2_K、IQ3_K、IQ4_K_R4、IQ1_S_R4等。这些类型的 MMVQ 实现在 iqk_mmvq.cu 中通过iqk_mul_mat_vec_q统一分发覆盖了二十余种 IK 特有量化格式mmvq.cu 中这些类型全部落入iqk_mul_mat_vec_q分支。#425 中用户使用的正是 ik_llama.cpp 特有的IQ4_K_R4量化说明该 bug 波及的是整个 MMVQ 家族而非某个特定量化类型。修复与验证一次典型的社区协作 debug 闭环从 issue 时间线可以看到整个修复过程的完整闭环所有记录均来自仓库的 github-data2025-05-07 ~ 05-15三个 issue 陆续提交。作者在 #398 中坦言这更可能是多 GPU 配置下的 bug我只有单 GPU 无法调试#425 中用户补充了NVRM Xid 31的内核日志指向真正的硬件级非法内存访问。社区逐步收窄范围#398 中用户测试发现单 GPU 下-fmoe正常、多 GPU 必崩#389 中发现-ser 7,1是触发器#425 中发现换用 IK 特有混合量化走dequantize - cuBLAS路径无量化 MMVQ 实现就不崩这从反面印证了问题出在量化矩阵乘实现上。2025-05-23ciprianveg在定位中起到关键作用作者发布 #446 修复修改mmvq.cu一行核心逻辑共 5 行变更。2025-05-24PR 合并。此前受困的用户如pt13762104反馈现在一切正常。这一过程也留下了对使用者有价值的经验当llama-server/llama-batched-bench在多 GPU MoE 模型 量化权重组合下出现illegal memory access或GGML_ASSERT崩溃时除了检查显存还应考虑是否命中了 MMVQ 的 batch 边界路径——升级到包含 #446 修复的构建版本是首要动作。结语PR #446 是理解 ik_llama.cpp 这类性能优化型 fork风险与收益的一个绝佳样本为了榨取量化矩阵乘的性能项目维护了覆盖 30 量化类型、按 batch 大小特化的 MMVQ 内核家族MMVQ_MAX_BATCH_SIZE 8而性能优化的复杂度直接转化为了隐藏的边界 bug 概率。该 bug 的隐蔽性在于日常 TG 场景的 batch 永远为 1恰好绕开了故障路径只有 MoE 推理在 PP 阶段把 batch 切成 2~3 个 token 分给某个专家时缺陷才被触发。修复虽小一个文件、5 行但对所有在消费级多 GPU 上运行 Qwen3-MoE、DeepSeek-V3 等模型的用户意义重大。对于希望进一步研究源码的读者建议按以下顺序阅读入口分发 mmvq.cu → 参数结构 mmvq-args.h → 内核模板 mmvq-templates.cuh → IK 量化分发 iqk_mmvq.cu配合 llama-bench.cpp 中-ser/-fmoe等参数的解析逻辑即可完整复原这条 bug 的产生—触发—定位—修复链路。赞分享人工智能大模型推理引擎本地部署模型量化模型优化【免费下载链接】ik_llama.cppllama.cpp fork with additional SOTA quants and improved performance项目地址https://gitcode.com/GitHub_Trending/ik/ik_llama.cpp点击查看免费下载相关推荐ik_llama.cpp CUDA 非法内存访问崩溃全记录:MoE 模型 MMVQ 内核 2/3 行越界 Bug 的排查与修复ik_llama.cpp CUDA 非法内存访问崩溃全记录:MoE 模型 MMVQ 内核 2/3 行越界 Bug 的排查与修复 本文基于 ik_llama.cp人工智能大模型推理引擎本地部署模型量化模型优化扫描文档中的手写签名提取从传统办公到数字化处理的智能解决方案扫描文档中的手写签名提取从传统办公到数字化处理的智能解决方案 在数字化转型浪潮中许多机构仍面临一个看似简单却异常棘手的挑战如何高效地从大量纸质文档中提取手人工智能大模型推理引擎本地部署模型量化模型优化ik_llama.cpp PR 430 解析MoE multi-add 优化的引入、崩溃与回退决策ik_llama.cpp PR 430 解析MoE multi add 优化的引入、崩溃与回退决策 导读 本文以 ik_llama.cpp 仓库中的 PR 4人工智能大模型推理引擎本地部署模型量化模型优化创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
深入解析 ik_llama.cpp PR 446：MMVQ 内核中隐藏的 MoE 崩溃 bug 修复
人工智能大模型推理引擎本地部署模型量化模型优化【免费下载链接】ik_llama.cppllama.cpp fork with additional SOTA quants and improved performance项目地址https://gitcode.com/GitHub_Trending/ik/ik_llama.cpp点击查看免费下载本文基于 ik_llama.cpp 仓库的 PR #446Fix bug in MMVQ kernel及其关闭的三个 issue#389、#398、#425完整复盘一次发生在 CUDA 量化矩阵-向量乘内核MMVQ中的隐蔽 bug 定位、根因分析与修复过程。读者将理解 MMVQ 内核在 MoE 模型推理中的关键作用、为什么一次只处理 2~3 个 token会成为触发崩溃的罕见条件以及当前仓库源码中该内核的实现结构。背景三个症状各异、根因相同的崩溃报告2025 年 5 月ik_llama.cpp 仓库陆续收到三份崩溃报告表面症状各不相同最终却都指向同一个根因#389 - llama-batched-bench 在 batch size 2 时崩溃用户QuPengfei在 Qwen3-235B-A22B128 个专家、8 个激活专家的 MoE 模型上以-ser 7,1等参数运行llama-batched-bench时程序在 PPprompt processing阶段刷出大量GGML_ASSERT(fms.S[j] 0) failed断言失败后 abortcore dumped。日志中反复出现的调用栈指向iqk_flash_attn_implissue 记录。#398 --fmoe引发 illegal memory access用户pt13762104在双 Tesla T4 上运行 Qwen3-30B-A3B 时发现只要开启-fmoe融合 MoE服务器运行片刻后必然报CUDA error: an illegal memory access was encountered错误定位在ggml_cuda_up_gate_unaryggml-cuda.cu中一次cudaMemcpyAsync且总是发生在 device 1issue 记录。#425 - 加载 DeepSeek-V3-0324 量化模型即崩溃用户nux用DeepSeek-V3-0324-IQ4_K_R4ik_llama.cpp 特有的 R4 量化运行llama-server时prompt processing 一开始就触发CUDA error: an illegal memory access was encountered同时内核日志出现NVRM: Xid 31 ... MMU Fault即使去掉-fmoe也无法避免issue 记录。这三份报告有个共同点问题都发生在多 GPU 环境双 T4、三块 3090 等而项目作者ikawrakow只有单 GPU 机器无法复现导致排查周期很长——直到社区成员ciprianveg在定位过程中提供了决定性帮助。根因定位MMVQ 内核在2~3 个 token路径上的缺陷PR #446 的描述给出了最终结论PR 记录The bug was in the CUDA matrix-vector multiplication kernel (a.k.a., MMVQ). It only shows up when the kernel processes 2 or 3 tokens. Hence, it was not observed during TG, and only showed up during PP when an expert in a MoE model ended up with having to process just 2 or 3 tokens from the batch (which is rare).关键信息有三点Bug 位于 MMVQmatrix-vector multiplication for quantized types内核即 CUDA 上的量化矩阵-向量乘实现只有内核一次处理 2 或 3 个 token 时才触发。因此在 TGtoken generation单 token 逐次解码阶段永远观察不到只有 PPprompt processing批量处理阶段才会暴露MoE 模型是放大器PP 阶段一个 batch 通常有大量 token但 MoE 的专家路由expert routing会把 token 分发到不同专家某个专家最终可能只分到 batch 中 2~3 个 token——此时 MMVQ 内核恰好处在最容易出错的路径上于是本应罕见的崩溃被 MoE 架构常态化了。这也解释了 #389 中-ser 7,1-ser即--smart-expert-reduction智能专家缩减参数为i,f组合见 llama-bench.cpp为什么会触发崩溃专家缩减改变了每个专家实际处理的 token 分布使单个专家只处理极少量 token的情况频繁出现从而撞上 MMVQ 的缺陷路径。PR #446 作者明确指出该 PR 修复的是真实存在的 bug无论能否复现这三个 issue都应该合并并认为此前 #442 中的其他改动可能并非必要I believe all other changes I made in #442 are not necessary。PR 于 2025-05-23 提交、次日确认合并仅修改了ggml/src/ggml-cuda/mmvq.cu一个文件5 行变更commit193a15b。源码级剖析MMVQ 内核的结构与可疑路径当前仓库中的 MMVQ 实现分散在ggml/src/ggml-cuda/下的一组文件中正好可以用来对照理解 bug 所在的代码区域。1. 参数聚合层mmvq.cu入口函数ggml_cuda_op_mul_mat_vec_q_implmmvq.cu把矩阵的维度信息、行区间row_low/row_high、batch 的 token 数src1_ncols等打包进mmvq_args再按权重张量的量化类型分发到对应的模板 kernel。值得注意的一处多 GPU 相关逻辑// the main device has a larger memory buffer to hold the results from all GPUs // nrows_dst nrows of the matrix that the kernel writes into const int64_t nrows_dst id ctx.device ? ne0 : row_diff;mmvq.cu从源码结构可以推断多 GPU 场景下非主设备offloaded 层所在设备只写入自己负责的行区间row_diff而主设备需要为整块ne0行准备结果缓冲区。这种每个设备写不同行数的布局一旦 kernel 对行号或结果指针的计算出现偏差就会造成越界写最终表现为illegal memory access——与 #398、#425 报错总是发生在非主设备device 1/device 2的现象吻合。ggml_cuda_op_mul_mat_vec_q_id中还有一个与 batch 大小直接相关的断言mmvq.cuGGML_ASSERT(src1-ne[1] MMVQ_MAX_BATCH_SIZE src1-ne[2] 1);其中src1-ne[1]就是一次送入 MMVQ 内核的 token 数batch 维度。2. batch 上限与内核特化mmvq.cuh与mmvq-templates.cuhMMVQ 并非一个万能内核而是针对不同 batch 大小做了模板特化。上限定义在 mmvq.cuh#define MMVQ_MAX_BATCH_SIZE 8 // Max. batch size for which to use MMVQ kernels.即当 token 数 ≤ 8 时矩阵-向量乘才会走 MMVQ 路径更大的 batch 会落入其他矩阵乘实现如 cuBLAS。在mul_mat_vec_q_cuda_Tmmvq-templates.cuh中ncols_y即 token 数从 1 到 8 分别实例化出独立的 kernelcase 1mul_mat_vec_qtype, 1, nwarps第 403-406 行case 2mul_mat_vec_qtype, 2, nwarps第 407-410 行case 3mul_mat_vec_qtype, 3, nwarps第 411-414 行case 4~8依此类推直到MMVQ_MAX_BATCH_SIZEncols_y 2和ncols_y 3正是 PR #446 描述的处理 2 或 3 个 token 时触发的模板实例。内核内还有一个随ncols_y变化的分支#if defined(GGML_USE_HIPBLAS) defined(__HIP_PLATFORM_AMD__) (defined(RDNA2) || defined(RDNA3)) constexpr int rows_per_cuda_block 1; #else constexpr int rows_per_cuda_block ncols_y 4 ? 1 : 2; #endifmmvq-templates.cuhrows_per_cuda_block表示一个 CUDA block 处理权重矩阵的几行token 数小于 4 时每 block 只处理 1 行否则处理 2 行。从源码结构可以推断ncols_y 2/ncols_y 3恰好落在batch 很小 每 block 单行这条特殊路径上任何针对行数、nrows_dst或结果写回索引dst[j*nrows_dst row0 threadIdx.x]见 mmvq-templates.cuh的边界计算错误都会在此处集中爆发——这与TG 不崩、PP 罕见崩的现象完全自洽。3. 量化类型全覆盖iqk_mmvq.cuik_llama.cpp 的另一特色是大量自研量化类型IQ2_K、IQ3_K、IQ4_K_R4、IQ1_S_R4等。这些类型的 MMVQ 实现在 iqk_mmvq.cu 中通过iqk_mul_mat_vec_q统一分发覆盖了二十余种 IK 特有量化格式mmvq.cu 中这些类型全部落入iqk_mul_mat_vec_q分支。#425 中用户使用的正是 ik_llama.cpp 特有的IQ4_K_R4量化说明该 bug 波及的是整个 MMVQ 家族而非某个特定量化类型。修复与验证一次典型的社区协作 debug 闭环从 issue 时间线可以看到整个修复过程的完整闭环所有记录均来自仓库的 github-data2025-05-07 ~ 05-15三个 issue 陆续提交。作者在 #398 中坦言这更可能是多 GPU 配置下的 bug我只有单 GPU 无法调试#425 中用户补充了NVRM Xid 31的内核日志指向真正的硬件级非法内存访问。社区逐步收窄范围#398 中用户测试发现单 GPU 下-fmoe正常、多 GPU 必崩#389 中发现-ser 7,1是触发器#425 中发现换用 IK 特有混合量化走dequantize - cuBLAS路径无量化 MMVQ 实现就不崩这从反面印证了问题出在量化矩阵乘实现上。2025-05-23ciprianveg在定位中起到关键作用作者发布 #446 修复修改mmvq.cu一行核心逻辑共 5 行变更。2025-05-24PR 合并。此前受困的用户如pt13762104反馈现在一切正常。这一过程也留下了对使用者有价值的经验当llama-server/llama-batched-bench在多 GPU MoE 模型 量化权重组合下出现illegal memory access或GGML_ASSERT崩溃时除了检查显存还应考虑是否命中了 MMVQ 的 batch 边界路径——升级到包含 #446 修复的构建版本是首要动作。结语PR #446 是理解 ik_llama.cpp 这类性能优化型 fork风险与收益的一个绝佳样本为了榨取量化矩阵乘的性能项目维护了覆盖 30 量化类型、按 batch 大小特化的 MMVQ 内核家族MMVQ_MAX_BATCH_SIZE 8而性能优化的复杂度直接转化为了隐藏的边界 bug 概率。该 bug 的隐蔽性在于日常 TG 场景的 batch 永远为 1恰好绕开了故障路径只有 MoE 推理在 PP 阶段把 batch 切成 2~3 个 token 分给某个专家时缺陷才被触发。修复虽小一个文件、5 行但对所有在消费级多 GPU 上运行 Qwen3-MoE、DeepSeek-V3 等模型的用户意义重大。对于希望进一步研究源码的读者建议按以下顺序阅读入口分发 mmvq.cu → 参数结构 mmvq-args.h → 内核模板 mmvq-templates.cuh → IK 量化分发 iqk_mmvq.cu配合 llama-bench.cpp 中-ser/-fmoe等参数的解析逻辑即可完整复原这条 bug 的产生—触发—定位—修复链路。赞分享人工智能大模型推理引擎本地部署模型量化模型优化【免费下载链接】ik_llama.cppllama.cpp fork with additional SOTA quants and improved performance项目地址https://gitcode.com/GitHub_Trending/ik/ik_llama.cpp点击查看免费下载相关推荐ik_llama.cpp CUDA 非法内存访问崩溃全记录:MoE 模型 MMVQ 内核 2/3 行越界 Bug 的排查与修复ik_llama.cpp CUDA 非法内存访问崩溃全记录:MoE 模型 MMVQ 内核 2/3 行越界 Bug 的排查与修复 本文基于 ik_llama.cp人工智能大模型推理引擎本地部署模型量化模型优化扫描文档中的手写签名提取从传统办公到数字化处理的智能解决方案扫描文档中的手写签名提取从传统办公到数字化处理的智能解决方案 在数字化转型浪潮中许多机构仍面临一个看似简单却异常棘手的挑战如何高效地从大量纸质文档中提取手人工智能大模型推理引擎本地部署模型量化模型优化ik_llama.cpp PR 430 解析MoE multi-add 优化的引入、崩溃与回退决策ik_llama.cpp PR 430 解析MoE multi add 优化的引入、崩溃与回退决策 导读 本文以 ik_llama.cpp 仓库中的 PR 4人工智能大模型推理引擎本地部署模型量化模型优化创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
