vLLM-MetaX MoE Triton 配置调优 —— 参赛文件包

参赛者：LindseyMei 目标仓库：MetaX-MACA/vLLM-metax（GitLink 主仓：metax-maca/vLLM-metax）贡献内容：为 MetaX C500（MXC500）补充 3 组缺失的 fused-MoE Triton tuned 配置 PR：MetaX-MACA/vLLM-metax#313（已关闭，未合并）重新提交 PR：MetaX-MACA/vLLM-metax#323（目标分支 v0.13.0-dev，源分支 feat/moe-c500-configs）后续主仓 PR：MetaX-MACA/vLLM-metax#319（将 3 个 tuned config 补齐到 releases/v0.13.0）包更新日期：2026-07-13

一句话概述

本文件包记录为 vLLM-MetaX（沐曦 MACA 后端） 做的 MoE（Mixture-of-Experts）Triton 内核配置调优 贡献：在真机 MetaX C500 上定位到 6 个热门 MoE 模型在 MXC500 上无 tuned 配置、回退默认 tile 的性能缺口，通过直接调优 MACA 实际运行的 vllm_metax fused-MoE 内核，给出针对性的两段式 Triton 配置，实现 kernel 级 1.17–3.38x 提速，并通过正确性与加载验证。

提交材料

贡献	GitHub PR	GitLink 留痕	状态
MoE Triton 配置调优（3 个 MXC500 缺口 shape）	MetaX-MACA/vLLM-metax#313	#221	CLOSED（未合并；已按管理员建议重新提交到 `v0.13.0-dev`）
将 3 个 tuned config 补齐到 releases/v0.13.0 主仓库	MetaX-MACA/vLLM-metax#319	#223	CLOSED（未合并）
MoE Triton 配置调优重新提交到 `v0.13.0-dev`	MetaX-MACA/vLLM-metax#323	#221	OPEN（目标 `v0.13.0-dev`，等待 review / merge）
扩展 MXC500 fused-MoE tuned 配置（DeepSeek-V2 / Mixtral-8x22B / DeepSeek-V3）	MetaX-MACA/vLLM-metax#329	#226	OPEN（决赛增量，目标 `v0.13.0-dev`）

1. 问题（真实缺口）

vllm_metax 的 fused-MoE 走 Triton 内核，其 tile 配置按 (H, E, N, device_name[, dtype]) 从 vllm_metax/model_executor/layers/fused_moe/configs/ 加载 tuned JSON；未命中则回退到通用 get_default_config，并打印 Using default MoE config. Performance might be sub-optimal!。

实测缺口：以下常见 MoE shape 在 MXC500 上没有对应的 tuned 配置文件：

Shape (H, E, N)	top_k	代表模型	架构	状态
(2048, 60, 1408)	4	`Qwen/Qwen1.5-MoE-A2.7B`	`Qwen2MoeForCausalLM`	PR #323
(2048, 64, 1408)	6	`deepseek-ai/DeepSeek-V2-Lite`	`DeepseekV2ForCausalLM`	PR #323
(2048, 128, 768)	8	`Qwen/Qwen3-30B-A3B`	`Qwen3MoeForCausalLM`	PR #323
(5120, 64, 1536)	6	`deepseek-ai/DeepSeek-V2`	`DeepseekV2ForCausalLM`	PR #329（本次新增）
(6144, 8, 2048)	2	`mistralai/Mixtral-8x22B` TP8	`MixtralForCausalLM`	PR #329（本次新增）
(7168, 256, 256)	8	`deepseek-ai/DeepSeek-V3` / `DeepSeek-R1` TP8	`DeepseekV3ForCausalLM`	PR #329（本次新增）

2. 方法（调优的是真正的 MACA 内核）

关键点：上游 benchmarks/kernels/benchmark_moe.py 导入的是上游 fused_experts，而 MACA 运行时用的是 vllm_metax 自己的 fused_experts（经 OOT 注册，二者是不同对象，已核实）。因此上游 tuner 调的是错误的内核。

本方案自建微基准（moe_tuning/tools/moe_tune.py）直接调优 MACA 内核：

