前情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走完整PythonLORA_DBG1日志环境变量数值探针不只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后埋探针 hsNaN - 上游层问题 hs正常 某projection输出NaN - 该projection的wrapper问题 hs正常 所有projection正常 输出NaN - attention或FFN计算问题 全正常但深层NaN - 逐层看hs 找第一个崩点根因演进假设驱动的排查提出假设设计实验验证或证伪缩小范围下一轮。实际做了5轮修正轮次假设怎么验证结果1GDN状态被LoRA污染零delta adapter证伪 和LoRA改了什么无关2FlashInfer GDN kernel不兼容换Triton backend证伪 不是GDN prefill的问题3in_proj_z的wrapper产生NaN代码注入探针逐层看数值确认 in_proj_z的apply产生NaN4绕过in_proj_z后 另一处也NaN绕过wrapper后继续探针确认 qkv_proj的apply也产生NaN5add_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_B0 则delta0 outputbase(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全为零delta0等价于不应用LoRA。结果仍然输出感叹号。推翻问题不在LoRA权重的数值里。假设4: PR #36976拆分路径的in_proj_z产生NaNPR #36976把GDN层融合的in_proj_qkvz拆成in_proj_qkvin_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}: fshape{list(t.shape)}, mean{mean_val:.6f}, std{std_val:.6f}, fhas_nan{has_nan}, has_inf{has_inf} )vLLM默认用torch.compileCUDA graph传统print会崩溃或静默不输出。#必须按照这个参数才会打印日志 CUDA_VISIBLE_DEVICES3 LORA_DBG1 vllm serve ... --enforce-eager探针拿到真实请求layer 0的数值[Layer0] input hidden_states: mean0.001234, std0.987654, has_nanFalse [Layer0] after in_proj_z: meannan, stdnan, has_nanTrue [Layer0] after attention: meannan, stdnan, has_nanTrueLayer 0输入干净。in_proj_z之后立刻NaN。但in_proj_z的base weight正常(mean0.000018, std0.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_linearvLLM用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, max0.5, mean0.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, max0.5, mean0.01) # 正常最终根因两个独立BugBug 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 1LoRA不生效但输出正常lora_a/b_stacked0 delta0 退化为普通推理只有Bug 2随机触发NaN真实权重被加载但Triton kernel在特定条件污染bufferBug 1Bug 2100%稳定复现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是否被适配器覆盖是否参与NaNin_proj_zColumnParallelLinearColumnParallelLinearWithLoRA否(权重0)Bug 2 → NaNin_proj_qkvMergedColumnParallelLinear不包装(3路output≥3)否安全in_proj_baMergedColumnParallelLinearMergedColumnParallelLinearWithLoRA否(权重0)Bug 2 → NaNqkv_projQKVParallelLinearQKVParallelLinearWithLoRA是(有真实权重)Bug 2 → NaNo_projRowParallelLinearRowParallelLinearWithLoRA是(有真实权重)Bug 2 → NaNGDN层的in_proj_z和in_proj_ba被LoRA wrapper包装但适配器不覆盖它们权重全为零这是Bug 1的内容。任何经过add_lora_linear的模块无论权重是真值还是零都会被Bug 2污染。为什么长序列触发而短序列不触发二分法测试稳定阈值是序列总token数约400。短序列(token400)时GDN内部状态更新次数少NaN虽然产生了但幅度不够大最终output_embedding没有NaN tokenizer正常解码。长序列(token400)时前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_qkvin_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), dtypetorch.float32, devicex.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乘第一步xA^T的结果上 和乘最后等价 减少第二步bufB^T中的数值范围问题.to(y.dtype)保证输出dtype和y一致 否则报错性能影响LoRA delta: rank128 hidden4096 每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零输入天然输出零。最终验证维度条件结果GPUH20 SM90通过LoRA rank1024通过输入长度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}写到最后