vLLM部署Qwen3.6-27B hybrid模型启用LoRA，长序列和视频输入产生感叹号输出。关掉LoRA正常。问题在LoRA动态加载机制。tokenizer、prompt、解码参数都排除了，嫌疑锁定在模型推理本身产生了异常hidden state。本篇是排查过程全记录，一共修正了5次对根因的判断。

前情

vLLM部署Qwen3.6-27B hybrid模型启用LoRA 短文本正常长序列和视频输入输出感叹号。关掉LoRA正常。问题在LoRA动态加载机制上。API返回200没有error，没有warning。短文本正常，长文本和视频触发，合并模型正常。状态一旦被污染后续所有请求全坏，重启才能恢复。

本篇是排查全过程以及最终修复方案。

调试工具链

代码注入

直接在site-packages改vLLM源码插入诊断代码。

vLLM三层加速让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会启动一次自校验。

筛选真实请求: grep '[DBG]' log | grep -v 'hs=(0.0, 0.0,'逐层NaN追踪决策树：

根因演进

假设驱动的排查：提出假设，设计实验，验证或证伪，缩小范围，下一轮。

实际做了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权重全部置零。

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。

结果：仍然输出感叹号，不是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。

vLLM默认用torch.compile+CUDA graph传统print会崩溃或静默不输出。

探针拿到真实请求layer 0的数值：

Layer 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。

绕过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。

结果：视频输出正常JSON 反复5遍都正常。模型forward本身正常。LoRA delta的计算过程 add_lora_linear是NaN源。

进一步验证——手动把lora_A和lora_B都置零再调add_lora_linear:

输入全零，权重全零，输出NaN。punica的Triton kernel在H20上存在bug。

最终根因：两个独立Bug

场景	现象	原因
只有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架构

Qwen3.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推理被污染。

状态污染也是同样的机制——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路径：

修复后PyTorch matmul路径：

关键细节：

.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一致否则+=报错

性能影响

Triton kernel比PyTorch快10倍整体吞吐差异0.0018% 淹没在GPU调度噪声里。

torch.compile兼容性

第一次修复尝试：if la.abs().sum() == 0: return——在--enforce-eager下完美，去掉后启动崩溃。

Dynamo在第一次执行时追踪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
多轮	长文本后短文本	正常无状态污染

修复后视频输出：

写到最后

是在往前走就好 bothsavage.github.io

将进酒

vLLM下LoRA挂载感叹号解决路程-1-背景介绍