[Bug]: Triton block quantized (e.g. MXFP4) MoE kernels producing NaNs due to OOB reads on scale values

快速结论：该报错发生在使用 vLLM 且启用 Hopper GPU（如 H100、B200）的 MXFP4 MoE 后端时。由于 Triton 内核在加载 scale 值时未对尾部进行掩码（mask），导致越界读取 0xff 值被解释为 NaN，进而产生 NaN 输出。优先排查 vLLM 中 vendored Triton 内核的版本（当前为 triton v3.5.1），并应用补丁修复。

问题场景

用户在 vLLM（v0.24.0 及 main 分支）上运行 MXFP4（例如 GPT-OSS 20B）MoE 模型，并在 Hopper 架构 GPU（H100、B200）上使用 Triton 默认内核时触发。问题在启用约束解码（constrained decoding）或其它内存分配操作（如位掩码分配）时更易复现，导致模型输出随机乱码或 0% 准确率。

报错原文

# 无显式报错信息，以 NaN 输出和 0% 准确率表现
# 内核复现失败：
$ python repro_matmul_ogs.py
# 设置 FILL_VALUE=0xff（默认）导致 NaN

# 复现无失败：
$ FILL_VALUE=0xfe python repro_matmul_ogs.py
$ FILL_VALUE=0x0 python repro_matmul_ogs.py

原因分析

可能原因是 Triton 内核（位于 vllm/third_party/triton_kernels/matmul_ogs_details/_matmul_ogs.py）中的 tl.load(WMxScalePtrs) 未使用掩码加载 scale 值。当 EVEN_K 为 False 时，越界读取的内存（例如其他操作分配的 0xff）会被解释为 NaN，导致 MoE 内核产生 NaN 输出。该问题在 Triton v3.5.1 中已有一个潜在补丁，但尚未集成到 vLLM 的 vendored 版本中。

环境排查

• vLLM 版本：v0.24.0 及 main 分支，commit abc71548ef029132c3316b902207f254a246d5
• PyTorch：2.11.0+cu130
• Triton：3.6.0（但 vLLM 使用的 vendored Triton 内核版本为 triton v3.5.1）
• GPU：H100、B200（Hopper 架构）
• 复现依赖：参考 repro_matmul_ogs.py 内核复现脚本

解决步骤

1. 确认问题：在 vLLM 环境中运行内核复现脚本 repro_matmul_ogs.py，观察是否产生 NaN。若设置 FILL_VALUE=0xfe 或 FILL_VALUE=0x0 后不再出现 NaN，则确认问题。
2. （可优先尝试）应用临时补丁：修改 vllm/third_party/triton_kernels/matmul_ogs_details/_matmul_ogs.py 文件，将 tl.load(WMxScalePtrs) 替换为带掩码的加载方式。具体补丁内容参考 Issue 中给出的 diff：

# 在文件 _matmul_ogs.py 中，找到以下代码行：
w_scales = unswizzle_mxfp4_scale_hopper(tl.load(WMxScalePtrs), mx_axis=1, num_warps=num_warps)

# 替换为：
if EVEN_K:
    hopper_scale_mask = tl.full([PACKED_MX_BLOCK], True, dtype=tl.int1)
else:
    hopper_scale_mask = (offs_k_scale // 32) * MX_PACK_DIVISOR < k
w_scales = unswizzle_mxfp4_scale_hopper(
    tl.load(WMxScalePtrs, mask=hopper_scale_mask[None, :], other=0),
    mx_axis=1,
    num_warps=num_warps,
)

3. 更新至修复版本：等待 Triton v3.8.0+ 正式发布，并更新 vLLM 中对应的 vendored 内核。注意新版本可能涉及接口变更（如 matmul 接口重命名），需要额外代码适配。
4. 切换 MoE 内核：作为临时工作区，可以切换到非 Triton 的 MoE 内核（例如使用其他 backend），以避免触发该问题。

验证方法

重新运行内核复现脚本 repro_matmul_ogs.py，确保设置任意 FILL_VALUE（包括默认的 0xff）时输出均不含 NaN。或者运行端到端评估脚本（如 eval_offline_harmony.py），确认准确率恢复正常（例如从 0% 提升至 100%）。

参考来源

vllm-project/vllm #47303

AI 工具推荐

想把多个 AI 模型放在一个入口？

GamsGo AI 集成 ChatGPT、DeepSeek、Gemini、Claude、Midjourney、Veo 等常用模型，适合写作、绘图、视频和日常 AI 工作流。

了解 GamsGo AI

推广链接：通过此链接购买，我可能获得佣金，不影响你的价格。