前情
vLLM部署Qwen3.6-27B hybrid模型启用LoRA 短文本正常 长序列和视频输入输出感叹号。关掉LoRA正常。问题在LoRA动态加载机制上。API返回200没有error，没有warning。短文本正常，长文本和视频触发，合并模型正常。状态一旦被污染后续所有请求全坏，重启才能恢复。
本篇是排查全过程以及最终修复方案。
#正常输出{"objects":["ceramic jar","bronze vessel"],"chamber_type":"primary burial chamber"}#异常："!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!"调试工具链
代码注入
直接在site-packages改vLLM源码插入诊断代码。
pip show vllm | grep Location cp target.py target.py.bak vim target.py #改完记得清pyc缓存，要不然不会生效 find site-packages/vllm -name "*.pyc" -deletevLLM三层加速让print失效
torch.compile把Python编译成优化计算图，print不存在于图里，所以导致CUDA Graph重放不走print。Triton kernel手写GPU汇编 完全不可观测。
三件套：
- torch.compiler.is_compiling()，guard 编译时跳过print块
- -enforce-eager，关闭CUDA graph forward走完整Python
- LORA_DBG=1，日志环境变量
数值探针
不只print文本 打印(mean, std, has_nan, has_inf)四元组。has_nan定位异常 has_inf检测溢出 mean/std发现值漂移早期信号。需要区分profile run——vLLM启动时用全零tensor会启动一次自校验。
def is_profile_run(hidden_states): return (hidden_states.abs().sum() < 0.01).item()筛选真实请求: grep ‘[DBG]’ log | grep -v 'hs=(0.0, 0.0,'逐层NaN追踪决策树：
在某层forward后埋探针 hs=NaN -> 上游层问题 hs正常 某projection输出NaN -> 该projection的wrapper问题 hs正常 所有projection正常 输出NaN -> attention或FFN计算问题 全正常但深层NaN -> 逐层看hs 找第一个崩点根因演进
假设驱动的排查：提出假设，设计实验，验证或证伪，缩小范围，下一轮。
实际做了5轮修正：
| 轮次 | 假设 | 怎么验证 | 结果 |
|---|---|---|---|
| 1 | GDN状态被LoRA污染 | 零delta adapter | 证伪 和LoRA改了什么无关 |
| 2 | FlashInfer GDN kernel不兼容 | 换Triton backend | 证伪 不是GDN prefill的问题 |
| 3 | in_proj_z的wrapper产生NaN | 代码注入探针逐层看数值 | 确认 in_proj_z的apply产生NaN |
| 4 | 绕过in_proj_z后 另一处也NaN | 绕过wrapper后继续探针 | 确认 qkv_proj的apply也产生NaN |
| 5 | add_lora_linear是NaN源 | 纯base patch 完全绕过LoRA delta | 确认 而且零权重+零输入也NaN |
假设1: GDN状态污染
GDN层维护conv_state和ssm_state两个内部状态 在长序列中持续累积。猜测LoRA修改了8层attention的QKV输出后 意外值喂给GDN层，导致状态累积偏差，长序列时雪崩。
验证方法：零delta adapter。把lora_B权重全部置零。
for name, param in lora_params.items(): if 'lora_B' in name: param.data.zero_()LoRA公式: output = base(x) + lora_B @ lora_A @ x * scaling。lora_B=0 则delta=0 output=base(x) 数学上和不用LoRA等价。但代码路径不同——LoRA wrapper仍然会走一遍add_lora_linear。
结果：仍然输出感叹号。推翻，问题在LoRA的存在本身，和它改了什么内容无关。
假设2: FlashInfer GDN kernel不兼容
vLLM用FlashInfer作为默认GPU kernel库。FlashInfer的GDN prefill kernel是JIT编译的，怀疑它在处理LoRA拆分路径后的数据布局时有数值问题。
验证方法：切换到Triton backend。
--gdn-prefill-backend triton结果：仍然输出感叹号，不是FlashInfer的问题。
假设3: LoRA delta数值错误
怀疑LoRA adapter的权重(lora_a x lora_b)产生了异常delta 叠加到base weight后输出NaN。
验证方法就是零delta adapter: lora_a和lora_b全为零，delta=0等价于不应用LoRA。
结果：仍然输出感叹号。推翻，问题不在LoRA权重的数值里。
假设4: PR #36976拆分路径的in_proj_z产生NaN
PR #36976把GDN层融合的in_proj_qkvz拆成in_proj_qkv+in_proj_z。怀疑in_proj_z被ColumnParallelLinearWithLoRA错误包装后apply产生NaN。
验证方法：代码注入探针 在forward_cuda的每个中间张量上打印mean/std/has_nan/has_inf。
def debug_tensor(t, name, layer_idx): if not torch.compiler.is_compiling() and os.environ.get('LORA_DBG'): t_float = t.float() has_nan = torch.isnan(t_float).any().item() has_inf = torch.isinf(t_float).any().item() mean_val = t_float.mean().item() if not has_nan else float('nan') std_val = t_float.std().item() if not has_nan else float('nan') logger.warning( f"[Layer{layer_idx}]{name}: " f"shape={list(t.shape)}, mean={mean_val:.6f}, std={std_val:.6f}, " f"has_nan={has_nan}, has_inf={has_inf}" )vLLM默认用torch.compile+CUDA graph传统print会崩溃或静默不输出。
#必须按照这个参数才会打印日志 CUDA_VISIBLE_DEVICES=3 LORA_DBG=1 vllm serve ... --enforce-eager探针拿到真实请求layer 0的数值：
[Layer0] input hidden_states: mean=0.001234, std=0.987654, has_nan=False [Layer0] after in_proj_z: mean=nan, std=nan, has_nan=True [Layer0] after attention: mean=nan, std=nan, has_nan=TrueLayer 0输入干净。in_proj_z之后立刻NaN。但in_proj_z的base weight正常(mean=0.000018, std=0.016766, no NaN)——权重本身没问题。in_proj_z的类型是ColumnParallelLinearWithLoRA 被LoRA wrapper包装后 apply过程在权重正常时产生了NaN。
进一步验证：绕过wrapper直接调base_layer。
z, _ = self.in_proj_z.base_layer(hidden_states) # 绕过wrapper绕过wrapper后 layer 0/1/2的z不再NaN。但layer 3(full_attention层)的qkv_proj又输出NaN layer 4+全部NaN。layer 3的qkv_proj被QKVParallelLinearWithLoRA包装——另一种wrapper 也在apply里产生NaN。
两种不同wrapper ColumnParallelLinearWithLoRA和QKVParallelLinearWithLoRA 都在apply出问题，怀疑收敛到它们共同调用的add_lora_linear。
假设5: 锁定punica add_lora_linear
vLLM用punica库执行add_lora_linear底层用Triton kernel加速。
纯base patch验证：在LoRA wrapper的forward里完全不碰LoRA delta。
def forward(self, x): return self.base_layer(x) # 完全不调add_lora_linear #注意坑：函数名*apply_base_forward看名字像纯base 但内部又调了*apply_lora_to_output应用LoRA delta。函数名不可信 必须读完整实现。结果：视频输出正常JSON 反复5遍都正常。模型forward本身正常。LoRA delta的计算过程 add_lora_linear是NaN源。
进一步验证——手动把lora_A和lora_B都置零再调add_lora_linear:
result = add_lora_linear(x, lora_A_zero, lora_B_zero, ...) # has_nan = True输入全零，权重全零，输出NaN。punica的Triton kernel在H20上存在bug。
[LORA_DBG] Before add_lora_linear: base_output: tensor([...], min=-0.5, max=0.5, mean=0.01) # 正常 lora_a_stacked: all zeros lora_b_stacked: all zeros [LORA_DBG] After add_lora_linear (Triton): output: tensor([nan, nan, nan, ...]) # NaN! [LORA_DBG] After add_lora_linear (PyTorch matmul): output: tensor([...], min=-0.5, max=0.5, mean=0.01) # 正常最终根因：两个独立Bug
Bug 1: LoRA权重未加载 Bug 2: Triton kernel NaN lora_a_stacked = 0 add_lora_linear lora_b_stacked = 0 Triton kernel在H20上 (应为微调后的权重) 即使权重为零也输出NaN | | +----------- 叠加形成 ---------------+ | LoRA delta = NaN | base_output + NaN = NaN | MambaCache被NaN污染 | 后续所有层输出NaN | lm_head argmax(NaN) → token 0 → "!"| 场景 | 现象 | 原因 |
|---|---|---|
| 只有Bug 1 | LoRA不生效但输出正常 | lora_a/b_stacked=0 delta=0 退化为普通推理 |
| 只有Bug 2 | 随机触发NaN | 真实权重被加载但Triton kernel在特定条件污染buffer |
| Bug 1+Bug 2 | 100%稳定复现 | Bug 1让权重为零 Bug 2让零权重也输出NaN |
两个Bug中任意一个单独存在都不会导致如此严重的问题。两个叠加才完美不可用。
源码分析
vLLM LoRA架构
LoRAModelManager 管理所有已加载的LoRA adapter 负责adapter的激活切换 LoRALayer (vllm/lora/layers/) BaseLayerWithLoRA: 所有LoRA层的基类 ColumnParallelLinearWithLoRA: 列并行层 用于in_proj_z、o_proj等 MergedColumnParallelLinearWithLoRA: 融合列并行层 用于in_proj_ba QKVParallelLinearWithLoRA: QKV并行层 用于full_attention的qkv_proj RowParallelLinearWithLoRA: 行并行层 PunicaWrapper (vllm/lora/punica_wrapper/) punica_gpu.py: GPU加速的LoRA计算 关键函数 add_lora_linear() 用Triton kernel计算 base + lora_A @ lora_BQwen3.5-27B GDN层的LoRA包装
| 模块 | 基础层类型 | LoRA Wrapper | 是否被适配器覆盖 | 是否参与NaN |
|---|---|---|---|---|
| in_proj_z | ColumnParallelLinear | ColumnParallelLinearWithLoRA | 否(权重=0) | Bug 2 → NaN |
| in_proj_qkv | MergedColumnParallelLinear | 不包装(3路output≥3) | 否 | 安全 |
| in_proj_ba | MergedColumnParallelLinear | MergedColumnParallelLinearWithLoRA | 否(权重=0) | Bug 2 → NaN |
| qkv_proj | QKVParallelLinear | QKVParallelLinearWithLoRA | 是(有真实权重) | Bug 2 → NaN |
| o_proj | RowParallelLinear | RowParallelLinearWithLoRA | 是(有真实权重) | Bug 2 → NaN |
GDN层的in_proj_z和in_proj_ba被LoRA wrapper包装但适配器不覆盖它们，权重全为零，这是Bug 1的内容。任何经过add_lora_linear的模块，无论权重是真值还是零，都会被Bug 2污染。
为什么长序列触发而短序列不触发
二分法测试，稳定阈值是序列总token数约400。
短序列(token<400)时GDN内部状态更新次数少，NaN虽然产生了但幅度不够大，最终output_embedding没有NaN tokenizer正常解码。
长序列(token>400)时前400个token让GDN内部状态累积微小异常，第400+个token时状态中的异常值触发kernel bug kernel输出NaN NaN进入GDN的conv_state/ssm_state所有后续step推理被污染。
短序列: [token1] -> GDN -> [state_clean] -> [token2] -> GDN -> [state_clean] 长序列: [token398] -> GDN -> [state_slightly_bad] -> [token399] -> GDN -> [state_worse] -> [token400] -> GDN -> [state_NaN] -> [token401] -> GDN -> [state_NaN] -> ^ 从这一刻起全NaN状态污染也是同样的机制——GDN状态被NaN写入后永久损坏，不重启服务无法清除。
修复方案
Bug 1修复：上游PR #36976 拆分GDN层融合投影in_proj_qkvz为in_proj_qkv+in_proj_z。
Bug 2修复：用PyTorch torch.matmul替换punica的Triton kernel。
为什么Triton kernel会NaN而torch.matmul不会
Triton kernel用torch.empty创建buffer——不初始化内存 buffer内容是GPU上次计算的残留垃圾值。正常流程是Triton kernel随后填充buffer。但当lora_a/lora_b都为0时 kernel的优化路径跳过buffer填充(认为零输入不需要计算) buffer保持垃圾值。后续用垃圾值参与矩阵运算 结果NaN。
torch.empty不用torch.zeros是为了性能——kernel设计假设是一定会填充buffer 不需要预先清零。
H20(Hopper架构 SM90)上特别容易触发：Tensor Core异步Warp调度(Hopper引入wgmma指令) Shared Memory布局变化 TMA(Tensor Memory Accelerator)专用硬件单元 这些新特性使Triton在Hopper上生成的代码执行时序与Ampere不同 零值优化路径更激进。
PyTorch的torch.matmul保证零输入一定产生零输出——矩阵乘法基本数学性质。
修复代码
原始Triton kernel路径：
def add_lora_linear(y, x, lora_a_stacked, lora_b_stacked, scale, ...): buffer = torch.empty((x.shape[0], rank), dtype=torch.float32, device=x.device) self.add_shrink(buffer, x, lora_a_stacked, scale) # 零权重时跳过填充 self.add_expand(y, buffer, lora_b_stacked, ...) # 用垃圾值计算修复后PyTorch matmul路径：
def add_lora_linear(y, x, lora_a_stacked, lora_b_stacked, scale, ...): la = lora_a_stacked[slice_idx].squeeze(1) # (max_loras, rank, input) lb = lora_b_stacked[slice_idx].squeeze(1) # (max_loras, output, rank) for li in range(la.shape[0]): buf = torch.matmul(x.to(torch.float32), la[li].T.to(torch.float32)) * scale delta = torch.matmul(buf, lb[li].T.to(torch.float32)).to(y.dtype) y[..., offset:offset + sz] += delta break关键细节：
- .squeeze(1)去掉stacking维度 lora_a_stacked shape是(num_slices, num_layers, rank, hidden_in)
- .to(torch.float32)保证精度 GPU推理用bf16/fp16 但小数值乘加低精度易累积误差
- scale乘第一步x@A^T的结果上 和乘最后等价 减少第二步buf@B^T中的数值范围问题
- .to(y.dtype)保证输出dtype和y一致 否则+=报错
性能影响
LoRA delta: rank=128 hidden=4096 每token约1M次浮点运算 模型base forward: Qwen3.6-27B 每token约54B次浮点运算 LoRA占比: ~0.002%Triton kernel比PyTorch快10倍 整体吞吐差异0.0018% 淹没在GPU调度噪声里。
torch.compile兼容性
第一次修复尝试：if la.abs().sum() == 0: return——在–enforce-eager下完美，去掉后启动崩溃。
torch._dynamo.exc.Unsupported: Data-dependent branching is not supported The condition `la.abs().sum() == 0` depends on tensor data which Dynamo cannot trace at compile timeDynamo在第一次执行时追踪Python代码路径 编译成优化计算图。if依赖tensor运行时值 Dynamo不知道该走哪个分支 直接拒绝。vLLM默认开启torch.compile 只在–enforce-eager下通过的修改等于没通过。
正确做法：不做数据依赖分支 依赖数学性质 torch.matmul零输入天然输出零。
最终验证
| 维度 | 条件 | 结果 |
|---|---|---|
| GPU | H20 SM90 | 通过 |
| LoRA rank | 1024 | 通过 |
| 输入长度 | 128/1024/4096 | 通过 |
| torch.compile | 开启/关闭 | 通过 |
| 数据类型 | bf16/fp16 | 通过 |
| 短文本 | 约50 tokens | 正常 |
| 长文本 | 约450 tokens | 正常 不再感叹号 |
| 视频 | 约1951 visual tokens | 正常JSON |
| 多轮 | 长文本后短文本 | 正常 无状态污染 |
修复后视频输出：
{"objects":["ceramic jar","bronze vessel","jade ornament"],"chamber_type":"primary burial chamber","estimated_period":"Warring States"}
