执行摘要

• 一句话：新增 DeepSeek V4 完整模型支持
• 推荐动作：值得深入精读。该 PR 展示了大规模模型集成的完整流程，特别关注 deepseek_v4_attention.py 中的 MLA 实现、mhc.py 的 TileLang kernel 设计，以及量化策略的权衡。评审中关于 API 设计和硬件兼容性的讨论也值得借鉴。

功能与动机

DeepSeek V4 是 DeepSeek 系列的最新模型，拥有全新架构（如压缩 KV cache 的 MLA、MegaMoE、MXFP4 量化等），需要 vLLM 提供原生支持。关联 Issue #40902 跟踪完整 Roadmap，本 PR 实现核心模型支持，包括 FP4 Indexer、MegaMoE 初始支持和 MTP 推测解码。

实现拆解

1. 模型核心定义：新增 deepseek_v4.py，包含 DeepseekV4ForCausalLM 类，继承 DeepseekV2 但重写量化配置（DeepseekV4FP8Config，MoE 层路由到 MXFP4）、MegaMoE 门控、权重加载映射。
2. MLA 注意力层：deepseek_v4_attention.py 定义 DeepseekV4MultiHeadLatentAttentionWrapper，集成 Indexer 和 Sparse SWA 后端，调用 fused kernel（RMSNorm+Rope+Quant）。
3. 多 Token 预测：deepseek_v4_mtp.py 实现 MTP draft model，包含独立的 e_proj/h_proj 线性层和 hc_head。
4. 超压缩模块：mhc.py 使用 TileLang 实现深度融合的 pre/post mHC kernel。
5. KV cache 压缩后端：deepseek_compressor.py 实现 CompressorBackend，管理压缩 KV cache 的元数据构建。
6. 专用 Tokenizer：deepseek_v4_encoding.py 支持 DSML 协议、工具调用编码/解码和思考模式。
7. MXFP4/MegaMoE 量化集成：修改 mxfp4.py、oracle/mxfp4.py、deep_gemm_moe.py 等，支持 MXFP4 MoE 内核选择。
8. KV cache 配置扩展：kv_cache_utils.py 添加 V4 专用的 block size 解析和组统一逻辑。
9. 模型注册与测试：更新 model_registry.py 和 tests/models/registry.py，但测试限于 Hopper/Blackwell GPU。
10. 基础设施：更新 requirements/cuda.txt 添加 tilelang 依赖，新增 csrc CUDA kernel（如 topk_softplus_sqrt）。

关键文件：

• vllm/model_executor/models/deepseek_v4.py（模块 模型层；类别 source；类型 data-contract；符号 DeepseekV4FP8Config, init, get_name, override_quantization_method）: 核心模型定义，包含量化配置、MegaMoE 门控、权重映射等。
• vllm/model_executor/layers/deepseek_v4_attention.py（模块 注意力层；类别 source；类型 data-contract；符号 DeepseekV4MLAModules, DeepseekV4MultiHeadLatentAttentionWrapper, init, forward）: 实现 MLA 注意力机制，包含 fused kernel 调用、Indexer 和 Sparse SWA 后端。
• vllm/model_executor/models/deepseek_v4_mtp.py（模块 推测解码；类别 source；类型 data-contract；符号 DeepSeekV4MultiTokenPredictorLayer, init, forward, DeepSeekV4MultiTokenPredictor）: Multi-Token Prediction 推测解码 draft model，支持 V4 特有架构。
• vllm/model_executor/layers/mhc.py（模块 融合内核；类别 source；类型 data-contract；符号 compute_num_split, mhc_pre_big_fuse_tilelang, mhc_pre, _mhc_pre_fake）: Hyper-Compression 模块，使用 TileLang 实现深度融合 kernel。
• vllm/model_executor/layers/deepseek_compressor.py（模块 缓存层；类别 source；类型 data-contract；符号 CompressorBackend, init, get_name, get_supported_kernel_block_sizes）: KV cache 压缩后端，实现 CompressorBackend 和 MetadataBuilder。
• vllm/tokenizers/deepseek_v4_encoding.py（模块 分词器；类别 source；类型 dependency-wiring；符号 to_json, tools_from_openai_format, tool_calls_from_openai_format, tool_calls_to_openai_format）: DeepSeek V4 专用 tokenizer，支持 DSML、工具调用、思考模式。
• vllm/v1/core/kv_cache_utils.py（模块 缓存管理；类别 source；类型 dependency-wiring；符号 resolve_kv_cache_block_sizes, _get_kv_cache_config_deepseek_v4, group_and_unify_kv_cache_specs, _approximate_gcd）: KV cache 配置扩展，支持 DeepSeek V4 的 block size 和分组逻辑。
• tests/models/registry.py（模块 测试注册；类别 test；类型 test-coverage）: 模型注册表添加 DeepSeek V4 测试条目（硬件限制）。

