vLLM-MUSA推理部署实战

从开发快照切换到正式 Release 基线

vLLM-MUSA v0.24.0 于 2026-09-04 正式发布，对齐上游 vLLM 0.24.0，并迁移到 PyTorch 2.11.x MUSA 栈。8月历史章节中的开发分支、PR 状态和带日期镜像只代表当时快照，不再作为本节推荐基线。v0.24.0 正式发布

registry.mthreads.com/mcconline/inference/vllm/vllm-openai:v0.24.0

上述为 GitHub Release 明确指定的“最终发布镜像”。官方文档另列带 amd64-ph1-musa5.2.0 后缀的标签；在未校验 registry digest 之前，不宣称两者是相同镜像。生产验收建议锁定实际拉取的 digest，并保留回滚镜像。最终镜像说明；S5000 文档镜像

该表是 S5000 发布栈，不是 S80、M1000 或所有 GPU 的通用兼容矩阵。依赖与 serving cookbook

如何验证新能力，而不堆叠开关

正式说明包含 FP8 E4M3 KV cache、V1 / ModelRunnerV2 与图编译改进、Qwen3.5 shared-expert 融合、DeepSeek-V4-Flash-Base sparse MLA 和推测解码等更新。是否启用及其参数须匹配具体模型的 cookbook，不能将不同模型命令拼接。主要更新

1. 选择官方 cookbook 中与硬件、模型相符的一组配置，固定权重 revision、TP 数和设备可见性。
2. 先使用该配置完成健康检查及一条短请求，记录镜像 digest、完整启动参数和输出。
3. 用约 4K 输入 / 1K 输出及有界并发复现 cookbook 的负载形态，再测试自己的业务长度分布；该形态不是性能保证。
4. 分别测首 token 时延、单请求生成速率、总吞吐、显存、失败率与输出正确性；前缀缓存冷热场景分开。
5. 在满足官方模型配置和显存安全约束后，再做单变量性能对照。官方要求联动调整的参数作为完整配置方案变更，不为遵守单变量形式而拆开；记录全部差异，不满足验收阈值则恢复旧镜像和旧配置。

官方强调显式设置 MUSA_VISIBLE_DEVICES，并使用项目的 MUSA platform helper 判断可用性，不依赖 torch.cuda.is_available()。本页未实际拉取镜像或运行 GPU benchmark。部署注意事项

历史教程与基础知识

一、什么是 vLLM-MUSA？

vLLM 是 UC Berkeley Sky Computing Lab 开源的 LLM 推理引擎，凭借 PagedAttention 和连续 batching，在高并发场景下实现了远超 HuggingFace Transformers 的吞吐。vLLM-MUSA 则是摩尔线程维护的官方插件，让 vLLM 可以直接运行在 MTT S4000 / S5000 等 GPU 上，提供完全一致的 OpenAI 兼容 API。

当前稳定版本组合：

• vLLM：0.17.0
• PyTorch：2.7.1（配合 torch_musa 2.9.0）
• MUSA SDK：>= 4.3.2
• engine：仅支持 V1 engine（即 VLLM_USE_V1=1），V0 暂不支持

二、依赖与硬件

启动 vLLM-MUSA 之前，你需要按顺序安装这些组件：

硬件方面推荐 MTT S4000（48GB） 或 MTT S5000（80GB），7B 级别模型单卡可跑，32B/72B 需要 TP（张量并行）多卡。

三、Step-by-Step 安装

# 1. 安装 torch_musa + torchada（二者都是必需的）
pip install torch_musa torchada mthreads-ml-py

# 2. 克隆 vllm-musa 仓库
git clone https://github.com/MooreThreads/vllm-musa.git
cd vllm-musa

# 3. 编译安装（可能耗时较久，首次会编译若干 MUSA kernel）
pip install -e .

# 4. 验证插件注册
python -c "import vllm; print('vLLM version:', vllm.__version__)"
python -c "import torchada, vllm; print('vLLM-MUSA plugin loaded OK')"

四、启动 OpenAI 兼容服务

vLLM 最大的卖点之一是内置的 OpenAI 兼容 HTTP 服务。下面这条命令把 DeepSeek-R1-Distill-Qwen-7B 跑成一个在 :8000 端口监听的 API：

# 选择一张卡
export MUSA_VISIBLE_DEVICES=0
# 显式启用 V1 engine（vllm-musa 仅支持 V1）
export VLLM_USE_V1=1

vllm serve deepseek-ai/DeepSeek-R1-Distill-Qwen-7B \
  --dtype float16 \
  --max-model-len 4096 \
  --gpu-memory-utilization 0.9 \
  --port 8000

启动成功后，终端会打印出 INFO Starting OpenAI-compatible server on 0.0.0.0:8000。用 curl 验证一下：

curl http://localhost:8000/v1/chat/completions \
  -H "Content-Type: application/json" \
  -d '{
    "model": "deepseek-ai/DeepSeek-R1-Distill-Qwen-7B",
    "messages": [
      {"role": "user", "content": "用三句话解释什么是 PagedAttention。"}
    ],
    "max_tokens": 256,
    "temperature": 0.7
  }'

至此，你已经拥有了一个与 OpenAI API 完全兼容的本地 LLM 服务 —— 任何支持 OPENAI_BASE_URL 的客户端（LangChain、OpenWebUI、ChatBox 等）都可以直接连上来。

