前情

在部署基于 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" -delete

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：日志环境变量

数值探针

不只是简单打印文本，而是打印 (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 状态污染

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=True

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：

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 中任意一个单独存在都不会导致如此严重的问题。两个叠加，才让系统变得完全不可用。

源码分析

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_B

Qwen3.5-27B GDN 层的 LoRA 包装

GDN 层的 in_proj_z 和 in_proj_ba 被 LoRA wrapper 包装，但 adapter 不覆盖它们，权重全为零，这就是 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 time

Dynamo 在第一次执行时追踪 Python 代码路径，编译成优化计算图。if 依赖 tensor 运行时值，Dynamo 不知道该走哪个分支，直接拒绝编译。vLLM 默认开启 torch.compile，只在 --enforce-eager 下通过的修改等于没通过。

正确做法是不做数据依赖分支，依赖数学性质——torch.matmul 零输入天然输出零。

最终验证

修复后视频输出：

{
  "objects":["ceramic jar","bronze vessel","jade ornament"],
  "chamber_type":"primary burial chamber",
  "estimated_period":"Warring States"
}

写到最后

这次排查的教训其实很直白：两个看似独立的 bug，单独存在时都只是小问题，但叠加在一起就变成了灾难。Bug 1 让 LoRA 权重未被正确加载，本该是微调后的权重变成了零；Bug 2 让即使权重为零，Triton kernel 也输出 NaN。任何一个先修好，问题都不会这么棘手。但现实就是，两个 bug 同时撞上了。

这也说明了一个道理：调试时不要只盯着自己假设的单点问题，有时候系统性的问题需要从整体链路去排查。而 vLLM 的 LoRA 动态加载机制，虽然功能强大，但底层模块间的耦合确实是个需要格外小心的区域。