[Bug]: LoRA `lora_expand` Triton kernel outputs NaN on Hopper (sm_90) with `block_n=128`, producing garbled LoRA output

问题场景

在 Hopper 架构（sm_90，如 H20/H100）上用 vLLM 加载 Qwen 系列的 LoRA 适配器（如 Qwen3-VL-8B-Instruct、Qwen3.6-35B-A3B MoE），基础模型不加 LoRA 时正常，但带 LoRA 后输出完全乱码，出现重复感叹号、混合语言或代码样式 token。长 prompt 场景（超过约 128 token）更容易暴露：短 prompt 时 LoRA 看起来正常，长 prompt 后输出变成 !!!，或 logprobs 返回 NaN 导致接口报错。同一适配器用 HF Transformers + PEFT 加载时输出正确，说明问题出在 vLLM 的 Triton kernel 路径。用户用 VLLM_ENABLE_V1_MULTIPROCESSING=0 并 monkey-patch 层 forward 定位到首个 NaN 出现在 layer 6 的 GDN 线性注意力，且 NaN 源于 LoRA lora_expand kernel，lora_shrink 正常。

报错原文

[Bug]: LoRA `lora_expand` Triton kernel outputs NaN on Hopper (sm_90) with `block_n=128`, producing garbled LoRA output
{"error":"Out of range float values are not JSON compliant: nan"}

原因分析

已确认的根因是 Triton/ptxas 的已知问题（triton-lang/triton#9433）。在 add_lora_linear 路径中，lora_expand 收到的是 fp32 的 shrink buffer，kernel 需要在核内做类型转换（CAST_TYPE=True），这会让 wgmma 从寄存器读取 A 操作数；当 num_stages ≥ 2 时，下一次迭代开始前上一次迭代的第二个 wgmma 仍未完成，旧于 13.3 的 ptxas 会复用其 A 寄存器给下一轮 shared memory 加载，HGMMA 把 fp32 位当作 bf16 读取，得到约 1e37 和 NaN。单 slice expand 默认 block_n=128、num_stages=2、num_warps=4 恰好命中全部条件。触发条件需同时满足：SM90、fp32 expand 输入、rank ≥ 64、num_stages ≥ 2、token 数足够多（H20 上约 T≈256–384 起，decode 阶段不受影响）。纯 bf16 输入的最小复现脚本反而不会触发，因为该路径正常。

环境排查

• 确认显卡架构是否为 sm_90（H20/H100 等 Hopper），Blackwell（如 RTX 5070，sm_120）不触发此问题。
• 确认 vLLM 版本：0.23.0 / 0.24.0 / 0.25.0 / 0.26.0 均可复现，升级无效。
• 确认 PyTorch、CUDA、triton 版本：Issue 复现环境为 PyTorch 2.11.0+cu130、CUDA 13.0.88、triton 3.6.0；main 上 torch 2.14、triton 3.8.0 也可复现。
• 确认 LoRA 配置：rank ≥ 64、多 LoRA key（comments 中为 780 keys）更容易命中；attention 的 out_proj 等单 slice 投影走 block_n=128 路径。
• 排查时可用 VLLM_ENABLE_V1_MULTIPROCESSING=0 并 monkey-patch 目标层 forward，定位首个 NaN 出现的层和 token 位置。
• 已排除的因素：--enforce-eager 开关、FlashAttn v2/v3/TORCH_SDPA 后端、flashinfer sampler 开关、bf16/fp8/bitsandbytes-4bit 量化、HAMi/vGPU 虚拟化、清空 ~/.triton/cache，均无法避免。

解决步骤

1. 打开 vllm/lora/ops/triton_ops/utils.py，找到 get_lora_op_configs 中设置 expand 默认配置的位置。
2. 把 "block_n": 64 if num_slices > 1 else 128, 改为 "block_n": 64 if num_slices > 1 else 32,，保留多 slice 分支的 64 不变。
3. 重新启动 vLLM 服务，确保加载的是修改后的代码路径，并清空或确认没有加载旧版缓存。
4. 用之前会触发乱码的长 prompt（例如超过 128 token，含约 2335 字符的 SciCode 风格 prompt）重新发起请求。
5. 如果暂不修改源码，可优先尝试在排错阶段用纯 bf16 输入走 reference 路径对比，确认 NaN 只出现在 fp32 expand 输入路径，以缩小问题范围。

验证方法

先用 fp32 输入的最小复现脚本调用 lora_expand，检查输出中 NaN 数量和与 torch reference 的偏差是否归零（Issue 中复现时 4.2M 元素里有数千到上万 NaN、数十万偏差，改 32 后应不再出现）。再端到端验证：长 prompt 不再输出 !!!，/v1/completions 带 logprobs 不再返回 NaN JSON 错误；LoRA 效果探针中输出与 base model 有差异；用同一测试集与 HF Transformers + PEFT 基线对比，Issue 中修复后 SciCode v51 训练集 50 样本得到 54.0% exact match、95.1% 相似度，与基线（56.0% / 95.3%）接近，约 2% 差距属 bf16 精度差异。

参考来源

vllm-project/vllm #48590

AI 工具推荐

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

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

了解 GamsGo AI

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