五、多卡张量并行

跑 32B / 72B 级别模型时，单卡显存不够，需要张量并行（TP）。vLLM 通过 --tensor-parallel-size 参数开启：

export MUSA_VISIBLE_DEVICES=0,1,2,3
export VLLM_USE_V1=1

vllm serve Qwen/Qwen2.5-72B-Instruct \
  --dtype float16 \
  --tensor-parallel-size 4 \
  --max-model-len 8192 \
  --gpu-memory-utilization 0.9

六、DeepSeek 模型专项适配：EPLB 与 MTP

4/20 周报里，vllm-musa 合入了两个关键 PR，显著增强了 DeepSeek 系列模型的推理质量与吞吐：

6.1 EPLB（Expert-Level Parallel Balancing）

DeepSeek-V3 / R1 都是 MoE 模型，有几百个专家。训练完成后各专家的调用频率 不均匀，简单的专家并行会让热门专家成为瓶颈。EPLB 动态地把热门专家复制到多张卡上，把冷门专家聚合在少数卡上，整体吞吐可提升 15~30%。

# 在 vllm serve 时额外传入（来自 vllm-musa PR #26）
export VLLM_MUSA_ENABLE_EPLB=1
export VLLM_MUSA_EPLB_REBALANCE_INTERVAL=500   # 每 500 step 重新平衡

vllm serve deepseek-ai/DeepSeek-V3 \
  --tensor-parallel-size 8 \
  --enable-expert-parallel \
  --dtype bfloat16

6.2 MTP（Multi-Token Prediction）

MTP 是 DeepSeek-V3 原生支持的投机解码方式：主模型一次预测多个 token，极大降低 decode 阶段的延迟。vllm-musa 通过 --speculative-config 暴露该能力。

vllm serve deepseek-ai/DeepSeek-V3 \
  --tensor-parallel-size 8 \
  --speculative-config '{"method":"mtp","num_speculative_tokens":3}'

七、常见问题与排障

八、吞吐量压测

压测脚本示例（需要 pip install openai）：

import time, asyncio
from openai import AsyncOpenAI

client = AsyncOpenAI(base_url="http://localhost:8000/v1", api_key="EMPTY")

async def one_req(i):
    resp = await client.chat.completions.create(
        model="deepseek-ai/DeepSeek-R1-Distill-Qwen-7B",
        messages=[{"role": "user", "content": f"写一首关于编号 {i} 的五言绝句。"}],
        max_tokens=128,
    )
    return resp.choices[0].message.content

async def main(n=128, concurrency=16):
    sem = asyncio.Semaphore(concurrency)
    async def bound(i):
        async with sem:
            return await one_req(i)
    t0 = time.time()
    results = await asyncio.gather(*[bound(i) for i in range(n)])
    dt = time.time() - t0
    print(f"{n} 个请求耗时 {dt:.1f}s，QPS = {n/dt:.2f}")

asyncio.run(main())

九、MUSA 5.1.0 兼容性提示（2026-04-27 补充）

本周 vllm-musa 仓库合入了一个针对 MUSA 5.1.0 软件栈的修复（PR #28）。如果你的环境已升级到 MUSA SDK 5.1.0，需要注意以下两点：

• 现象：在 MUSA 5.1.0 上从源码编译 paged attention kernel 时，mtcc 会报 auto-vectorization 拓展优化错误。
• 修复：升级到最新 main 分支（包含 PR #28），编译脚本会自动禁用该项拓展优化。官方说明该调整不会带来性能损失。

# 检查当前 MUSA 软件栈版本
mcc --version           # 查看 mtcc 版本
musa-sdk --version       # 查看 MUSA SDK 版本

# 如果是 5.1.0，从源码重新编译
cd vllm-musa
git pull origin main
pip install -e . --no-build-isolation

十、升级到 vLLM 0.20.0（2026-05-04 新增）

2026-04-28，vllm-musa 仓库合入了 PR #27 "support vllm releases/v0.20.0"，正式推进上游 vLLM v0.17.0 → v0.20.0 的版本跳迁。官方验证过两个关键路径：

• Qwen3-8B 使用 FLASH_ATTN backend 运行通过。
• DeepSeek-Coder-V2-Lite-Instruct 使用 FLASHMLA backend 运行通过（验证了 FlashMLA 在 MUSA 上的端到端可用性）。

升级到 v0.20.0 带来的主要变化：

# 1. 升级到支持 v0.20.0 的 vllm-musa
cd vllm-musa
git pull origin main
pip install -e . --no-build-isolation

# 2. 验证 Qwen3 (FLASH_ATTN)
VLLM_ATTENTION_BACKEND=FLASH_ATTN \
  vllm serve Qwen/Qwen3-8B \
    --tensor-parallel-size 1 \
    --dtype bfloat16

# 3. 验证 DeepSeek-Coder-V2-Lite (FLASHMLA)
VLLM_ATTENTION_BACKEND=FLASHMLA \
  vllm serve deepseek-ai/DeepSeek-Coder-V2-Lite-Instruct \
    --tensor-parallel-size 2 \
    --dtype bfloat16

十一、torchada Fused MoE Triton 实战（2026-05-04 新增）