用 vllm.model_executor.layers.fused_moe.override_config 强制指定 tile，计时单次 vllm_metax ... fused_experts()（随机权重，因此无需下载模型）；
CUDA event + torch.cuda.synchronize()，warmup + 多次取中位数；
网格搜索 BLOCK_SIZE_M/N/K、GROUP_SIZE_M、num_warps、num_stages；
将胜出的扁平配置包装成 MetaX 两段式 schema（stage1/stage2 + ACCF32/SPLIT_K/pipeline/scenario），对称填充（低风险起点）。

正确性验证：moe_tuning/tools/moe_verify.py 使用 torch.allclose(rtol=2e-2, atol=2e-2) 比较「默认 tile」与「调优 tile」的 MoE 输出，并调用 get_moe_configs() 确认配置文件会被运行时正确加载。

3. 结果（kernel 级，真机 sGPU 切片）

Qwen1.5-MoE-A2.7B（H=2048, E=60, N=1408, top_k=4）

M	默认 (ms)	调优后 (ms)	加速
64	1.154	0.915	1.26x
256	1.291	1.025	1.26x
1024	2.340	1.567	1.49x
2048	4.559	2.589	1.76x

DeepSeek-V2-Lite（H=2048, E=64, N=1408, top_k=6）

M	默认 (ms)	调优后 (ms)	加速
128	1.260	1.046	1.21x
512	1.642	1.311	1.25x
1024	3.163	1.720	1.84x
2048	6.793	3.081	2.20x

Qwen3-30B-A3B（H=2048, E=128, N=768, top_k=8）

M	默认 (ms)	调优后 (ms)	加速
256	1.505	1.189	1.27x
1024	2.650	1.850	1.43x
2048	4.424	2.755	1.61x

DeepSeek-V2（H=5120, E=64, N=1536, top_k=6）

M	默认 (ms)	调优后 (ms)	加速
128	3.183	2.381	1.34x
256	3.582	2.647	1.35x
512	4.678	2.972	1.57x
1024	9.488	3.831	2.48x
2048	24.488	7.248	3.38x

Mixtral-8x22B TP8（H=6144, E=8, N=2048, top_k=2）

M	默认 (ms)	调优后 (ms)	加速
16	0.806	0.647	1.25x
64	0.904	0.709	1.27x
256	1.770	1.017	1.74x
512	2.855	1.375	2.08x
1024	5.299	2.174	2.44x
2048	9.727	3.573	2.72x

DeepSeek-V3 / DeepSeek-R1 TP8（H=7168, E=256, N=256, top_k=8）

M	默认 (ms)	调优后 (ms)	加速
512	3.470	2.919	1.19x
1024	4.648	3.421	1.36x
2048	11.342	4.624	2.45x

小 M（decode）保留默认 tile 或近似配置，避免延迟回退；大 M（prefill / 大 batch）获得显著加速。

speedup summary

图：6 个已覆盖 shape 在 MXC500 上的 fused-MoE kernel 加速比（默认 / 调优）。

原始数据：moe_tuning/results/moe_tune_result.json。

4. 正确性与加载验证

正确性：同输入下「默认 tile」与「调优 tile」的 MoE 输出 max|Δ|≈1e-3，torch.allclose(rtol/atol=2e-2)=True —— 调优只改 tile 选择，不改数学。
加载：对 3 个 shape 分别调用 get_moe_configs(E, N, None, 0, 0, H) 均能命中本配置（device_name=MXC500），返回全部 10 个 M-条目 → 运行时会被正确采用。

验证命令：

source /data/workspace/vLLM-metax-Project/env_vllm.sh
cp moe_tuning/tools/moe_verify.py /tmp/ && cd /tmp && python3 /tmp/moe_verify.py

5. 重要说明（诚实边界）

本次调优/验证在 C500 sGPU 切片（16GB/25% 算力） 上完成。tile 形状（BLOCK_SIZE/warps/stages）是每-SM 特性，预期可迁移到整卡 C500；grid 级参数（GROUP_SIZE_M=1, SPLIT_K=1）取值保守。整卡 C500 上的最优性未验证，欢迎维护者在整卡上复跑 moe_tuning/tools/moe_tune.py 复核。
因 16GB 限制，完整的 28GB Qwen1.5-MoE bf16 等模型无法做整模型 E2E；本贡献采用 kernel 级基准（与上游 benchmark_moe.py 同类证据，是 MoE 配置 PR 的标准依据）。
PR 仅新增 3 个 JSON 配置文件，无代码改动，风险极低。

