Eval bug: ~20% slower MoE prompt processing on RDNA1 (gfx1010) since #29948 (MMQ grid order change)

快速结论：在 RDNA1（gfx1010）上使用 ROCm/HIP 后端运行 MoE 模型且设置 -ot exps=CPU 时，自 #29948 起 prompt processing（pp512）约下降 20%，而 token generation 不受影响。优先排查 MMQ 的 grid 顺序：该提交把 launch_mul_mat_q 的 tiling grid 从 (nty, ntx, ntzw) 改成 (ntx, nty, ntzw)，这一变化在 HIP 上引入了性能回退。

适用环境：llama.cpp 0.8.0-dev（build 207，commit 713319e28，对应 b12066），Linux（Ubuntu 24.04，kernel 6.8），AMD Ryzen 9 3900X + Radeon RX 5700 XT 8 GB（gfx1010，RDNA1），ROCm 7.2.4，cmake 参数为 -DGGML_HIP=ON -DGPU_TARGETS=gfx1010 -DGGML_HIP_NO_VMM=ON -DCMAKE_BUILD_TYPE=Release，使用 GNU 13.3.0 编译。模型为 Qwen3.6-35B-A3B Q4_K_M（MoE，对照组为全 GPU 的 dense 模型 Qwen3.5-9B Q6_K）。

最快修复方案：暂无确认的一步修复方案。Issue 上报者使用的本地 workaround 是：在 ggml/src/ggml-cuda/mmq.cuh 中，仅在 GGML_USE_HIP 下恢复 #29948 之前的 grid 顺序，并将 kernel 内 blockIdx.x/y 到 it/jt 的映射同步改回。实测 pp512 从 180.8 t/s 恢复到 223.7 t/s，tg 不变。

注意事项：该补丁属于本地 workaround，尚未合并，也未经上游验证；仅由原始 Issue 作者在 gfx1010 上测试过，其他 AMD 架构是否受影响尚不明确，因此应作为“可优先尝试”方案。补丁在 HIP 范围内生效，并未只限定 RDNA1；#29948 引入的 fused gate/up/GLU kernel 本身是 NVIDIA 专用的。token generation 与全 GPU 的 dense 模型不受此回退影响，验证时应重点比较 MoE 模型的 pp512。

问题场景

用户在 gfx1010（RDNA1）上通过 HIP/ROCm 后端运行 llama.cpp，测试 MoE 模型 Qwen3.6-35B-A3B Q4_K_M，并设置 -ngl 99 -ot exps=CPU -fa 1，让专家权重留在 CPU 上、其余算子卸载到 GPU。使用 llama-bench 构造 prompt 处理负载（pp512）时，发现自 #29948 之后 prompt processing 性能明显下降。用 -nopo 1（不 offload 算子）后两个版本性能一致，说明回退出在 op offload 下的 GPU 专家 matmul 上。对照组使用全 GPU 的 dense 模型 Qwen3.5-9B Q6_K。

报错原文

Eval bug: ~20% slower MoE prompt processing on RDNA1 (gfx1010) since #29948 (MMQ grid order change)

原因分析

最可能的原因：#29948（commit c49edf096，“CUDA: fuse FFN gate/up matmuls and GLU in MMQ with interleaved gate/up rows”）修改了 launch_mul_mat_q 中的 MMQ tiling grid，由 (nty, ntx, ntzw) 改为 (ntx, nty, ntzw)（column tiles innermost），以让并发 block 共享同一块 x 的 tile 并从 L2 读取。但该提交真正要服务的 fused gate/up kernel 是 NVIDIA-only，而这次 grid 顺序变更作用于所有 MMQ kernel，因此在 HIP/RDNA1 上对 MoE 专家 matmul 产生约 20% 的 pp512 回退。原始 Issue 作者通过恢复旧 grid 顺序在 GGML_USE_HIP 下复现了性能回升；至于其他 AMD 架构是否同样受影响，目前没有证据。

环境排查

确认 GPU 架构是否为 gfx1010（RDNA1），以及 ROCm 版本（Issue 中为 7.2.4）。
确认 llama.cpp 版本与提交：回退是否出现在 b12066 / 713319e28 之后，以及 bad commit 是否落在 c49edf096（#29948）附近。
确认构建参数包含 -DGGML_HIP=ON 与 -DGPU_TARGETS=gfx1010，即确实走 HIP/ROCm 后端而非 CUDA。
确认测试命令与 Issue 一致（-ngl 99 -ot exps=CPU -fa 1 -p 512 -n 64 -r 3），这样 pp512 才具备可比性。
确认回退仅出现在 MoE prompt processing：对比 tg64 是否不变，以及全 GPU dense 模型是否不变。
可用 -nopo 1 和 GGML_CUDA_DISABLE_FUSION=1 辅助判断问题是否位于 op offload 下的 GPU 专家 matmul；Issue 中这两项都未改变结果。

解决步骤

先复现基线：在 gfx1010 上运行 llama-bench -m Qwen3.6-35B-A3B-Q4_K_M.gguf -ngl 99 -ot exps=CPU -fa 1 -t 12 -p 512 -n 64 -r 3，确认 pp512 是否明显低于 #29948 之前的版本。
确认 token generation（tg64）与全 GPU dense 模型未变，把排查范围收窄到 MoE 的 op offload matmul。
“可优先尝试”：在 ggml/src/ggml-cuda/mmq.cuh 的 mul_mat_q kernel 中，仅在 GGML_USE_HIP 下把 jt/it 取回 blockIdx.y/blockIdx.x；非 HIP 仍保持 #29948 的 blockIdx.x/blockIdx.y。
“可优先尝试”：在 launch_mul_mat_q 中，仅在 GGML_USE_HIP 下把 block_nums_xy_tiling 恢复为 (nty, ntx, ntzw)，非 HIP 保持 (ntx, nty, ntzw)。两处改动需同时进行，保证 grid 布局与 kernel 内映射一致。
重新编译后再次执行第 1 步的 llama-bench 命令，与未打补丁的 b12066 及 #29948 之前的构建对比。

验证方法

在同一 gfx1010 机器、同一模型与同一 bench 参数下，打完补丁后 MoE 模型的 pp512 应基本回到 #29948 之前的水平（Issue 中为 180.8 t/s → 223.7 t/s，约 224 t/s），tg64 保持约 36.6 t/s 不变，全 GPU dense 模型 pp512 保持约 189 t/s 不受影响。可同时用 -nopo 1 做交叉验证：不 offload 算子时应无明显差异。