同一周 torchada 仓库合入了 PR #60 "feat(fused_moe_triton): add configuration-driven MoE kernel tuning"，提供了一套面向 MUSA 的 fused MoE Triton kernel 子系统，它是近期跑 DeepSeek / Mixtral / Qwen-MoE 等专家混合模型的首选加速路径。

核心能力：

• fused MoE 核心 kernel，支持 FP8 / INT8 / INT4 三种量化模式。
• per_token_group_quant_fp8：逐 token-group 的 FP8 量化助手函数。
• moe_align_block_size：expert 路由 + token 对齐工具。
• 支持 TMA（Tensor Memory Accelerator）加速路径。
• “configuration-driven”：其同仓库提供了针对不同 MoE shape 预调到的 Triton kernel 参数 JSON，启动时自动匹配。

# torchada fused_moe_triton 使用示例
import torch
import torchada
from torchada.fused_moe_triton import (
    fused_moe,
    per_token_group_quant_fp8,
    moe_align_block_size,
)

# 场景：DeepSeek-V2 风格的 MoE，top-2 路由
batch_size, seq_len = 4, 4096
hidden_size = 5120
intermediate_size = 1408
num_experts = 64
top_k = 6

hidden_states = torch.randn(
    batch_size * seq_len, hidden_size,
    dtype=torch.bfloat16, device="cuda"
)

# 1. (可选) 逐 token-group 的 FP8 量化输入
hidden_q, hidden_scale = per_token_group_quant_fp8(
    hidden_states, group_size=128
)

# 2. 路由输出 (已经 softmax 之后的 top-k logits + indices)
topk_weights = torch.rand(batch_size * seq_len, top_k, device="cuda")
topk_ids     = torch.randint(0, num_experts, (batch_size * seq_len, top_k), device="cuda", dtype=torch.int32)

# 3. 调用 fused MoE
out = fused_moe(
    hidden_states,
    w1, w2,                       # 两个专家矩阵 (FP8 / INT8 / INT4)
    topk_weights, topk_ids,
    inplace=True,
    use_fp8_w8a8=True,            # 开启 FP8 路径
)

在 vLLM 侧启用该 kernel 只需一个环境变量（仓库包会自动在 fused MoE 路径中优先使用 torchada）：

# 启用 torchada fused MoE Triton kernel 加速
export VLLM_MUSA_USE_TORCHADA_FUSED_MOE=1

# DeepSeek-V2 / V3 类 MoE 推荐走 FP8 路径
vllm serve deepseek-ai/DeepSeek-V2 \
  --tensor-parallel-size 8 \
  --enable-expert-parallel \
  --quantization fp8 \
  --dtype bfloat16

十二、FP8 推理实战（2026-05-11 新增）

2026 年 5 月 7 日 vllm-musa PR #34 "Support fp8 on musa" 合入主分支，这是国内 GPU 生态里把 FP8 推理真正落到 vLLM 主干上的一次关键升级。背景是 DeepSeek-V4 这类前沿模型已经采用 "FP4+FP8" 混合精度策略，而 MTT S5000 作为国内首个原生支持 FP8 的全功能 GPU，能直接吃到这块红利。

12.1 FP8 在 vllm-musa 里具体改了什么

• csrc/musa/quantization/per_token_group_quant.cu：新增 FP8 per-token group 量化 CUDA/MUSA 内核，是后续 GEMM/MoE 的前置算子。
• vllm_musa.model_executor.kernels.linear.scaled_mm.deep_gemm 中的 MUSADeepGemmFp8BlockScaledMMKernel：通过 DeepGEMM 路径执行 FP8 block-scaled 矩阵乘，对应 dense Linear 层的主体计算。
• vllm_musa.model_executor.layers.fused_moe.fused_moe.fused_experts_impl：让 fused MoE 路径也吃到 FP8，配合上一节的 torchada fused MoE Triton，DeepSeek/Mixtral 类 MoE 的全部专家 FFN 都走 FP8。

12.2 起一个 FP8 推理服务

# 拉最新 vllm-musa 主分支
cd vllm-musa && git pull origin v0.20.0-dev && pip install -e . --no-build-isolation

# 启动 FP8 推理服务（以 DeepSeek-V4-Flash-mthreads-FlagOS 为例）
export VLLM_MUSA_USE_TORCHADA_FUSED_MOE=1
vllm serve /models/DeepSeek-V4-Flash-mthreads-FlagOS \
    --served-model-name deepseek-v4 \
    --quantization fp8 \
    --kv-cache-dtype fp8 \
    --tensor-parallel-size 2 \
    --gpu-memory-utilization 0.85 \
    --max-model-len 32768

# 简单测一下吞吐
curl http://localhost:8000/v1/completions \
  -H "Content-Type: application/json" \
  -d '{"model":"deepseek-v4","prompt":"写一段 MUSA kernel 的 hello world","max_tokens":256}'

十三、TurboQuant：KV-Cache 极致压缩（2026-05-11 新增）

2026 年 5 月 7 日 vllm-musa PR #35 "Support turboquant on musa" 合入，给 MUSA 平台带来了一个全新的 attention backend：TurboQuant。它解决的是长上下文推理里最贵的一块——KV-Cache 显存。

13.1 TurboQuant 的能力