6. 目录结构

vLLM-metax-Contributions/
├── README.md                 # 本文件
├── LINKS.md                  # PR / issue 链接汇总
├── docs/
│   └── PROGRESS_vLLM_metax.md   # 路线 A 全栈安装 + 恢复文档
└── moe_tuning/               # 主贡献：MoE Triton 配置调优
    ├── README.md
    ├── configs/
    │   ├── H=2048,E=60,N=1408,device_name=MXC500.json   # Qwen1.5-MoE
    │   ├── H=2048,E=64,N=1408,device_name=MXC500.json   # DeepSeek-V2-Lite
    │   ├── H=2048,E=128,N=768,device_name=MXC500.json   # Qwen3-30B-A3B
    │   ├── H=5120,E=64,N=1536,device_name=MXC500.json   # DeepSeek-V2
    │   ├── H=6144,E=8,N=2048,device_name=MXC500.json    # Mixtral-8x22B TP8
    │   └── H=7168,E=256,N=256,device_name=MXC500.json   # DeepSeek-V3 / R1 TP8
    ├── tools/
    │   ├── moe_tune.py       # MACA 内核微基准调优器
    │   └── moe_verify.py     # 正确性 + 配置加载验证
    └── results/
        └── moe_tune_result.json  # 原始调优数据

7. 复现

source /data/workspace/vLLM-metax-Project/env_vllm.sh

# 调优（无需下载模型，随机权重）；注意从不含 vllm/ 子目录的 cwd 运行
cp moe_tuning/tools/moe_tune.py /tmp/ && cd /tmp
python3 /tmp/moe_tune.py --H 2048 --E 60 --N 1408 --top-k 4 --Ms 1 8 16 32 64 128 256 512 1024 2048

# 正确性 + 加载验证
cp moe_tuning/tools/moe_verify.py /tmp/ && python3 /tmp/moe_verify.py

# 部署到插件：把 configs/*.json 复制到
#   vllm_metax/model_executor/layers/fused_moe/configs/
# 或者用 VLLM_TUNED_CONFIG_FOLDER 覆盖测试

8. 后续可扩展方向

在整卡 C500 上复跑 moe_tuning/tools/moe_tune.py 复核 tile。
补全更多 MoE 配置缺口（如 Qwen2-57B-A14B、DeepSeek-V2 / Qwen3-235B-A22B 的量化 shape）。
在验证过程中发现并修复 metax 后端的真实 bug。
**决赛前非 MoE 配置方向的详细执行计划见 FINAL_STAGE_PLAN.md**（含 attention tile 调优、mctlass 量化路径验证、DeepSeek FlashMLA 调参等）。

9. 决赛增量（2026-08-03 起）

9.1 方向 A：Triton decode attention tile 自动调优

代码改动：
- vllm_metax/attention/ops/triton_decode_attention.py：为 _fwd_kernel_stage1 / _fwd_grouped_kernel_stage1 增加 @triton.autotune
- 搜索空间：非 grouped kernel BLOCK_N ∈ {8,16}、grouped kernel BLOCK_N ∈ {16}；num_warps ∈ {1,2}、num_stages ∈ {1,2}
- 新增环境开关 VLLM_MACA_DECODE_ATTN_AUTOTUNE（默认 1）
扫描脚本：attention_tuning/tools/attn_decode_sweep.py
文档：attention_tuning/README.md
状态：代码已完成；2026-08-09 已补齐真机数据与 baseline 对比：

Model	baseline_ms	autotune_ms	speedup	best_config
deepseek-v2-lite	4.150	1.142	3.63x	BLOCK_N=8, warps=1, stages=1
qwen3-30b-a3b	0.545	0.234	2.33x	BLOCK_N=16, warps=1, stages=1
deepseek-v2	0.280	0.310	0.90x	BLOCK_N=16, warps=4, stages=1
deepseek-v3	0.281	0.310	0.91x	BLOCK_N=16, warps=4, stages=1

PR 已创建：#331 feat(attn): add autotuned triton decode attention tiles for MXC500。

9.2 方向 B：MACA 环境变量解析 bug 修复 + mctlass 路径基准