关键符号：DeepseekV4FP8Config.get_quant_method, DeepseekV4MultiHeadLatentAttentionWrapper.attention_impl, DeepSeekV4MultiTokenPredictorLayer.forward, CompressorBackend.get_kv_cache_shape, mhc_pre_big_fuse_tilelang, _deepseek_v4_stage_mega_moe_inputs_kernel, tools_from_openai_format, hc_head

评论区精华

1. CUDA/ROCm 兼容性：tjtanaa 建议在 topk_softplus_sqrt_kernels.cu 中使用 VLLM_SHFL_XOR_SYNC_WIDTH 宏，jeejeelee 同意。
2. 依赖版本：mgoin 建议在 requirements/cuda.txt 中添加 tilelang 和 apache-tvm-ffi 版本下限，尚未解决。
3. 测试硬件限制：mgoin 担心测试在非 H100/B200 上失败，jeejeelee 已通过 GPU 能力检查禁用（仅 >= 9.0 启用）。
4. MegaMoE 后端选择：mgoin 建议将 VLLM_DEEPSEEK_V4_USE_MEGA_MOE 集成到 --moe_backend 参数，WoosukKwon 解释当前未使用 fused moe 抽象，zyongye 表示后续改进。
5. Oracle 选择器重复：BowenBao 质疑 select_gpt_oss_mxfp4_moe_backend 和 select_mxfp4_moe_backend 应合并，未明确解决。

• CUDA/ROCm 兼容性：topk softplus sqrt kernel 中 shuffle 宏的使用 (correctness): 已接受建议，后续采用宏替换。
• 依赖版本：tilelang 和 apache-tvm-ffi 需要版本下限 (other): 待处理，未在 PR 中解决。
• 测试硬件限制：DeepSeek V4 测试仅限 Hopper/Blackwell (testing): 已通过条件判断解决。
• MegaMoE 后端选择：VLLM_DEEPSEEK_V4_USE_MEGA_MOE 应集成到 moe_backend 参数 (design): 暂维持 env 变量，未来集成到 moe_backend。
• Oracle 选择器重复：select_gpt_oss_mxfp4_moe_backend 与 select_mxfp4_moe_backend 应合并 (design): 未明确解决，待后续讨论。

风险与影响

• 风险：
  • 硬件兼容性：新 kernel 主要针对 Hopper/Blackwell（SM >= 90），旧 GPU 可能失败，虽有 fallback 但测试不足。
  • 依赖风险：新增 tilelang 和 apache-tvm-ffi 依赖，版本兼容性未经充分验证。
  • 正确性：MegaMoE 和压缩 KV cache 逻辑复杂，边界情况（如部分序列长度）可能存在错误。
  • 性能退化：MTP 推测解码与现有调度器交互可能引入死锁或额外开销。
  • 测试覆盖：大量新代码缺少单元测试，仅少量 e2e 测试且硬件受限。
• 影响：
  • 用户：可部署 DeepSeek V4 系列模型（Base/Flash/Pro），获得完整功能支持，包括工具调用和推测解码。
  • 系统：增加约 16k 行代码，新依赖 tilelang 增加构建时间和二进制体积。
  • 团队：后续需持续优化 MegaMoE kernel、修复兼容性问题，管理 Roadmap 中未完成项（如 PD offload、长上下文支持）。
  • 风险标记：新模型架构重大变更, 依赖 tilelang 版本风险, 测试硬件受限, MegaMoE 正确性待验证

关联脉络

• PR #40760 [Feat] DeepSeek V4: 当前 PR #40860 为该 PR 的 rebase 版本，包含相同变更集。
• PR #40902 [Roadmap] DeepSeek V4: Roadmap Issue 跟踪 V4 实现的完整计划，当前 PR 是其中一部分。
• PR #40923 [Kernel] Marlin MoE: include SM 12.x in default arch list: 后续修复，为 SM 12x 添加 Marlin MoE 架构支持（由 tonyliu312 在评论中提及）。
• PR #40833 MegaMoE 继续优化工作: Roadmap 中提到的 MegaMoE 后续优化 PR。