• 新增 MUSA TurboQuant attention backend，注册到 vLLM 的 backend 选择逻辑里。
• 支持的 KV-Cache 量化档位（命名约定：k<k-bits>v<v-bits>_nc）：
  • turboquant_4bit_nc：K、V 都 4-bit，最大压缩比。
  • turboquant_k3v4_nc：K 3-bit / V 4-bit，平衡精度与吞吐。
  • turboquant_3bit_nc：K、V 都 3-bit，超长上下文场景。
• 不支持 k8v4（PR 中已显式 reject），原因是配置在 MUSA 上未做验证。
• 平台初始化时 torchada 必须在 torch 之前 import，PR 修正了这个顺序；同时 patch 了 _flash_attn_varlen 的参数兼容 TurboQuant 的 decode 路径。

13.2 启动一个 Qwen3-8B + TurboQuant K3V4 服务

下面这条命令直接来自 vllm-musa PR #35 的官方验证脚本，能在 2 张 MTT S5000 上稳定跑通：

# 关键点：torchada 必须先于 torch 被 import，vllm-musa 在 platform 初始化里已经处理好
vllm serve /home/dist/Qwen3-8b \
    --served-model-name qwen \
    --gpu-memory-utilization 0.7 \
    -tp 2 \
    -cc.mode=NONE \
    -cc.cudagraph_mode=FULL \
    --cudagraph-capture-sizes 4 \
    --kv-cache-dtype turboquant_k3v4_nc \
    --data-parallel-size 1

十四、FA3 + DCP：多卡长序列推理的关键一环（2026-05-18 追加）

2026-05-12 合入主线的 vllm-musa PR #37 support fa3 dcp，补齐了 vllm-musa 中 FlashAttention 后端在 DCP（Decode Context Parallel）模式下缺失的能力。之前 self.dcp_world_size > 1 的路径会直接抛 NotImplementedError，使多机推理场景无法使用 flash_attn。由于 DCP 是 vLLM 在长序列计算阶段拆分 KV 计算的核心机制，补上之后，MTT S5000 上能跑起 32K+ 长上下文的 Qwen3 多卡部署。

PR #37 在 flash_attn.py 中增加了 _forward_with_dcp 函数，实现了两种通信后端：a2a（all-to-all）和 ag_rs（all-gather + reduce-scatter）。官方验证命令：

# FA3 + DCP 多卡长序列推理
vllm serve /home/dist/Qwen3-30B-A3B-Instruct-2507 \
  --served-model-name qwen-dcp \
  --gpu-memory-utilization 0.7 \
  --max-model-len 2048 \
  --max-num-seqs 2 \
  --max-num-batched-tokens 1024 \
  -tp 8 \
  -dcp 2 \
  --dcp-comm-backend a2a \
  --dcp-kv-cache-interleave-size 16 \
  -ac.backend FLASH_ATTN

参数说明：

• -tp 8：张量并行 8 路（权重拆分到 8 张卡）。
• -dcp 2：DCP 拆分为 2 路，在 decode 阶段把 KV 计算再拆一层，适合长序列 + 多卡。
• --dcp-comm-backend a2a：选择 all-to-all 通信后端；另一选项 ag_rs 适合 NVLink 拓扑其他形状的集群。
• --dcp-kv-cache-interleave-size 16：KV cache 的交错粒度，多卡同步开销与计算量的平衡点。

十五、DeepSeek/Qwen FP8 综合优化栈（2026-05-18 追加）

2026-05-15 合入主线的 vllm-musa PR #39 Consolidate DeepSeek/Qwen FP8 performance optimizations 是一份集中交付的大包（3116 行添加 / 36 个文件），把近月 S5000 上针对 DeepSeek-V2-Lite-Chat-FP8 与 Qwen FP8 模型的多项优化一次性合入。你可以把它看作“FP8 路径从跑通到跑顺”的里程碑。

主要增量可以分五大块：

• FP8 运行时加固：FP8 量化工具、DeepGEMM dispatch、scaled-mm fallback 处理、fused MoE chunking、layernorm/activation 包装与自定义 op 绑定。
• torch.compile / CUDAGraph 兼容：为 vLLM 的 compile / graph 路径追加 MUSA 专用 patch 与 graph policy，让“FP8 + 编译加速”同时生效。
• FP8 GEMV 与 MoE 调优：register FP8 group quant fast path，Qwen / DeepSeek FP8 MoE GEMV 专项调优。
• TP padding：DeepSeek/Qwen FP8 服务路径需要 TP padding 与 CUDAGraph policy 处理，避免 block_k=128 在某些 shape 上的划分不兼容。
• 连贯测试：tests/test_musa.py、tests/test_patches.py、tests/test_fused_moe_chunking.py 覆盖 79 passed / 1 skipped，运行仅 9.33s。

官方给出的 单 batch 端到端性能变化（DeepSeek-V2-Lite-Chat-FP8，MUSA-0044 vs MUSA-0035）：

十六、MiniMax-M2.5 + Eagle3 SOTA 推理（2026-05-25 追加）

2026-05-18 合入主线的 vllm-musa PR #40 与紧随其后的 PR #41 (enhance) 共同把 MiniMax-M2.5 这款 230B+ MoE 大模型的推理路径与 Eagle3 投机解码 一次性带到了 MTT S5000 上，4885 行 / 32 文件的合入直接将单 batch 4K 输入 / 1K 输出场景下的解码速率推到 95.2 tok/s，是当下国产 GPU 上同规模 MoE 推理的官方 SOTA。

16.1 启动配方

PR #40 的官方 SOTA 启动命令（8× MTT S5000）：

