ik_llama.cpp CUDA 非法内存访问崩溃全记录:MoE 模型 MMVQ 内核 2/3 行越界 Bug 的排查与修复

ik_llama.cpp CUDA 非法内存访问崩溃全记录:MoE 模型 MMVQ 内核 2/3 行越界 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 仓库 issue #425 的完整排查档案,复盘一起典型的 CUDA error: an illegal memory access was encountered 崩溃:多用户在使用 DeepSeek-V3、Qwen3-235B 等 MoE 大模型时,llama-server在处理 prompt 后开始解码时随机崩溃。通过 gdb 调用栈、CUDA 调用追踪和 compute-sanitizer,最终定位到根因并非多卡互联或数据拷贝,而是矩阵-向量乘法(MMVQ)内核在处理 2 或 3 行输入时发生的越界读。读完本文,你将掌握 CUDA 异步错误报错点≠出错点的本质、一套可复用的 GPU 崩溃排查流程,以及 MoE 模型 u-batch、-ot张量覆盖与图切分(graph splits)调优的实战经验。一、故障现象:解码启动瞬间的非法内存访问issue 由用户 nux 于 2025-05-15 提交。使用单张 GPU 运行 DeepSeek-V3 IQ4_K_R4 量化模型时,llama-server在 prompt 处理完成、进入解码阶段后崩溃:INFO [update_slots] kv cache rm [p0, end) | id_slot0 id_task3 p00 CUDA error: an illegal memory access was encountered current device: 0, in function ggml_backend_cuda_synchronize at .../ggml/src/ggml-cuda.cu:3067 cudaStreamSynchronize(cuda_ctx-stream()) .../ggml/src/ggml-cuda.cu:110: CUDA error内核日志中还能看到 NVIDIA 驱动的 MMU 故障记录,进一步佐证是 GPU 侧的非法虚拟地址读取:NVRM: Xid (PCI:0000:01:00): 31, pid80798, namellama-server, ... MMU Fault: ENGINE GRAPHICS GPC1 GPCCLIENT_T1_3 faulted 0x7e9f_4f200000. Fault is of type FAULT_PDE ACCESS_TYPE_VIRT_READ出错的原始命令(DeepSeek-V3 MLA 融合 MoE CPU 专家覆盖):llama-server --model DeepSeek-V3-0324-IQ4_K_R4.gguf \ --ctx-size 32768 -mla 2 -fa -amb 512 -fmoe \ --n-gpu-layers 63 --override-tensor expsCPU \ --parallel 1 --threads 32 --host 0.0.0.0 --port 8081后续又有两位用户报告了几乎一致的崩溃:ciprianveg:2×RTX-3090 2×A4000 四卡混合,运行 Qwen3-235B-A22B UD 量化模型,llama-sweep-bench完全正常,但通过 open-webui 发送真实 prompt 即崩;Lissanro:4×RTX-3090 EPYC 7763,运行 DeepSeek-R1T-Chimera IQ4_K_R4,崩溃呈周期性出现,且从未在第一次请求时发生,通常在重新生成(regenerate)消息时触发。二、排除法:模型、参数、驱动都无辜排查初期维护者 ikawrakow 与用户共同做了大量交叉实验,逐步排除了多个嫌疑人:实验结果排除对象nux 去掉-fmoe再跑仍然崩溃融合 MoE 路径nux 设-n-gpu-layers 0全 CPU 运行正常—(说明纯 CPU 不受影响)nux 设-n-gpu-layers 1仅 1 层在 GPU崩溃并非需要大量层在 GPU 才触发nux 用同样的 unsloth UD-Q4_K_XL 模型跑原版 llama.cpp正常(24.18 tok/s PP)模型文件本身nux 换 Qwen3-30B-A3B-Q4_K_M 跑 ik_llama.cpp正常并非所有模型都触发ciprianveg 回退到 4 天前的 main 提交0c57f84重建仍然崩溃近期回归ciprianveg 把 4 张 GPU 中任意 3 张组合轮换测试3 卡均可运行,4 卡崩溃GPU 硬件故障ciprianveg 在 sweep-bench 中加-ub 873、-b 1234等非整除 batch正常u-batch 边界 bugLissanro 更换 prompt(去掉某行 PHP 正则代码)、调整 thinking 模式崩溃概率无确定变化prompt 内容其中 nux 还注意到一个可疑细节:某次把 prompt 中一行 PHP 正则代码删掉后没崩,但单独发送该正则又不崩——最终被证明只是时间上的巧合。这一阶段的结论(维护者原话):it is the bug that happens with multiple GPUs and partial offload,即多卡 部分 offload 场景,且维护者本人单机环境无法复现,排查一度陷入僵局。三、用 gdb 追到调度器:崩溃发生在张量拷贝阶段维护者要求用户在RelWithDebInfo构建下用 gdb 抓取崩溃现场:cmake --build ./build --config RelWithDebInfo -j $(nproc) gdb --args ./build/bin/llama-server your_commandnux 提供的调用栈还原了崩溃时的执行路径,自底向上为:server_queue::start_loop → server_context::update_slots → llama_decode → llama_decode_internal → llama_graph_compute → ggml_backend_sched_graph_compute_async → ggml_backend_sched_compute_splits ← 崩溃帧 → ggml_backend_synchronize → ggml_backend_cuda_synchronize → cudaStreamSynchronize → 非法内存访问其中ggml_backend_sched_compute_splits是后端调度器按图切分(graph split)在多后端间搬运张量的核心函数,当前仓库中可对照 ggml_backend_sched_compute_splits。在崩溃帧内用p *input查看正在处理的张量,得到关键信息:(gdb) p *input $1 {type GGML_TYPE_F32, backend GGML_BACKEND_TYPE_CPU, ..., name ffn_moe_weighted-60, ...}即崩溃时调度器正在处理第 60 层的 MoE 加权求和张量(其前驱输出为l_out-42之类的层间结果)。对 ciprianveg 的四卡环境,进一步打印sched-n_splits(值为 93)、各 split 的输入张量,确认输入张量inp_pos、KQ_mask (copy)此前已成功拷贝 42 次,嫌疑集中到层间结果上。但仅凭调用栈仍无法定论,因为报错位置只是同步点,不是真正的出错点——这一点是破局的关键。四、核心认知:CUDA 错误的异步性维护者在结案时给出了整个排查中最有教育价值的一段解释:I realized only yesterday that checking for an error after launching a CUDA kernel does not tell us that the kernel was successfully executed, but only tells us that the kernel was successfullyqueuedfor execution. If there is a bug in the kernel (e.g., illegal memory access), the resulting error will get reported in some later call.也就是说:cudaLaunchKernel成功只表示内核被成功入队,不代表内核已执行完毕;内核内部的非法访存,其错误会在后续任意 CUDA 调用(通常是cudaStreamSynchronize、下一次拷贝)才浮出水面;因此所有崩溃日志都看起来发生在 ggml_cuda_error 报告的ggml_backend_cuda_synchronize/ggml_backend_sched_compute_splits处——这正是前几轮排查被误导去怀疑后端、设备间拷贝、peer-to-peer 访问的原因。这也解释了为什么崩溃随机:错误何时被上报,取决于崩溃内核与下一次同步调用之间的时序,于是同一台机器上有时崩、有时不崩。五、compute-sanitizer 实锤:MMVQ 内核越界读 2 字节确认设备间无法开启 peer access(日志中反复出现Failed to enable peer access ... not supported between these two devices)后,维护者建议改用 CUDA 官方 sanitizer 直接检查内核访存。ciprianveg 的compute-sanitizer输出一次性暴露了真相: Invalid __global__ read of size 2 bytes at void mul_mat_vec_q(ggml_type)12, (int)2, (int)4(...) by thread (8,0,0) in block (779,0,0) Address 0x7f788a3f7a0c is out of bounds and is 50,701 bytes after the nearest allocation of size 18,224,165,888 bytes ... ERROR SUMMARY: 1263 errors模板实参(ggml_type)12, (int)2, (int)4中,第二个参数ncols_y 2正是本次内核处理的行数(token 数)。结合此前临时加入的张量元数据检查(ggml_cuda_up_gate_unary打印的blk.42.ffn_up_exps.weight, q4_K, 4096 x 1536 x 128等),元数据本身是正确的——越界读是内核自身的寻址问题。六、根因:MMVQ 内核在处理 2 或 3 行时越界维护者最终确认(2025-05-23):The bug was in the matrix-vector multiplication kernel. It only shows up when the number of rows being processed (i.e., tokens) is 2 or 3 (the matrix-vector kernel confusingly processes up to 8 rows). This is not used during TG, and only triggers if an expert ends up with 2 or 3 rows, which is rare.翻译成工程细节,可对照当前仓库的 ggml/src/ggml-cuda/mmvq-templates.cuh 理解:MMVQ(matrix-vector, quantized)内核按ncols_y(待处理的 token 行数)模板化,支持 1~8 行,入口处有GGML_ASSERT(args.ncols_y MMVQ_MAX_BATCH_SIZE)约束(mul_mat_vec_q_cuda_T);内核主体 k_mul_mat_vec_q 中rows_per_cuda_block ncols_y 4 ? 1 : 2(L86),即 1~3 行时每个 CUDA block 处理 1 行;修复前,当ncols_y恰为 2 或 3 时,内核内部的行寻址会越过激活缓冲区边界,产生 sanitizer 所报告的小幅越界读(2 字节);触发条件是某个 MoE 专家恰好只分到 2 或 3 个 token。DeepSeek-V3/R1 每个 token 激活 8 个专家,235B 模型有 128 个专家,MoE 路由下个别专家拿到 2/3 行虽然rare,但在长对话、KV cache 部分匹配、重新生成消息等场景下会周期性出现——这与 Lissanro 从未第一次崩溃、regenerate 时高发 的观察完全吻合;该路径只在prompt 处理的分批解码(一次处理少量 token 的mul_mat_id)中走 MMVQ 内核,正常逐 token 的 TG(token generation)不经过它,所以llama-sweep-bench和llama-cli一直表现正常。修复是 PR #442 中一个独立的 Fix bug in MMVQ kernel 提交(提交号b79be8a),用户 Lissanro 验证:只应用这一个提交即可消除崩溃,无需 #442 中的其他追踪/补丁改动。随后该修复随主线发布,issue 关闭。#442 分支中的其余内容(CUDA 调用追踪器、设备切换日志等)属于调试工具,不是修复本体。七、排查过程中沉淀的调优经验这场事故虽然根因在内核,但排查过程同时验证/产出了多条 ik_llama.cpp 多卡 MoE 部署的实战经验,均出自 issue 讨论:1. u-batch 对 MoE 的 PP 速度是杠杆,对 dense 模型不是。MoE 模型下,u-batch512 时 DeepSeek-V3/R1 每个 ubatch 共激活512×84096个专家,摊到 256 个专家上平均每个专家只有 16 行,矩阵乘效率极低;增大 u-batch 可显著提升 PP 速度而 TG 几乎不受损。维护者建议 MoE 场景直接-b 4096 -ub 4096并不要超过 4096(8192 曾暴露另一类 bug)。Lissanro 实测从 35 tok/s 提到 100~105 tok/s(ciprianveg 从 80 提到 180 tok/s)。dense 模型则相反,默认 batch2048、u-batch512 已接近最优。2. 大 u-batch 吃 CUDA compute buffer,用-DGGML_SCHED_MAX_COPIES1省显存。编译期加-DGGML_SCHED_MAX_COPIES1可显著缩小调度器为每个切分预留的拷贝缓冲,腾出的显存足以支持 4096 的 u-batch 和更长的 q8_0 KV cache。3.-ot覆盖正则要写精确,多余切分会拖慢 TG。ciprianveg 的-ot blk.(?:[x]|[5-9][0-9]).ffn.*CPU会把ffn_gate_inp、ffn_norm等小张量也留在 CPU,制造大量图切分。维护者用 Qwen3-30B-A3B 单卡 RTX 4080 测得:把正则收紧为只覆盖.*_exps后,图切分从 74 个降到 38 个,TG 从 66.7 提到 70.4 t/s。4. 专家放置与-mla的取舍。对 DeepSeek/Qwen3 这类大专家数 MoE,大 u-batch 下把三类专家(ffn_up_exps/ffn_gate_exps/ffn_down_exps)一起 offload 到 GPU 更优;对 Llama-4 Maverick(128 专家仅激活 1 个),offload 到 GPU 反而比纯 CPU 慢;若显存不够跑大 u-batch,可每卡少放一层换 PP 速度,TG 损失温和;-mla 3是当时新支持的模式,省掉 FA 内核中的两次矩阵乘,长上下文 TG 更优(nux 最初用-mla 2)。5. 设备间拷贝告警不是病因。attempt to copy from device X to device Y without access enabled这类告警在无法 P2P 的主机上属正常回退路径(数据经 CPU 中转),ciprianveg 的四卡正是这种拓扑;真正的主线 llama.cpp 里也只有在 split-mode row 下才显式开关 peer access,这里的检查并非必需。八、给使用者的排查清单结合本 issue,遇到 CUDA error: an illegal memory access was encountered 时,建议按以下顺序操作(全部为仓库内可执行的查看/构建步骤):看报错函数而非报错内容:ggml_backend_cuda_synchronize之类的同步函数只是报案地点。错误来自之前入队的某个内核,优先怀疑上一批计算而非拷贝本身。RelWithDebInfo gdb 抓栈,定位崩溃帧,再frame Np *input等打印张量元数据(type、ne、nb、name),确认张量名与所在层,判断属于 attention、MLA 还是 MoE 路径。用 compute-sanitizer 抓内核级越界:compute-sanitizer ./build/bin/llama-server ...,输出会直接给出内核名、模板实参与越界字节数——本案中一行mul_mat_vec_q(ggml_type)12, (int)2, (int)4就锁定了2 行这个稀有分支。MoE 模型注意 2/3 行分支:任何按 token 数模板化的内核(本例 MMVQ 支持 1~8 行)在行数落在小整数上时才越界,表现为周期性、regenerate 时高发、与上下文长短无关。控制变量:保持模型不变,切换 0/1/全部 GPU 层、开关-fmoe、更换同架构小模型,可快速区分后端路径问题与模型数据问题。本次事故的完整档案(含全部调用栈、sanitizer 输出与逐轮实验)见 issue #425;修复后的内核实现位于 ggml/src/ggml-cuda/mmvq-templates.cuh,错误上报路径位于 ggml/src/ggml-cuda.cu,调度切分逻辑位于 ggml/src/ggml-backend.cpp。【免费下载链接】ik_llama.cppllama.cpp fork with additional SOTA quants and improved performance项目地址: https://gitcode.com/GitHub_Trending/ik/ik_llama.cpp创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考