[Bug]: The Deepseek v4 model compiles kenrel during the inference process

问题场景

用户在 4 张 NVIDIA H100 上通过 vllm serve 加载 deepseek-ai/DeepSeek-V4-Flash，启用 fp8 KV cache、block size 256、张量并行 4、DeepSeek V4 的 tokenizer/tool-call/reasoning parser，并通过 --speculative-config 开启 MTP 投机解码（num_speculative_tokens=1）。在推理过程中，四个 TP worker 依次开始编译 TileLang kernel（mhc_pre_big_fuse_with_norm_tilelang、mhc_post_tilelang、hc_head_fuse_tilelang 等），随后 EngineCore 报出共享内存广播 60 秒内无可用块的提示。

报错原文

[TileLang:tilelang.jit.kernel:INFO] (kernel.py:133): TileLang begins to compile kernel `mhc_pre_big_fuse_with_norm_tilelang` with `out_idx=None`
[TileLang:tilelang.jit.kernel:INFO] (kernel.py:141): TileLang completes to compile kernel `mhc_pre_big_fuse_with_norm_tilelang`
[TileLang:tilelang.jit.kernel:INFO] (kernel.py:133): TileLang begins to compile kernel `mhc_post_tilelang` with `out_idx=None`
[TileLang:tilelang.jit.kernel:INFO] (kernel.py:141): TileLang completes to compile kernel `mhc_post_tilelang`
(EngineCore pid=900) INFO 07-06 16:04:19 [shm_broadcast.py:705] No available shared memory broadcast block found in 60 seconds. This typically happens when some processes are hanging or doing some time-consuming work (e.g. compilation, weight/kv cache quantization).
[TileLang:tilelang.jit.kernel:INFO] (kernel.py:133): TileLang begins to compile kernel `hc_head_fuse_tilelang` with `out_idx=None`
...

原因分析

根据 Issue 讨论中的分析，DeepSeek V4 Flash 引入了 12 个模型专属 kernel，分布在 Sparse MLA（Triton）、Eagle 投机解码（Triton）、MTP RMSNorm（Triton）和 mHC（TileLang）中，但上游 warmup 覆盖不完整：

• 4 个 Eagle/MTP kernel 完全没有 warmup（eagle_prepare_next_token_padded_kernel、eagle_prepare_inputs_padded_kernel、eagle_step_slot_mapping_metadata_kernel、_mtp_shared_head_rmsnorm_kernel），上游无对应 warmup 文件。
• 4 个 mHC TileLang kernel 仅通过 15 个固定 token size 间接触发，其中 NVIDIA 专用 fused 路径 mhc_fused_post_pre_tilelang 完全未被覆盖，且未产生独立缓存条目。
• 2 个 Sparse MLA kernel 缺少关键特化：_build_c128a_topk_metadata_kernel 未枚举 7 个 int 参数的 int==1 特化、未枚举指针可整除性，且使用 torch.empty 初始化缓冲（残留垃圾值可能构成非法地址），也未加 _is_safe() 越界保护；_combine_topk_swa_indices_kernel 仅覆盖 %16==0 路径。
• 结果：12 个 kernel 总计约 94 个缓存条目，其中 6 个贡献零独立条目；首次请求仍需现场 JIT 编译，导致首请求延迟尖峰，并存在未初始化缓冲与越界导致的潜在崩溃风险。

因此，日志中的 TileLang 编译记录属于“预期外的运行时编译”，而共享内存广播超时是该编译耗时期间的连锁现象。具体是否由某一 kernel 直接导致异常，Issue 中未给出定论，仅作为可能原因。

环境排查

• 确认 vLLM 版本是否已合入 DeepSeek V4 Flash 的 kernel warmup 补丁，尤其关注 vllm/v1/spec_decode/utils.py、vllm/models/deepseek_v4/common/ops/fused_mtp_input_rmsnorm.py、vllm/models/deepseek_v4/sparse_mla.py、vllm/models/deepseek_v4/common/ops/cache_utils.py 是否存在 warmup 调用或特化枚举。
• 确认 TileLang 与 Triton 版本：Issue 未记录具体版本号，需以实际环境为准，避免自行假定。
• 确认 Python、CUDA、PyTorch 版本（Issue 未提供 collect_env.py 输出，建议补跑该脚本以便对照）。
• 确认 GPU 数量与并行配置一致：4 × H100、--tensor-parallel-size 4，并确认每张卡显存与共享内存配置未被限制。
• 确认启动参数中是否启用 MTP 投机解码（--speculative-config），这直接决定是否触发 4 个零 warmup 的 Eagle/MTP kernel。
• 确认是否命中 --block-size 256 与 --kv-cache-dtype fp8 对应的 kernel 特化路径。

解决步骤

1. 先复现并记录首请求延迟与日志：保留完整 TileLang begins to compile kernel 段落，确认编译发生在推理请求期间而非启动 warmup 阶段。
2. 临时移除以 --speculative-config 开头的参数（即关闭 MTP 投机解码），重启服务，观察是否还会出现 Eagle/MTP 相关 kernel 的现场编译与共享内存广播超时。这是用于定位“可优先尝试”的隔离步骤，不是已确认的最终修复。
3. 若关闭投机解码后现场编译消失，说明问题集中在未 warmup 的 Eagle/MTP kernel，需要升级到包含完整 warmup 的 vLLM 版本或应用对应补丁。
4. 若关闭后仍出现 mHC TileLang kernel 编译，检查是否使用了 NVIDIA fused 路径 mhc_fused_post_pre_tilelang；该路径在上游 warmup 中完全未覆盖，需确认版本是否已补充调用点。
5. 若怀疑崩溃由 _build_c128a_topk_metadata_kernel 引起，关注是否存在 int==1 参数或指针未对齐场景，并确认版本是否已加入 _is_safe() 越界保护与缓冲初始化修正。
6. 升级或打补丁后，重新以原始参数启动，先预热一次完整请求，再观察日志中是否不再出现推理期间的 TileLang 编译记录及共享内存广播超时。

验证方法

以原始启动参数运行服务，发送一次或多次推理请求，确认：1）推理日志中不再出现 TileLang begins to compile kernel（尤其 mhc_*、hc_head_fuse_tilelang）；2）不再出现 No available shared memory broadcast block found in 60 seconds；3）首请求延迟与后续请求延迟差异明显收敛；4）服务在 MTP 投机解码开启时能稳定完成多次推断且不发生崩溃。若任一条件仍不满足，说明对应 warmup 路径仍未被覆盖。

参考来源

vllm-project/vllm #48809

AI 工具推荐

想把多个 AI 模型放在一个入口？

GamsGo AI 集成 ChatGPT、DeepSeek、Gemini、Claude、Midjourney、Veo 等常用模型，适合写作、绘图、视频和日常 AI 工作流。

了解 GamsGo AI

推广链接：通过此链接购买，我可能获得佣金，不影响你的价格。