export VLLM_ATTENTION_BACKEND=TREE_ATTN
export VLLM_WORKER_MULTIPROC_METHOD=spawn
export VLLM_DISABLE_COMPILE_CACHE=1
export PYTHONUNBUFFERED=1
export VLLM_MUSA_SPHERE_SILU_FP8=1   # MUSA-0101 sphere kernel（默认开启）

vllm serve <model-root>/MiniMax/MiniMax-M2.5     --served-model-name MiniMax/MiniMax-M2.5     --trust-remote-code --tensor-parallel-size 8     --enable-auto-tool-choice --tool-call-parser minimax_m2     --reasoning-parser minimax_m2_append_think     --host 0.0.0.0 --port 8000     --no-enable-prefix-caching --no-async-scheduling     --compilation-config '{"cudagraph_mode": "PIECEWISE", "pass_config": {"fuse_attn_quant": true}}'     --speculative-config '{"model": "<model-root>/thoughtworks/MiniMax-M2.5-Eagle3", "num_speculative_tokens": 3, "method": "eagle3", "speculative_token_tree": "[(0,), (1,), (0,0), (1,1), (0,0,0), (1,1,0)]"}'

16.2 配方中的关键开关

• VLLM_ATTENTION_BACKEND=TREE_ATTN：Eagle3 投机解码必须用树形 attention，确保多分支 draft token 的并行 attend。
• VLLM_MUSA_SPHERE_SILU_FP8=1：启用 MUSA-0101 sphere kernel，把 SwiGLU 的 SiLU 阶段直接吃 FP8 输入。
• cudagraph_mode: PIECEWISE + fuse_attn_quant: true：分段 graph 捕获 + 把 attention 后的量化操作融到 graph 里，避免 PyTorch op dispatch 的延迟。
• --no-enable-prefix-caching --no-async-scheduling：SOTA 性能基准为单请求纯解码场景，关闭前缀缓存和异步调度后稳态延迟更稳定。
• speculative_token_tree：定义 Eagle3 draft token 的 6 分支树形 layout，匹配 M2.5 的 acceptance 分布。

16.3 性能压测

SOTA benchmark 命令（random 4K/1K，BS=1 greedy）：

vllm bench serve     --backend openai-chat --base-url http://localhost:8000     --endpoint /v1/chat/completions     --model MiniMax/MiniMax-M2.5     --dataset-name random     --random-input-len 4096 --random-output-len 1024     --num-prompts 1 --max-concurrency 1     --temperature 0     --save-result --result-dir /tmp/mm25_perf --result-filename run.json

十七、EagleFullLoopRunner：把 Eagle3 draft 链做成一张 CUDAGraph（2026-05-25 追加）

仍在开放中的 vllm-musa PR #43（MUSA-0109）把 Eagle3 投机解码的 整条 draft 链 捕获成一张固定形状 CUDAGraph，把先前的「每步 PIECEWISE graph + Python 调度」改成「首步 eager + 链尾一张 graph」。在 MiniMax-M2.5 上：

17.1 关键修复

把 draft 链做成一张 graph 看起来简单，实际有两个隐藏的正确性坑被这次 PR 一次性修掉：

• Draft KV 在链各步没写入：原 _run_n_step_inner 没把 slot_mapping 传给 set_forward_context，vLLM 的 KV-cache 写入会读 forward_context.slot_mapping[layer_name]，None 时直接跳过。结果链上每一步都读到旧的 draft KV，token acceptance 在 position 0 之后整体崩塌。修复方式：构造每层 slot-mapping dict 并显式传入。
• graph 里固化了过期的 kv_cache 指针：CUDAGraph 在捕获时会把 device 指针冻结。runner 原本在 boot 阶段捕获，那时 draft attention 的 kv_cache 还是 8 block 的 dummy；回放时却去访问已经释放的内存（真实 cache ≈ 55k blocks）。修复方式：把 boot 捕获 延迟到 draft attention kv_cache 的 num_blocks 切换成真实大尺寸之后，并在日志里输出 captured-at-boot = 8（每个 TP rank 一次）作为可观测信号。

17.2 默认开关与回退

PR 默认把 VLLM_MUSA_EAGLE_RUNNER 翻成 on，开发者只需要：

# 默认开启（无需任何环境变量）
# 临时回退到旧的迭代路径——零风险
export VLLM_MUSA_EAGLE_RUNNER=0
vllm serve MiniMax/MiniMax-M2.5 ...

# 另有 MUSA-0124 选项（默认 off），把 Eagle3 draft 模型建在 TP=1
# 注意：runner 已经吃掉了它要解决的 launch overhead，叠加无收益
# export VLLM_MUSA_DRAFT_TP1=1   # 不推荐和 runner 同时开

十八、TurboQuant FP8 dtype 修复（PR #42，2026-05-25 追加）

本周还顺手修了 PR #42：MUSA 平台不支持 float8_e4b15 dtype，但 vLLM 0.20 在 TurboQuant 路径下会按 dtype 选择 kernel。新增的 _musa_use_fp8_e4b15 函数直接在 MUSA 上返回 0，把 dtype 选择路由到 float8_e4m3fn / float8_e5m2。验证场景：Qwen3-8B + --kv-cache-dtype turboquant_k8v4 正常启动。