代码改动：
- vllm_metax/envs.py：修复 USE_PRECOMPILED_KERNEL / MACA_DP_OPT 的 bool(os.environ.get(...)) 解析 bug，统一为 bool(int(...))
- vllm_metax/v1/attention/backends/triton_attn.py：MIN_LAUNCH_GRID_SIZE_2D / NUM_PAR_SOFTMAX_SEGMENTS 改为环境变量可覆盖
验证：MACA_DP_OPT=0 返回 False，USE_PRECOMPILED_KERNEL=0 返回 False，默认值符合预期。
mctlass 基准脚本：moe_tuning/tools/mctlass_benchmark.py
状态：bug 修复已验证；2026-08-07 已补齐 6 个 shape 的 default vs mctlass A/B 数据（mctlass best_ms / default best_ms）：

Shape	M=1	M=16	M=64	M=256	M=1024
Qwen1.5-MoE-A2.7B	1.02x	1.01x	1.00x	1.00x	1.06x
DeepSeek-V2-Lite	1.13x	1.26x	1.29x	1.20x	1.52x
Qwen3-30B-A3B	1.19x	1.00x	0.98x	0.83x	0.78x
DeepSeek-V2	0.77x	0.89x	0.71x	1.00x	1.23x
Mixtral-8x22B TP8	0.69x	0.77x	0.73x	0.63x	0.57x
DeepSeek-V3/R1 TP8	0.84x	0.76x	0.85x	1.00x	1.03x

PR 已创建：#332 fix(envs): correct bool parsing for MACA env vars and make attn thresholds overridable。

9.3 当前状态

GPU 环境已恢复，torch.zeros(10, device='cuda') 可正常执行。
方向 A/B 的真机性能数据与 baseline 对比已于 2026-08-09 补齐：
- attention_tuning/results/attn_decode_sweep.json
- attention_tuning/results/attn_decode_baseline.json
- moe_tuning/results/mctlass_ab_summary.json
两个决赛增量 PR 均已创建：#331、#332。
详细调试与数据补齐过程见 WORKLOG_2026-08-07.md。

10. 后续提交记录

2026-08-02：扩展 MXC500 fused-MoE tuned 配置到 DeepSeek-V2 / Mixtral-8x22B / DeepSeek-V3

GitHub PR：https://github.com/MetaX-MACA/vLLM-metax/pull/329
目标分支：MetaX-MACA:v0.13.0-dev
源分支：LindseyMei:feat/moe-more-c500-configs
提交内容：
- 新增 3 个 MXC500 tuned config：DeepSeek-V2（5120,64,1536）、Mixtral-8x22B TP8（6144,8,2048）、DeepSeek-V3/R1 TP8（7168,256,256）
- 更新 config_list.txt
- 所有配置通过正确性与加载验证
状态：OPEN（决赛增量贡献）

2026-07-13：按管理员建议将 MoE 调优重新提交到 `v0.13.0-dev`

GitHub PR：https://github.com/MetaX-MACA/vLLM-metax/pull/323
GitLink 留痕：https://www.gitlink.org.cn/ccf-ai-infra/GPUApps/issues/221
目标分支：MetaX-MACA:v0.13.0-dev
源分支：LindseyMei:feat/moe-c500-configs
提交内容：
- 原 PR #313 已关闭（未合并）
- 基于 v0.13.0-dev 重新整理并提交相同的 3 个 MXC500 tuned config
- 管理员留言：you could merge this pr to the dev branch of v0.13.0 (v0.13.0-dev)
状态：OPEN（等待上游 review / merge）

2026-07-06：将 tuned config 补齐到 vLLM-MetaX 主仓库

GitHub PR：https://github.com/MetaX-MACA/vLLM-metax/pull/319
GitLink 留痕：https://www.gitlink.org.cn/ccf-ai-infra/GPUApps/issues/223
目标分支：MetaX-MACA:releases/v0.13.0
源分支：LindseyMei:feat/moe-mxc500-missing-configs-v0.13.0
提交内容：
- 将本文件包中的 3 个 tuned config 复制到主仓库 vllm_metax/model_executor/layers/fused_moe/configs/
- 更新 config_list.txt
- 修复 pyproject.toml license 字段以兼容新版 setuptools
状态：CLOSED（未合并）

注：PR #313 为原始调优贡献（GitLink 留痕 #221），PR #319 为将配置正式合并到主仓库的后续补齐（GitLink 留痕 #223）。