llama.cpp b11178合并MUSA补丁：修复MTT S5000算子失败与构建问题
该版本针对摩尔线程MTT S5000（PH1，mp_31，MUSA SDK 5.2.0）修复FlashAttention拷贝宽度、CUB路径、MMQ误用、TOPK_MOE挂起等问题，并将MUSA CI与文档默认目标从MTT S80迁至PH1。
AI解读：这条更新解决的是llama.cpp在摩尔线程MTT S5000上“跑得起来但跑不对”的问题。此前MUSA后端把FlashAttention的每线程拷贝宽度算成 8 字节，而sm_70+的CUDA目标是 16 字节，等于每次搬运只干一半的活；同时CUB路径因为CUDART_VERSION在MUSA SDK里根本没定义而被整体编译掉。
对使用S5000跑本地推理的人来说，最直接的变化是量化矩阵乘法不再走错误的MMQ路径。测试显示PH1上 260 个失败用例全部集中在MMQ，改用cuBLAS/dequant路径后，27B模型默认参数下的困惑度从nan变为有限值，全套 22237 个用例 0 失败——之前是数值直接崩掉，现在至少结果是对的。
几个挂起和超时也修了：融合TOPK_MOE内核在行数不是 8 的倍数时，提前返回的warp会卡住块级屏障，导致内核挂死、启动超时，现在改为让这些warp钳位到最后一行并重新计算。GATED_DELTA_NET也重新启用，此前该算子被强制丢给CPU，每次前向有 48 个节点，在S5000上关闭FA时pp512从 964.51 提升到 2119.26 t/s。
另一个实际收益是构建门槛降低：CI任务、ci/run.sh默认值和构建文档的目标从第一代MTT S80（MUSA_ARCHITECTURES=21）改为PH1（31），并使用PH1的MUSA SDK devel镜像。该镜像需要额外装python3-venv给ccache-buckets步骤用，且构建命令不能再带time前缀，因为容器任务用sh执行且镜像没有time二进制。
需要留意的是作者明确说MMQ本身的缺陷没有在这里修，只是绕开；实验性的GGML_MUSA_GRAPHS选项仍然默认关闭，因为在S5000上实测捕获路径更慢（pp512 693对 772 t/s），且边界MUL_MAT用例在两次运行间无法复现。FA关闭时的性能基本不变，FLASH_ATTN_EXT、困惑度等指标在多个补丁里都报告为一致或逐位相同。
llama.cpp发布b11178，其中包含公关 #29193 的MUSA后端补丁，集中修复PH1（MTT S5000）上的算子失败与构建问题。
补丁针对MUSA SDK 5.2.0和mp_31验证，覆盖FlashAttention拷贝宽度、CUB路径启用、MMQ路径误用、融合TOPK_MOE内核挂起、GATED_DELTA_NET启用以及CI与文档目标迁移。
FlashAttention拷贝宽度与CUB路径
ggml_cuda_get_max_cpy_bytes() 依据 __CUDA_ARCH__推导拷贝宽度，但mcc从不定义它，MUSA落入通用分支返回 8 字节，而所有sm_70+目标拿到的是 16 字节。该值决定FlashAttention K/V暂存代码（fattn-common、fattn-vec、fattn-tile、fattn-mma-f16共享内存加载）以及mmq-vec-dot的每线程拷贝单元，此前每个MUSA FlashAttention内核对每条指令搬运的字节数只有一半。
在MTT S5000（mp_31，MUSA SDK 5.2.0）上，Qwen3.8-27B-UD-Q4_K_M、-ngl 999、-p 512 -n 64、-fa on条件下，预填充从 751.15 提升到 794.73 t/s，解码从 15.59 提升到 15.69 t/s；-fa off不变（1050.05 到 1052.86 t/s），FLASH_ATTN_EXT不变（3984 ok / 0 fail / 1323 unsupported），困惑度不变。
CUB路径方面，GGML_CUDA_USE_CUB和USE_CUB由CUDART_VERSION >= 11070 选择，而MUSA SDK从不满足该条件：/usr/local/musa/include下未定义CUDART_VERSION，导致所有基于CUB的路径在MUSA上被编译掉，尽管该SDK自带CUB且内核可为mp_31构建。补丁改为同时依据GGML_USE_MUSA选择，设备级算法也可用：cub::DeviceSegmentedSort在mp_31上可编译并产生正确结果。这使ne[0] 从 100 ok / 0提升到CUDA对等水平，TOP_K从 0 ok / 354 not supported变为 527 ok / 0。其余 20 个逐算子测试套件不变，Qwen3-0.6B f16（14.4679）和Qwen3.8-27B iq4_nl（5.1724）困惑度不变。
MMQ在PH1上被绕开
在MTT S5000（mp_31，MUSA SDK 5.2.0）上运行test-backend-ops失败 260 个用例，全部走MMQ路径：批处理src1的MUL_MAT（任意bs/nr不等于 [1,1]）109 个用例失败，涵盖所有量化类型，例如n=16 时 13 个用例中失败 12 个，而普通 [1,1] 布局通过；所有量化MUL_MAT_ID失败 147 个，而相同形状的f16/f32变体通过；超过约 512 token的MUL_MAT失败 4 个（n=509 到 4096），小n用例通过。
cuBLAS/dequant路径对所有这些用例都正确，小批量使用的MMVQ路径不受影响，因此PH1上的量化矩阵乘法现在改走该路径，而不是返回错误值。默认参数下 27B困惑度从nan变为有限值，全套测试报告 22237 个用例 0 失败。补丁明确表示MMQ本身的缺陷（fastdiv、__umulhi、uint3内核参数和基于 __CUDA_ARCH__的MMA可用性均已检查且在该部件上正确）不在这里解决。
TOPK_MOE挂起与GATED_DELTA_NET启用
topk_moe_cuda对超出图尾的行提前返回，但一个块覆盖TOPK_MOE_ROWS_PER_BLOCK（8）行，因此当n_rows不是 8 的倍数时最后一个块只有部分填充。在MUSA上，已经返回的warp会阻塞下方的块级 __syncthreads()，导致内核挂起且启动超时；CUDA容忍已退出的warp，所以CUDA数据从未暴露该问题。补丁对MUSA将这些warp的行索引钳位到最后一行，使块内每个warp都到达屏障，它们重算最后一行并写入相同的值，CUDA代码路径不变。在MTT S5000（mp_31）上，融合TOPK_MOE用例从启动超时且无完成用例，变为 418 ok / 0 not supported / 0 failed，即与CUDA结果一致，其余 101 个逐算子套件不变（0 failed，计数无变化）。
GATED_DELTA_NET此前因mcc当时无法构建该内核而对所有MUSA目标关闭。当前工具链可以构建：mp_31与MUSA SDK 5.2.0下文件编译零错误，全部 36 个test-backend-ops GATED_DELTA_NET用例通过并对照CPU参考，27B困惑度不变。该算子被拒绝时，调度器只能将其放到CPU上执行，每次前向有 48 个GATED_DELTA_NET节点。在MTT S5000（Qwen3.8-27B Q4_K_M，-ngl 999，单容器，-r 3）上：pp512（FA off）从 964.51 提升到 2119.26 t/s，tg64（FA off）从 10.15 提升到 15.50 t/s。
CI与文档目标迁至PH1
MUSA CI任务和文档默认值仍指向第一代（MTT S80，MUSA_ARCHITECTURES=21），而当前MUSA SDK面向PH1（MTT S5000，31）。补丁将任务、ci/run.sh默认值和构建文档改为 31，并在PH1 MUSA SDK devel镜像中运行任务：registry.mthreads.com/mcconline/inference/pytorch:2.9.1.post1-py3.10-musa5.2.0-mp31-devel-ubuntu22.04-amd64。
该镜像需要两项此前没有的东西：python3-venv用于ccache-buckets步骤（该步骤为Hugging Face CLI构建虚拟环境），以及构建命令不加time前缀，因为容器任务用sh执行步骤且镜像不含time二进制。
其他运行时路径调整与实测数据
补丁还移除了一些过时的编译期保护：设备初始化改为查询cudaDevAttrCooperativeLaunch而非硬编码false，设备报告cooperativeLaunch=1 且musaLaunchCooperativeKernel可用；保留prop.warpSize而非覆盖为 32（设备本就报告 32，仅消除分歧）；CUDA_SET_SHARED_MEMORY_LIMIT和FA共享内存提升现在可用，musaFuncSetAttribute返回成功且sharedMemPerBlockOptin为 192 KiB，内核可使用超过默认 48 KiB；vendors/musa.h增加设备初始化所需的cudaDeviceGetAttribute和cudaDevAttrCooperativeLaunch映射。
在单张S5000上以Qwen3.8-27B Q4_K_M（-ngl 999，-r 3）实测：pp512从 968.27 变为 957.09 t/s，tg64从 10.09 变为 10.23 t/s，FLASH_ATTN_EXT扫描相同（3975/3982 两者一致），困惑度相同（80.2841 +/- 7.26772）。
另一组改动移除MUSA运行时已覆盖的编译期保护：mcc从不定义 __CUDA_ARCH__，因此这些按架构分流的回退在MUSA上本就已采用，叠加的GGML_USE_MUSA保护只是让上游文本不被编译。涉及wkv.cu中无效的 #pragma unroll抑制、common.cuh中MUSA专用的 __builtin_unreachable()、ssm-scan.cu中由GGML_CUDA_CC_IS_NVIDIA(cc) 和turing_mma_available(cc) 运行时把关的SSD（Mamba-2预填充）块（两者对PH1即cc 0x100310均为false），以及common.cuh中与warp_reduce_sum(half2) 同样受FP16_AVAILABLE保护的warp_reduce_max(half2)（MUSA专用保护使其没有return语句，目前无调用方）。
MTT S5000（mp_31，MUSA SDK 5.2.0），MUSA_ARCHITECTURES=31：构建rc=0。与同一代码树在同一张卡上的未修改版本相比，FLASH_ATTN_EXT（3984 ok / 0 fail / 1323 unsupported）、SSM_SCAN（15/0）、RWKV_WKV6（6/0）、GATED_DELTA_NET（38/0）和MUL_MAT（1299/0/385 unsupported）结果相同，-fa on时困惑度逐位相同（5.1639 +/- 0.36673，4 chunks）。
实验性的GGML_MUSA_GRAPHS选项仍默认关闭：argsort.cu和mean.cu在USE_CUDA_GRAPH块内调用cudaStreamCaptureStatus、cudaStreamIsCapturing和cudaStreamCaptureStatusNone，但MUSA兼容头没有这些名字的别名，开启该选项会在这两个文件产生 7 个构建错误。补丁在启用图代码的相同保护下将这三个名字映射到musa* 对应项，默认构建不受影响。在MTT S5000上实测捕获路径更慢（两个会话中pp512 693对 772 t/s，tg128 15.20对 15.39 t/s），且边界MUL_MAT用例在运行间无法复现。