十九、Qwen3 prefill/decode 路径分离 + tilelang RoPE（PR #44，2026-06-01 追加）

PR #44 把 Qwen3 在 vllm-musa 的 FlashAttention 后端拆成独立的 prefill / decode 两条路径，并在两条路径上分别支持 Eagle3 投机解码。同时把 RoPE 切换到 tilelang RoPE，目标是降低 TPOT。这次 PR 还做了一件「反直觉但正确」的事：删除 MUSA_SAFE_MAX_CUDAGRAPH_CAPTURE_SIZE 这个 env——它原本是用来「限制 CUDAGraph 捕获 size 来规避 OOM」，但实际反而触发了 compile-mode 下的内存异常；正确做法是从 capture/compile 路径自身解决，而不是限制 size。

二十、prefill RoPE：用 MUSA cpp + TVM 替代 tilelang JIT（PR #47 / PR #49，2026-06-01 追加）

tilelang 的 RoPE 在 prefill 阶段每次都触发 JIT，把首 token 延迟（TTFT）拖长。PR #47 直接把 RoPE 改写成 MUSA cpp + TVM，避免 prefill 期 JIT；PR #49 在此基础上把 RoPE JIT 包成 vLLM 自定义算子（custom op），让 torch.compile + CUDAGraph 能替换捕获到的 forward 路径，解锁所有用 MusaRotaryEmbedding 的模型在 compile 模式下启动。

更重要的是，PR #49 还删除了 PR #34 时代的 fused MoE 兜底覆写：

• PR #34 时代：fused_experts_impl 被一行覆写劫持到 v0.18 风格的 chunked GEMV workaround——为了让裸 v0.20.0-dev 上的 FP8 路径不至于报 KeyError: <PlatformEnum.OOT: 6>。
• PR #49 之后：vLLM modular_kernel 正确走 Fp8MoeBackend.DEEPGEMM → DeepGemmExperts.apply，进入 mate 的 m_grouped_fp8_gemm_nt_contiguous 高性能路径。

二十一、flashinfer sampler 接入（PR #48，2026-06-01 追加）

PR #48 一次性接入 5 个 flashinfer sampler 接口，让 vllm-musa 的采样路径不再依赖 PyTorch 慢速 fallback：

• musa_top_k_top_p_sampling_from_probs
• min_p_sampling_from_probs
• top_p_sampling_from_probs
• top_p_renorm_probs
• top_k_renorm_probs

同时把 vllm_musa/__init__.py 里 import vllm_musa._inductor 的位置往后挪，避免 musa_platform_plugin 解析失败。

二十二、dflash + Gemma-4-31B：1.47× → 4.08× 投机解码（PR #52，2026-06-01 追加）

2026-05-30 合入的 PR #52 是本周分量最重的 PR：在 vllm-musa 上启用 dflash（block-diffusion 投机解码），目标模型是 google/gemma-4-31B-it，并把 draft-loop FULL CUDAGraph 捕获默认打开。dflash 本身已经在 vLLM base 主线（vllm/v1/spec_decode/dflash.py）；PR #52 是 MUSA 兼容性 + 捕获改造。

22.1 性能数字

22.2 启动命令

# 关键：用 --attention-backend，而不是 VLLM_ATTENTION_BACKEND 环境变量
#       因为 PR #52 所在的 vLLM 分支已经废弃了该环境变量
vllm serve google/gemma-4-31B-it   --tensor-parallel-size 8   --attention-backend FLASH_ATTN   --speculative-config '{"method": "dflash", "draft_tensor_parallel_size": 1, "num_speculative_tokens": 4}'   --max-num-batched-tokens 8192   --max-model-len 32768

22.3 关键开关

• VLLM_MUSA_DFLASH_FULL_WRAP=1：默认开启，把 dflash 的 draft 链整段做 FULL CUDAGraph 捕获。
• VLLM_MUSA_DFLASH_FULL_WRAP=0：退回到 PIECEWISE 捕获（用于 A/B 对比或 debug）。
• --attention-backend FLASH_ATTN：mate 的 FLASH_ATTN backend 现在支持 non-causal + head_dim=512，dflash 的 block 块特别需要这两项。

二十三、ccache 加速 MUSA 原生编译（PR #51，2026-06-01 追加）

vllm-musa 的 csrc 构建从一次 2385 秒（约 40 分钟）下降到一次 119 秒（约 2 分钟）——只要你启用 ccache。PR #51 为 MUSA native 构建管线接入 ccache，二次以上 editable install 命中率 40/40，提速 20×。

# 安装 ccache（Ubuntu / Debian）
sudo apt-get install -y ccache

# vllm-musa editable install
cd vllm-musa
pip install -e . -v
# 第二次起进入分钟级编译

# 验证 ccache 命中
ccache -s | grep -E "Hits|Misses"

附录：2026-08-10 更新章节 —— 新镜像、/vllm-workspace、依赖命名与分支选择

0. 为什么这类"只是改了个名字"的变更最危险

在开始逐项看变化之前，先建立一个心理预期：推理框架的升级风险并不主要来自"新功能有 bug"，而是来自约定变了而你不知道。新功能出问题会报错、会有栈，你能查；而镜像内工作目录改了、wheel 名字变了、依赖下限提高了，这些变化的典型表现是 CI 在某个深夜构建失败，或者更糟——构建成功但运行时行为已经不同。

