Eval bug: CUDA MoE MMQ illegal memory access at ubatch 512 (src1 padding uses ne11 instead of gathered columns)

快速结论：这个报错通常出现在 llama.cpp 使用 CUDA 后端、以较大 ubatch（例如 512）对 MoE 模型做 MUL_MAT_ID 推理时，报错原文为 Eval bug: CUDA MoE MMQ illegal memory access at ubatch 512 (src1 padding uses ne11 instead of gathered columns)；优先排查 MoE + MMQ 量化路径下 src1 padding 是否按 gathered columns 计算，而不是按 ne11。

适用环境：llama.cpp 0.5.0-dev（build 3155，commit 5fc4f3c8c），Linux（Arch），CUDA 后端，2× NVIDIA GeForce RTX 5060 Ti 16 GB（Blackwell，sm_120），NVIDIA 驱动 610.57.04，CUDA 13.3，AMD Ryzen 5 5600X / 62 GB DDR4。

最快修复方案：暂无确认的一步修复方案。Issue 中作者通过修改 ggml/src/ggml-cuda/mmq.cu，把 ggml_cuda_mmq_get_J_max 的最后一个参数从 ne11 改为 ne_get_rows 后验证通过；该改动随后指向 PR #29941，但 Issue 本身标注为 bug-unconfirmed，应用前需要自行评估。

注意事项：该补丁属于用户自行验证的改动，作者也说明“不确定这是否是正确解法”，并非领域内确认的最终方案。是否已被上游合并、是否适用于其他 MoE 架构、其他量化类型或其他 CUDA 版本，Issue 中没有给出完整结论。

问题场景

用户在 Linux 上使用 llama.cpp 的 CUDA 后端运行 MoE 模型（Swift-1.5-Qwen3.8-Flash-Next，架构 qwen4exp，512 experts，每 token 使用 10 个 expert，GGUF 量化 IQ3_XXS），通过 llama-server -ub 512 启动。先发送一条长输出请求，再发送一条 prompt 超过 512 token 的请求，第二条请求在 prompt processing 阶段触发 MUL_MAT_ID 失败并出现 CUDA illegal memory access。也可以直接通过 tests/test-backend-ops.cpp 中新增的 MUL_MAT_ID 测试用例复现。

报错原文

Eval bug: CUDA MoE MMQ illegal memory access at ubatch 512 (src1 padding uses ne11 instead of gathered columns)

========= Invalid __global__ read of size 4 bytes
=========     at void mul_mat_q_process_tile<(ggml_type)17, (int)128, (bool)0, (bool)0, (ggml_prec)30>(...)+0x7f00 in mmq.cuh:940
=========     by thread (24,7,0) in block (2032,0,0)
=========     Access to 0xb00e00000 is out of bounds
=========     and is 1 bytes after the nearest allocation at 0xb00200000 of size 12582912 bytes
========= ERROR SUMMARY: 429 errors

原因分析

可能原因是 MoE MMQ 路径中，src1 的 padding 大小按 ne11（对 MoE 而言为 1 或 n_expert_used）计算，而不是按实际 gathered columns（ne_get_rows）计算。当 gathered column 数不是 2 的幂、且最后一个 tile 会读取超过已 gather 行数最多 J 列时，读取会越过分配边界，触发 CUDA illegal memory access。Issue 作者怀疑首个坏提交可能是 6eddde06a（#24519）。

环境排查

确认 llama.cpp 版本与 commit：./llama-cli --version，Issue 中为 0.5.0-dev（build 3155，commit 5fc4f3c8c）。
确认 GGML 后端为 CUDA，并确认是否启用了 MMQ 路径。
确认 GPU 型号与计算能力：Issue 中为 2× RTX 5060 Ti 16 GB（Blackwell，sm_120）。
确认 NVIDIA 驱动与 CUDA 版本：驱动 610.57.04，CUDA 13.3。
确认模型架构与量化：qwen4exp、512 experts、10 used per token、IQ3_XXS。
确认启动参数是否使用 -ub 512；Issue 中 -ub 256 可作为未打补丁时的临时规避。
确认复现请求序列：先一条长输出请求，再一条 prompt 长度超过 512 token 的请求。

解决步骤

先按 Issue 中的复现步骤确认问题：使用 Qwen3.8-Flash-Next GGUF 模型，以 llama-server -ub 512 启动，先发长输出请求，再发超过 512 token 的 prompt，观察是否在 prompt processing 阶段出现 MUL_MAT_ID 失败和 CUDA illegal memory access。
可优先尝试的验证方式：在 tests/test-backend-ops.cpp 中加入 Issue 提供的 MUL_MAT_ID 测试用例：test_mul_mat_id(GGML_TYPE_Q4_0, GGML_TYPE_F32, 512, 10, b, 640, 508, 2560)，其中 b 取 false 和 true。
如果测试稳定复现非法内存访问，可优先尝试 Issue 中作者验证过的补丁：修改 ggml/src/ggml-cuda/mmq.cu，将 nbytes_src1_q8_1 计算中的 ggml_cuda_mmq_get_J_max(src0->type, fallback, cc, ne11) 改为 ggml_cuda_mmq_get_J_max(src0->type, fallback, cc, ne_get_rows)，使 padding 按 gathered columns 计算。
应用补丁后重新编译 llama.cpp，并重新运行 test-backend-ops 中的 MUL_MAT_ID 与 MUL_MAT 测试。
在真实模型上重新运行 llama-server -ub 512，按同样请求序列验证是否不再出现 CUDA 错误。
如果未采用补丁，可使用 -ub 256 作为未打补丁时的临时规避方案；Issue 中该规避可避免触发，但吞吐低于 -ub 512 打补丁后的结果。

验证方法

可通过以下证据确认：test-backend-ops 中新增的 MUL_MAT_ID 用例在补丁前后表现不同（无补丁时 b=0 / b=1 均出现 illegal memory access，打补丁后均 OK）；compute-sanitizer memcheck 在相同 shape 下错误数从 429 / 201 降为 0；完整 test-backend-ops 中 MUL_MAT_ID 964/964 通过、MUL_MAT 1304/1304 通过；真实模型上连续 3 条请求不再出现 HTTP 500 或 CUDA error，且 8K-token prefill 在 -ub 512 下可正常完成。