PR 正文仅标注为 revert #48597，未明确说明具体原因。推测回滚是由于原始优化引入了稳定性、兼容性或回归问题。
建议阅读原始 PR #48597 以理解回滚的上下文。对于关注 Blackwell 性能的用户，可关注后续更稳定的优化合并。此 PR 本身无深入洞察价值，但可作为回滚操作的参考案例。
无代码 review 讨论记录。
撤销内核层优化：在 vllm/models/deepseek_v32/nvidia/kernels.py 中移除 PDL（依赖延迟发射）支持、cuTEDSL 融合 Q 自定义操作以及 fused_norm_rope 的相关注册，恢复为纯 Triton 内核实现。
删除特殊 GEMM 分发：从 vllm/models/deepseek_v32/nvidia/attention.py 和 vllm/models/deepseek_v32/nvidia/mtp.py 中移除 bf16_skinny_gemm 和 dsv3_fused_a_gemm 的硬件特定分发逻辑，以及相关的自定义操作注册；同时移除 vllm/_custom_ops.py 中的对应符号。
回退推测解码生命周期钩子：在 vllm/v1/worker/gpu/spec_decode/autoregressive/speculator.py 和对应的 MTP 子类中删除 on_prefill_begin/end、on_multi_step_decode_begin/end 钩子，这些钩子曾用于控制 indexer 跳过和 topk 缓存管理；同时调整 token 计数变量以消除跟踪差异。
清理配套文件：移除 vllm/v1/attention/backends/mla/sparse_utils.py 中的 phys_shadow 注册机制、删除测试文件 tests/kernels/test_fused_deepseek_v32_norm_rope.py 和 tests/kernels/test_bf16_skinny_gemm.py，以及对 Dockerfile、编译配置等基础设施的同步回退。
| 文件 | 模块 | 状态 | 重要度 |
|---|---|---|---|
vllm/models/deepseek_v32/nvidia/kernels.py |
内核层 | modified | 9.21 |
vllm/v1/worker/gpu/spec_decode/mtp/speculator.py |
推测解码 | modified | 8.12 |
vllm/models/deepseek_v32/nvidia/ops/fused_q_cutedsl.py |
NVIDIA 模型 | removed | 8.02 |
vllm/models/deepseek_v32/nvidia/kernels.py
data-contract
移除 cuTEDSL 融合 Q 算子、PDL 支持及相关自定义操作注册，显著简化内核接口。
# vllm/models/deepseek_v32/nvidia/kernels.py ( 回滚后 )
import torch
from vllm.triton_utils import tl, triton
# 原先的 current_platform、has_cutedsl、direct_register_custom_op 导入已被移除
_DUMMY_CACHE: dict[tuple, torch.Tensor] = {} # 用于 has_indexer=False 路径的虚拟张量缓存
def _dummy(shape: tuple, dtype: torch.dtype, device: torch.device) -> torch.Tensor:
key = (shape, dtype, device)
t = _DUMMY_CACHE.get(key)
if t is None:
t = torch.empty(shape, dtype=dtype, device=device)
_DUMMY_CACHE[key] = t
return t
# _can_use_fused_q_cutedsl、_is_arch_support_pdl、_fused_q_cutedsl_impl、_fused_q_cutedsl_fake 及对应的
# direct_register_custom_op 调用已被完全删除。
# 控制流 : 之前根据这些函数的返回值决定是否使用 cuTEDSL 内核，现在统一使用 Triton 实现。
@triton.jit
def _rms_norm(x, w, eps, HIDDEN_SIZE: tl.constexpr):
# ... 保持不变
@triton.jit
def _get_cos_sin(
cos_sin_cache_ptr,
cos_sin_cache_stride,
pos,
HALF_ROT_DIM: tl.constexpr,
):
# ... 保持不变
@triton.jit
def _fused_norm_rope_kernel(
# ... 参数列表
USE_PDL: tl.constexpr, # 之前存在，现在删除
):
# 原先 USE_PDL 路径下的 gdc_wait/gdc_launch_dependents 已被移除
# 控制流 : 删除 pid == 2 分支（Q RMS norm）以提高性能？但可能影响精度
# ... 其余逻辑
vllm/v1/worker/gpu/spec_decode/mtp/speculator.py
core-logic
删除 MTP 特有的 index_share_for_mtp_iteration 逻辑和 skip-topk 共享机制，简化提案器实现。
# vllm/v1/worker/gpu/spec_decode/mtp/speculator.py ( 回滚后 )
import torch.nn as nn
from vllm.v1.worker.gpu.spec_decode.autoregressive.speculator import (
AutoRegressiveSpeculator,
)
from vllm.v1.worker.gpu.spec_decode.eagle.utils import load_eagle_model
class MTPSpeculator(AutoRegressiveSpeculator):
def load_draft_model(
self,
target_model: nn.Module,
target_attn_layer_names: set[str],
) -> nn.Module:
# 原始版本中会读取 draft_hf_config.index_share_for_mtp_iteration
# 并设置 self.share_mtp_topk_indices，现在直接加载
return load_eagle_model(target_model, self.vllm_config)
# __init__、on_prefill_end、on_multi_step_decode_begin、on_multi_step_decode_end 等方法被移除。
# 控制流 : 原先通过这些钩子控制 indexer 跳过和 topk 缓存压缩，现在回归到自动回归版本的默认行为。
当前评论区没有形成足够清晰的争议点或结论，后续有更多讨论时会体现在这里。
回滚本身风险较低，已通过 CI 验证。但删除了重要的性能优化，可能对 Blackwell GPU 上 DeepSeek V3.2 系列模型的解码延迟产生 10-20% 的负面影响，具体取决于算子。另外，由于 revert 仅撤销代码，未填补因优化暴露的潜在设计缺陷，未来若重新引入需注意兼容性。
用户：使用 Blackwell GPU 运行 DeepSeek V3.2 或 GLM-5.2 模型的用户将失去解码延迟优化，性能回退至优化前水平。系统：代码库回归到优化前的状态，移除约 2100 行新增代码，模块间依赖更加简单。团队：后续迭代需重新评估是否以不同方式合并相关优化，对 NVIDIA 模型内核团队有短期回溯成本。
参与讨论