本轮四项变化里，有三项属于这种"契约型"变更。它们共同的特征是：单看每一项都很小，但只要你的部署脚本里有一行写死了旧约定，整条流水线就会断。所以本章节的阅读方式建议是：一边读，一边把自己仓库里的 Dockerfile、requirements、部署脚本与 K8s 清单打开做 grep，而不是读完再回头找。本节每一小节都附了对应的排查命令，就是为了配合这种读法。

1. 新镜像 tag

随 PR #161（Qwen Omni ModelOpt FP8 推理支持）一同发布的新推理镜像为：

registry.mthreads.com/mcconline/inference/vllm:v0.24.0-ph1-5.2.0-torch2.9.1.post1-20260801

tag 的构成本身就是一份版本声明，值得逐段读懂：v0.24.0 对应 vLLM-MUSA 分支；ph1 是平湖架构（MTT S5000）；5.2.0 是 MUSA SDK 版本；torch2.9.1.post1 是 PyTorch/torch_musa 版本；20260801 是构建日期。养成从 tag 反推软件栈的习惯，能避免大量"我以为我在跑 5.2.0"的误判。

PR #161 本身带来的能力是 Qwen Omni 的 ModelOpt FP8 推理支持，实现要点包括：MUSA 上的静态 scaled_fp8_quant 回退路径、E4M3 clamp，以及 seeded multinomial 限定在 vocab 151936 / 152064 / 248320 三种规模。最后一项是典型的"工程约束"——只对特定 vocab size 做了保证，用别的模型时需要自行验证。

2. Docker 工作目录改为 /vllm-workspace

PR #154 把镜像结构与上游 vLLM 对齐，两处变化：工作目录改为 /vllm-workspace，构建目标为 final 与 vllm-openai。

# 旧脚本里写死的路径需要全量排查替换
grep -rn "workspace\|/vllm" deploy/ scripts/ 2>/dev/null

# 进入容器后确认
docker run --rm -it \
  registry.mthreads.com/mcconline/inference/vllm:v0.24.0-ph1-5.2.0-torch2.9.1.post1-20260801 \
  bash -lc 'pwd && ls /vllm-workspace'

# 自行构建时指定目标
docker build --target vllm-openai -t my-vllm-musa:latest .
docker build --target final       -t my-vllm-musa:final .

与上游对齐的直接好处是：上游 vLLM 社区的示例脚本、compose 文件、K8s 清单可以近乎原样复用，不必再做路径改写。这也是为什么这类"看起来只是改了个目录"的 PR 值得单独讲——它降低的是长期维护成本。

3. 依赖命名契约：mthreads-ml-py 与 wheel 去 s5000 后缀

PR #149 调整了两处依赖约定：

• mthreads-ml-py>=2.4.1：GPU 管理库最低版本上调。这与 mthreads-ml-py PR #13 的 MTML 2.4.1 支持是同一条线，详见本站 m2-l5 的 2026-08-10 更新章节。
• MUSA wheel 名称去掉 s5000 后缀：这是一处容易被自动化脚本绊倒的变化。如果你的内网源、离线包清单或 requirements.txt 里按旧命名做了精确匹配，升级后会直接找不到包。

# 检查旧命名残留
grep -rniE "s5000" requirements*.txt constraints*.txt Dockerfile* 2>/dev/null

# 确认 pymtml 版本达标
pip install "mthreads-ml-py>=2.4.1"
python -c "import pymtml; print(getattr(pymtml,'__version__','unknown'))"

4. 分支选择：v0.24.0-dev 与新出现的 v0.26.0-dev

根据 vllm-musa 分支列表，2026-08-07 起仓库出现了新分支 v0.26.0-dev。两条分支的状态如下：

5. Step-by-Step 迁移与回滚

# Step 1 — 记录当前状态（回滚依据）
docker images | grep vllm
git -C /path/to/vllm-musa rev-parse --abbrev-ref HEAD
pip freeze > /tmp/before_20260810.txt

# Step 2 — 拉取新镜像（不删旧镜像）
docker pull registry.mthreads.com/mcconline/inference/vllm:v0.24.0-ph1-5.2.0-torch2.9.1.post1-20260801

# Step 3 — 全量排查脚本中的路径与命名
grep -rn "s5000" . 2>/dev/null | head
grep -rn "WORKDIR\|workdir" Dockerfile* 2>/dev/null

# Step 4 — 在新镜像中冒烟
docker run --rm -it --network host \
  -v /models:/models \
  registry.mthreads.com/mcconline/inference/vllm:v0.24.0-ph1-5.2.0-torch2.9.1.post1-20260801 \
  bash -lc 'cd /vllm-workspace && python -c "import vllm,torch,torch_musa;print(vllm.__version__, torch.__version__)"'

# Step 5 — 拉起服务并验证
#   VLLM_USE_V1=1 vllm serve /models/Qwen3-8B --port 8000
curl -s http://127.0.0.1:8000/v1/models | python -m json.tool

# Step 6 — 回滚：切回旧镜像 tag 并重启，pip 环境用 /tmp/before_20260810.txt 还原
# pip install -r /tmp/before_20260810.txt

6. 升级前的一张自查清单

把上面四项变化收敛成一张可以直接拿去评审的清单。建议把它贴进你团队的变更单模板，每次升级 vLLM-MUSA 都过一遍。

这里特别提醒一个实践细节：不要在同一次变更里同时升镜像和改部署脚本。正确的顺序是先把脚本里的硬编码路径与命名匹配改成兼容两个版本的写法（例如用变量而非字面量、用模糊匹配而非精确匹配），合并并验证旧镜像仍能跑通，再单独做一次只换镜像 tag 的变更。这样一旦出问题，你能立刻判断是脚本改错了还是镜像引入的，而不是对着一堆同时变动的变量猜。同理，分支切换与依赖升级也应该拆成两次独立的变更。

最后关于分支选择再补一句：v0.26.0-dev 的头提交时间反而早于 v0.24.0-dev，这说明它目前更像是一个从某个时间点切出去的开发分支，而非已经超越 v0.24.0-dev 的新基线。判断一个新分支是否已经可用，看分支名里的版本号是最不可靠的方式，应该看的是头提交内容、提交频率，以及官方是否已为它发布对应镜像。截至本章节写作时，v0.26.0-dev 三项都不满足，因此结论是观望。

本章节小结：这轮变化没有引入新功能概念，但每一项都会让旧命令失效。镜像升级到 ...-20260801；工作目录改为 /vllm-workspace，构建目标为 final / vllm-openai；依赖要求 mthreads-ml-py>=2.4.1，MUSA wheel 名称去掉 s5000 后缀；分支上 v0.24.0-dev 仍是生产基线，新出现的 v0.26.0-dev 暂作观察。迁移的核心纪律是：先记录、后升级、保留旧镜像、每步验证。

v0.28.0 升级边界：先读交付物

GitHub Release 于 2026-11-03（UTC）发布，标签 v0.28.0 指向 commit 070c0ea09ec266972fd6097937fde97d400d70ed，对齐上游 vLLM 0.28.0（pin 2cf0a69），vllm-musa 包版本为 0.1.28。发布说明称 v0.28.0-dev 相对 v0.24.0 分支（2733b6f）包含 70 个提交、9 位贡献者。vLLM-MUSA v0.28.0 Release

官方文档的 S5000 发布信息页已列出 0.28.0，但页面未写发布日期；本站以 GitHub Release 时间为准，不把抓取日期当作发布日期。vLLM-MUSA S5000 版本发布信息

软件栈对照：整体替换，不逐个混装

数据来自 Release 软件栈表。注意 torch_musa 仓库本轮读取到的最新正式 Release 仍是 2.11.0.post1，post2 随镜像交付；不要从 torch_musa Release 页单独拼装这一组合。vLLM-MUSA v0.28.0 Release

镜像标签：两处写法需当场核对

GitHub Release 写的最终镜像是 registry.mthreads.com/mcconline/inference/vllm/vllm-openai:v0.28.0；官方文档按 CPU 平台给出 v0.28.0-amd64-ph1-musa5.2.0（Intel/AMD）和 v0.28.0-hygon-ph1-musa5.2.0（Hygon）。vLLM-MUSA v0.28.0 Release vLLM-MUSA S5000 版本发布信息

本站未拉取验证这些标签。部署前先用 docker manifest inspect 或实际 pull 确认标签存在、架构正确，并记录镜像 digest；后续复现以 digest 为准。Release 中的 qwen38-flash-next、glm53-flash 等模型专用 tag 在说明中注明“未在本分支内验证”，不要作为通用基线。vLLM-MUSA v0.28.0 Release

已知限制与默认行为

• MUSA 默认仍使用 Model Runner V1；V2 需显式启用，DSpark 场景另有 V2 override。vLLM-MUSA v0.28.0 Release
• MLA prefill 固定使用 MATE FlashAttention，output_scale 尚未实现。vLLM-MUSA S5000 版本发布信息
• diffusion 模型固定使用 TRITON_ATTN，即使指定 FLASH_ATTN 也会被覆盖并告警。vLLM-MUSA v0.28.0 Release
• 显式设置 --block-size 会关闭混合 attention 的 64-token 快路径并告警，而不再被静默覆盖。vLLM-MUSA v0.28.0 Release
• FP8、MTP 和 CUDA Graph 路径对模型 revision、依赖版本和输入 shape 敏感，文档要求上线前用真实请求做语义和性能验证。vLLM-MUSA S5000 版本发布信息

Cookbook 表格面向 S5000、约 4K 输入 / 1K 输出的有界并发；Release 要求同一行的精度、TP 与推测解码设置整体复制，更换模型、并发或上下文后需重新校准。发布说明中的性能数字来自各提交的测量记录，并说明不是 v0.24.0 到 v0.28.0 的整体基准。vLLM-MUSA v0.28.0 Release

升级验收步骤

1. 保留当前 v0.24.0 容器与镜像 digest，作为回滚目标；新版本在独立容器启动。
2. 在容器内记录 pip list 中 vllm、vllm-musa、torch、torch_musa、mate、torchada 的版本，与上表对照；任何一项不符就停止，先查清镜像来源。
3. 选 Cookbook 中与现场卡数一致的一行，整行复制启动参数；先跑单请求正确性，再跑固定并发压测，用本模块的 TTFT/TPOT 指标记录。
4. 与 v0.24.0 在同一模型、同一输入集上对比输出语义和指标；差异超出预先定义的阈值即回滚。
5. 验收单中分列“Release 说明”“作者测量”“现场复测”，不要把提交说明中的倍数写成现场结论。