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

[leihuang-sketch]

Your current environment

The output of python collect_env.py

Your output of `python collect_env.py` here

🐛 Describe the bug

Nvidia H100 * 4

test cmd

vllm serve deepseek-ai/DeepSeek-V4-Flash \
  --trust-remote-code \
  --kv-cache-dtype fp8 \
  --block-size 256 \
  --tensor-parallel-size 4 \
  --no-enable-flashinfer-autotune \
  --tokenizer-mode deepseek_v4 \
  --tool-call-parser deepseek_v4 \
  --enable-auto-tool-choice \
  --reasoning-parser deepseek_v4 \
  --speculative-config '{"method":"mtp","num_speculative_tokens":1}'

Exception log

(Worker_TP1 pid=1099) 2026-07-06 16:03:24  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:133): TileLang begins to compile kernel `mhc_pre_big_fuse_with_norm_tilelang` with `out_idx=None`
(Worker_TP0 pid=1098) 2026-07-06 16:03:24  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:133): TileLang begins to compile kernel `mhc_pre_big_fuse_with_norm_tilelang` with `out_idx=None`
(Worker_TP2 pid=1100) 2026-07-06 16:03:24  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:133): TileLang begins to compile kernel `mhc_pre_big_fuse_with_norm_tilelang` with `out_idx=None`
(Worker_TP3 pid=1101) 2026-07-06 16:03:24  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:133): TileLang begins to compile kernel `mhc_pre_big_fuse_with_norm_tilelang` with `out_idx=None`
(Worker_TP1 pid=1099) 2026-07-06 16:03:34  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:141): TileLang completes to compile kernel `mhc_pre_big_fuse_with_norm_tilelang`
(Worker_TP0 pid=1098) 2026-07-06 16:03:34  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:141): TileLang completes to compile kernel `mhc_pre_big_fuse_with_norm_tilelang`
(Worker_TP2 pid=1100) 2026-07-06 16:03:34  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:141): TileLang completes to compile kernel `mhc_pre_big_fuse_with_norm_tilelang`
(Worker_TP3 pid=1101) 2026-07-06 16:03:34  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:141): TileLang completes to compile kernel `mhc_pre_big_fuse_with_norm_tilelang`
(Worker_TP1 pid=1099) 2026-07-06 16:03:48  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:133): TileLang begins to compile kernel `mhc_post_tilelang` with `out_idx=None`
(Worker_TP0 pid=1098) 2026-07-06 16:03:48  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:133): TileLang begins to compile kernel `mhc_post_tilelang` with `out_idx=None`
(Worker_TP3 pid=1101) 2026-07-06 16:03:49  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:133): TileLang begins to compile kernel `mhc_post_tilelang` with `out_idx=None`
(Worker_TP2 pid=1100) 2026-07-06 16:03:49  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:133): TileLang begins to compile kernel `mhc_post_tilelang` with `out_idx=None`
(Worker_TP1 pid=1099) 2026-07-06 16:03:52  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:141): TileLang completes to compile kernel `mhc_post_tilelang`
(Worker_TP0 pid=1098) 2026-07-06 16:03:52  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:141): TileLang completes to compile kernel `mhc_post_tilelang`
(Worker_TP3 pid=1101) 2026-07-06 16:03:53  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:141): TileLang completes to compile kernel `mhc_post_tilelang`
(Worker_TP2 pid=1100) 2026-07-06 16:03:53  [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).
(Worker_TP1 pid=1099) 2026-07-06 16:04:20  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:133): TileLang begins to compile kernel `hc_head_fuse_tilelang` with `out_idx=None`
(Worker_TP0 pid=1098) 2026-07-06 16:04:20  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:133): TileLang begins to compile kernel `hc_head_fuse_tilelang` with `out_idx=None`
(Worker_TP2 pid=1100) 2026-07-06 16:04:20  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:133): TileLang begins to compile kernel `hc_head_fuse_tilelang` with `out_idx=None`
(Worker_TP3 pid=1101) 2026-07-06 16:04:20  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:133): TileLang begins to compile kernel `hc_head_fuse_tilelang` with `out_idx=None`
(Worker_TP0 pid=1098) 2026-07-06 16:04:26  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:141): TileLang completes to compile kernel `hc_head_fuse_tilelang`
(Worker_TP2 pid=1100) 2026-07-06 16:04:26  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:141): TileLang completes to compile kernel `hc_head_fuse_tilelang`
(Worker_TP3 pid=1101) 2026-07-06 16:04:26  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:141): TileLang completes to compile kernel `hc_head_fuse_tilelang`
(Worker_TP1 pid=1099) 2026-07-06 16:04:26  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:141): TileLang completes to compile kernel `hc_head_fuse_tilelang`
(Worker_TP0 pid=1098) INFO 07-06 16:04:40 [indexer.py:276] DSA indexer decode path: use_flattening=False (next_n=2, use_fp4_indexer_cache=False)
(Worker_TP1 pid=1099) INFO 07-06 16:04:43 [gpu_model_runner.py:6483] Profiling CUDA graph memory: PIECEWISE=50 (largest=512), FULL=50 (largest=512)
(Worker_TP0 pid=1098) INFO 07-06 16:04:43 [gpu_model_runner.py:6483] Profiling CUDA graph memory: PIECEWISE=50 (largest=512), FULL=50 (largest=512)
(Worker_TP3 pid=1101) INFO 07-06 16:04:43 [gpu_model_runner.py:6483] Profiling CUDA graph memory: PIECEWISE=50 (largest=512), FULL=50 (largest=512)
(Worker_TP2 pid=1100) INFO 07-06 16:04:43 [gpu_model_runner.py:6483] Profiling CUDA graph memory: PIECEWISE=50 (largest=512), FULL=50 (largest=512)
(Worker_TP1 pid=1099) 2026-07-06 16:04:46  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:133): TileLang begins to compile kernel `mhc_pre_big_fuse_with_norm_tilelang` with `out_idx=None`
(Worker_TP0 pid=1098) 2026-07-06 16:04:46  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:133): TileLang begins to compile kernel `mhc_pre_big_fuse_with_norm_tilelang` with `out_idx=None`
(Worker_TP3 pid=1101) 2026-07-06 16:04:46  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:133): TileLang begins to compile kernel `mhc_pre_big_fuse_with_norm_tilelang` with `out_idx=None`
(Worker_TP2 pid=1100) 2026-07-06 16:04:46  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:133): TileLang begins to compile kernel `mhc_pre_big_fuse_with_norm_tilelang` with `out_idx=None`
(Worker_TP1 pid=1099) 2026-07-06 16:04:56  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:141): TileLang completes to compile kernel `mhc_pre_big_fuse_with_norm_tilelang`
(Worker_TP0 pid=1098) 2026-07-06 16:04:56  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:141): TileLang completes to compile kernel `mhc_pre_big_fuse_with_norm_tilelang`
(Worker_TP3 pid=1101) 2026-07-06 16:04:57  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:141): TileLang completes to compile kernel `mhc_pre_big_fuse_with_norm_tilelang`
(Worker_TP2 pid=1100) 2026-07-06 16:04:57  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:141): TileLang completes to compile kernel `mhc_pre_big_fuse_with_norm_tilelang`

others

(Worker_TP1 pid=387) WARNING 07-14 11:09:35 [jit_monitor.py:126] Triton kernel JIT compilation during inference: _build_c128a_topk_metadata_kernel (constexprs={BLOCK_SIZE=1024}; signature={BLOCK_SIZE='constexpr', block_size='i32', block_table_ptr='*i32', block_table_stride='i32', compress_ratio='i32', decode_lens_ptr='*i32', global_decode_ptr='*i32', global_decode_stride='i32', max_compressed_tokens='i32', num_decode_tokens='i32', positions_ptr='*i64', prefill_local_ptr='*i32', prefill_local_stride='i32', slot_mapping_ptr='*i64', token_to_req_indices_ptr='*i32'}; extra_compile_info={configs=[{(0,): [['tt.divisibility', 16]], (1,): [['tt.divisibility', 16]], (2,): [['tt.divisibility', 16]], (3,): [['tt.divi..., device=1, enable_fp_fusion=True, extern_libs=(('libdevice', '/usr/local/lib/python3.12/dist-packages/triton/backends/nvidia/lib/libdevice.10.bc'),), is_warmup=False, launch_cooperative_grid=False, num_ctas=1, num_stages=3, num_warps=4, specialization_data='{"name": "vllm.models.deepseek_v4.sparse_mla._build_c128a_topk_metadata_kernel", "signature": {"global_decode_ptr": ...}; key="[('*i32', 'D'), ('i32', 'D'), ('*i32', 'D'), ('*i32', 'D'), ('i32', 'D'), ('*i64', 'D'), ('i32', 'D'), ('i32', 'D'),...). This causes a latency spike; consider extending warmup to cover this shape/config.
(Worker_TP2 pid=388) WARNING 07-14 11:09:35 [jit_monitor.py:126] Triton kernel JIT compilation during inference: _build_c128a_topk_metadata_kernel (constexprs={BLOCK_SIZE=1024}; signature={BLOCK_SIZE='constexpr', block_size='i32', block_table_ptr='*i32', block_table_stride='i32', compress_ratio='i32', decode_lens_ptr='*i32', global_decode_ptr='*i32', global_decode_stride='i32', max_compressed_tokens='i32', num_decode_tokens='i32', positions_ptr='*i64', prefill_local_ptr='*i32', prefill_local_stride='i32', slot_mapping_ptr='*i64', token_to_req_indices_ptr='*i32'}; extra_compile_info={configs=[{(0,): [['tt.divisibility', 16]], (1,): [['tt.divisibility', 16]], (2,): [['tt.divisibility', 16]], (3,): [['tt.divi..., device=2, enable_fp_fusion=True, extern_libs=(('libdevice', '/usr/local/lib/python3.12/dist-packages/triton/backends/nvidia/lib/libdevice.10.bc'),), is_warmup=False, launch_cooperative_grid=False, num_ctas=1, num_stages=3, num_warps=4, specialization_data='{"name": "vllm.models.deepseek_v4.sparse_mla._build_c128a_topk_metadata_kernel", "signature": {"global_decode_ptr": ...}; key="[('*i32', 'D'), ('i32', 'D'), ('*i32', 'D'), ('*i32', 'D'), ('i32', 'D'), ('*i64', 'D'), ('i32', 'D'), ('i32', 'D'),...). This causes a latency spike; consider extending warmup to cover this shape/config.
(Worker_TP3 pid=389) WARNING 07-14 11:09:35 [jit_monitor.py:126] Triton kernel JIT compilation during inference: _build_c128a_topk_metadata_kernel (constexprs={BLOCK_SIZE=1024}; signature={BLOCK_SIZE='constexpr', block_size='i32', block_table_ptr='*i32', block_table_stride='i32', compress_ratio='i32', decode_lens_ptr='*i32', global_decode_ptr='*i32', global_decode_stride='i32', max_compressed_tokens='i32', num_decode_tokens='i32', positions_ptr='*i64', prefill_local_ptr='*i32', prefill_local_stride='i32', slot_mapping_ptr='*i64', token_to_req_indices_ptr='*i32'}; extra_compile_info={configs=[{(0,): [['tt.divisibility', 16]], (1,): [['tt.divisibility', 16]], (2,): [['tt.divisibility', 16]], (3,): [['tt.divi..., device=3, enable_fp_fusion=True, extern_libs=(('libdevice', '/usr/local/lib/python3.12/dist-packages/triton/backends/nvidia/lib/libdevice.10.bc'),), is_warmup=False, launch_cooperative_grid=False, num_ctas=1, num_stages=3, num_warps=4, specialization_data='{"name": "vllm.models.deepseek_v4.sparse_mla._build_c128a_topk_metadata_kernel", "signature": {"global_decode_ptr": ...}; key="[('*i32', 'D'), ('i32', 'D'), ('*i32', 'D'), ('*i32', 'D'), ('i32', 'D'), ('*i64', 'D'), ('i32', 'D'), ('i32', 'D'),...). This causes a latency spike; consider extending warmup to cover this shape/config.
(Worker_TP0 pid=386) WARNING 07-14 11:09:35 [jit_monitor.py:126] Triton kernel JIT compilation during inference: _build_c128a_topk_metadata_kernel (constexprs={BLOCK_SIZE=1024}; signature={BLOCK_SIZE='constexpr', block_size='i32', block_table_ptr='*i32', block_table_stride='i32', compress_ratio='i32', decode_lens_ptr='*i32', global_decode_ptr='*i32', global_decode_stride='i32', max_compressed_tokens='i32', num_decode_tokens='i32', positions_ptr='*i64', prefill_local_ptr='*i32', prefill_local_stride='i32', slot_mapping_ptr='*i64', token_to_req_indices_ptr='*i32'}; extra_compile_info={configs=[{(0,): [['tt.divisibility', 16]], (1,): [['tt.divisibility', 16]], (2,): [['tt.divisibility', 16]], (3,): [['tt.divi..., device=0, enable_fp_fusion=True, extern_libs=(('libdevice', '/usr/local/lib/python3.12/dist-packages/triton/backends/nvidia/lib/libdevice.10.bc'),), is_warmup=False, launch_cooperative_grid=False, num_ctas=1, num_stages=3, num_warps=4, specialization_data='{"name": "vllm.models.deepseek_v4.sparse_mla._build_c128a_topk_metadata_kernel", "signature": {"global_decode_ptr": ...}; key="[('*i32', 'D'), ('i32', 'D'), ('*i32', 'D'), ('*i32', 'D'), ('i32', 'D'), ('*i64', 'D'), ('i32', 'D'), ('i32', 'D'),...). This causes a latency spike; consider extending warmup to cover this shape/config.
(Worker_TP0 pid=386) WARNING 07-14 11:09:35 [jit_monitor.py:126] Triton kernel JIT compilation during inference: eagle_step_slot_mapping_metadata_kernel (constexprs={PAD_ID=-1, batch_size=1, block_size=64, max_model_len=1048576, n_blocks_per_req=16384}; signature={PAD_ID='constexpr', batch_size='constexpr', block_size='constexpr', block_table_ptr='*i32', block_table_stride='i32', max_model_len='constexpr', n_blocks_per_req='constexpr', out_clamped_positions_ptr='*i64', out_slot_mapping_ptr='*i64', positions_ptr='*i64', seq_lens_ptr='*i32'}; extra_compile_info={configs=[{(0,): [['tt.divisibility', 16]], (1,): [['tt.divisibility', 16]], (2,): [['tt.divisibility', 16]], (3,): [['tt.divi..., device=0, enable_fp_fusion=True, extern_libs=(('libdevice', '/usr/local/lib/python3.12/dist-packages/triton/backends/nvidia/lib/libdevice.10.bc'),), is_warmup=False, launch_cooperative_grid=False, num_ctas=1, num_stages=3, num_warps=4, specialization_data='{"name": "vllm.v1.spec_decode.utils.eagle_step_slot_mapping_metadata_kernel", "signature": {"positions_ptr": "*i64",...}; key="[('*i64', 'D'), ('*i32', 'D'), ('i32', 'D'), ('*i32', 'D'), ('*i64', 'D'), ('*i64', 'D'), ('constexpr', 64), ('const...). This causes a latency spike; consider extending warmup to cover this shape/config.
(Worker_TP3 pid=389) WARNING 07-14 11:09:35 [jit_monitor.py:126] Triton kernel JIT compilation during inference: eagle_step_slot_mapping_metadata_kernel (constexprs={PAD_ID=-1, batch_size=1, block_size=64, max_model_len=1048576, n_blocks_per_req=16384}; signature={PAD_ID='constexpr', batch_size='constexpr', block_size='constexpr', block_table_ptr='*i32', block_table_stride='i32', max_model_len='constexpr', n_blocks_per_req='constexpr', out_clamped_positions_ptr='*i64', out_slot_mapping_ptr='*i64', positions_ptr='*i64', seq_lens_ptr='*i32'}; extra_compile_info={configs=[{(0,): [['tt.divisibility', 16]], (1,): [['tt.divisibility', 16]], (2,): [['tt.divisibility', 16]], (3,): [['tt.divi..., device=3, enable_fp_fusion=True, extern_libs=(('libdevice', '/usr/local/lib/python3.12/dist-packages/triton/backends/nvidia/lib/libdevice.10.bc'),), is_warmup=False, launch_cooperative_grid=False, num_ctas=1, num_stages=3, num_warps=4, specialization_data='{"name": "vllm.v1.spec_decode.utils.eagle_step_slot_mapping_metadata_kernel", "signature": {"positions_ptr": "*i64",...}; key="[('*i64', 'D'), ('*i32', 'D'), ('i32', 'D'), ('*i32', 'D'), ('*i64', 'D'), ('*i64', 'D'), ('constexpr', 64), ('const...). This causes a latency spike; consider extending warmup to cover this shape/config.
(Worker_TP1 pid=387) WARNING 07-14 11:09:35 [jit_monitor.py:126] Triton kernel JIT compilation during inference: eagle_step_slot_mapping_metadata_kernel (constexprs={PAD_ID=-1, batch_size=1, block_size=64, max_model_len=1048576, n_blocks_per_req=16384}; signature={PAD_ID='constexpr', batch_size='constexpr', block_size='constexpr', block_table_ptr='*i32', block_table_stride='i32', max_model_len='constexpr', n_blocks_per_req='constexpr', out_clamped_positions_ptr='*i64', out_slot_mapping_ptr='*i64', positions_ptr='*i64', seq_lens_ptr='*i32'}; extra_compile_info={configs=[{(0,): [['tt.divisibility', 16]], (1,): [['tt.divisibility', 16]], (2,): [['tt.divisibility', 16]], (3,): [['tt.divi..., device=1, enable_fp_fusion=True, extern_libs=(('libdevice', '/usr/local/lib/python3.12/dist-packages/triton/backends/nvidia/lib/libdevice.10.bc'),), is_warmup=False, launch_cooperative_grid=False, num_ctas=1, num_stages=3, num_warps=4, specialization_data='{"name": "vllm.v1.spec_decode.utils.eagle_step_slot_mapping_metadata_kernel", "signature": {"positions_ptr": "*i64",...}; key="[('*i64', 'D'), ('*i32', 'D'), ('i32', 'D'), ('*i32', 'D'), ('*i64', 'D'), ('*i64', 'D'), ('constexpr', 64), ('const...). This causes a latency spike; consider extending warmup to cover this shape/config.
(Worker_TP2 pid=388) WARNING 07-14 11:09:35 [jit_monitor.py:126] Triton kernel JIT compilation during inference: eagle_step_slot_mapping_metadata_kernel (constexprs={PAD_ID=-1, batch_size=1, block_size=64, max_model_len=1048576, n_blocks_per_req=16384}; signature={PAD_ID='constexpr', batch_size='constexpr', block_size='constexpr', block_table_ptr='*i32', block_table_stride='i32', max_model_len='constexpr', n_blocks_per_req='constexpr', out_clamped_positions_ptr='*i64', out_slot_mapping_ptr='*i64', positions_ptr='*i64', seq_lens_ptr='*i32'}; extra_compile_info={configs=[{(0,): [['tt.divisibility', 16]], (1,): [['tt.divisibility', 16]], (2,): [['tt.divisibility', 16]], (3,): [['tt.divi..., device=2, enable_fp_fusion=True, extern_libs=(('libdevice', '/usr/local/lib/python3.12/dist-packages/triton/backends/nvidia/lib/libdevice.10.bc'),), is_warmup=False, launch_cooperative_grid=False, num_ctas=1, num_stages=3, num_warps=4, specialization_data='{"name": "vllm.v1.spec_decode.utils.eagle_step_slot_mapping_metadata_kernel", "signature": {"positions_ptr": "*i64",...}; key="[('*i64', 'D'), ('*i32', 'D'), ('i32', 'D'), ('*i32', 'D'), ('*i64', 'D'), ('*i64', 'D'), ('constexpr', 64), ('const...). This causes a latency spike; consider extending warmup to cover this shape/config.

Before submitting a new issue...

• Make sure you already searched for relevant issues, and asked the chatbot living at the bottom right corner of the documentation page, which can answer lots of frequently asked questions.

[leihuang-sketch]

DeepSeek V4 Flash introduces 12 model-specific kernels across Sparse MLA (Triton), Eagle speculative decoding (Triton), MTP RMSNorm (Triton), and mHC (TileLang). Upstream warmup coverage is incomplete: 4 Eagle/MTP kernels have zero warmup, 4 mHC TileLang kernels are only exercised indirectly with 15 fixed token sizes (and the NVIDIA fused path mhc_fused_post_pre_tilelang is entirely uncovered), and 2 Sparse MLA kernels miss key int==1 and %16!=0 divisibility specializations. Only 2 kernels are fully covered. In total, upstream produces ~94 cache entries for 12 kernels — 6 contribute zero independent entries — leaving first-request latency spikes from on-the-fly JIT compilation and latent crash risk from uninitialized buffers and missing OOB guards in _build_c128a_topk_metadata_kernel.

#	Kernel	Type	Source File	Upstream Warmup Status	Upstream Entries	What's Missing / Not Covered
1	_combine_topk_swa_indices_kernel	Triton JIT	vllm/models/deepseek_v4/common/ops/cache_utils.py:567	Partial (exists but incomplete)	~48	constexpr only 4 hardcoded cases (no dynamic derivation from config); M/N only covers %16==0 divisibility ('D'), misses %16!=0 (''); pointer divisibility only 3 offset variants out of 64 possible bitmask combinations; called via wrapper so cannot control per-pointer alignment
2	_build_c128a_topk_metadata_kernel	Triton JIT	vllm/models/deepseek_v4/sparse_mla.py:341	Partial (exists but incomplete)	6	No int==1 specialization enumeration for 7 int args (global_decode_stride, prefill_local_stride, compress_ratio, max_compressed_tokens, num_decode_tokens, block_table_stride, block_size); pointer divisibility not enumerated at all (assumes all 7 pointers aligned); no OOB guard _is_safe(); buffer initialized with torch.empty (garbage values → illegal address); no per-launch synchronize() (async errors accumulate → context corruption)
3	_compute_prefill_metadata_kernel	Triton JIT	vllm/v1/attention/backends/mla/sparse_swa.py:632	Complete (upstream = patch, unchanged)	24	Pointer divisibility for 3 pointers not enumerated (gap exists in both upstream and patch)
4	_build_prefill_chunk_metadata_kernel	Triton JIT	vllm/v1/attention/backends/mla/indexer.py:868	Complete (upstream = patch, unchanged)	~16	Only 2 pointer-divisibility variants (offset / no-offset); could be expanded
5	eagle_prepare_next_token_padded_kernel	Triton JIT	vllm/v1/spec_decode/utils.py	None (no warmup file upstream)	0	Entire kernel not warmed up; no BLOCK_SIZE_TOKENS, num_sampled_tokens_per_req, num_reqs, stride_sampled_token_ids specialization; pointer divisibility not enumerated
6	eagle_prepare_inputs_padded_kernel	Triton JIT	vllm/v1/spec_decode/utils.py	None (no warmup file upstream)	0	Entire kernel not warmed up; num_reqs int==1 specialization missing; pointer divisibility not enumerated
7	eagle_step_slot_mapping_metadata_kernel	Triton JIT	vllm/v1/spec_decode/utils.py:28	None (no warmup file upstream)	0	Entire kernel not warmed up; block_size, max_model_len, n_blocks_per_req, PAD_ID constexpr and batch_size, block_table_stride int==1 specialization all missing; pointer divisibility not enumerated
8	_mtp_shared_head_rmsnorm_kernel	Triton JIT	vllm/models/deepseek_v4/common/ops/fused_mtp_input_rmsnorm.py	None (no warmup file upstream)	0	Entire kernel not warmed up; HIDDEN constexpr (=4096) never enumerated
9	mhc_pre_tilelang	TileLang	torch.ops.vllm.mhc_pre_tilelang	Indirect only (incomplete)	0 (indirect)	Called indirectly via layer.hc_pre() / layer.mhc_pre() with only 15 fixed token sizes (_DEFAULT_TOKEN_SIZE_CANDIDATES); no independent cache entries; RMSNorm weight not passed (attn_norm.weight + variance_epsilon); full token-size range 1..max_tokens not covered
10	mhc_post_tilelang	TileLang	torch.ops.vllm.mhc_post_tilelang	Indirect only (incomplete)	0 (indirect)	Same as above: only 15 fixed token sizes via indirect layer.hc_post() / layer.mhc_post(); no independent cache; full range not covered
11	mhc_fused_post_pre_tilelang	TileLang	torch.ops.vllm.mhc_fused_post_pre_tilelang	None (path not covered)	0	NVIDIA-only fused path (RMSNorm fused into TileLang kernel) entirely uncovered; upstream warmup has no fused-path call site at all
12	hc_head_fused_kernel_tilelang	TileLang	torch.ops.vllm.hc_head_fused_kernel_tilelang	Indirect only (partial sizes)	0 (indirect)	Upstream calls hc_head_op indirectly with only partial token sizes; full range 1..max_tokens not covered

[leihuang-sketch]

We've been seeing DeepSeek V4 flash startup latency spikes — several JIT-compiled kernels trigger compilation on the first served request. The gaps span multiple subsystems (attention, spec decode, compiler, warmup dispatcher), so I split the fix into 5 independent PRs, all forked from main, with no cross-dependencies. Each can be reviewed and merged on its own:
PR Scope Diff
#47807 (#47807) DeepSeek V4 mHC TileLang kernel warmup 1 file, +177/-51
#48804 (#48804) Eagle spec-decode Triton kernel warmup 2 files, +388
#48805 (#48805) Sparse-MLA Triton metadata kernel warmup 1 file, +529/-71
#48806 (#48806) Fix TRITON_CACHE_DIR being silently overwritten 1 line
#48807 (#48807) Wrap each kernel_warmup call in try/except 1 file, +66/-27
Why separate? Each PR touches a different subsystem (attention backends / spec decode / compiler / dispatcher resilience) and likely a different reviewer. Every warmup is wrapped in its own try/except inside the kernel_warmup() dispatcher, so they are independent at runtime — merging one doesn't require the others.
#47807 streamlines the mHC warmup: removes the hard 16K token-size cap, restores sparse power-of-2 warmup sizes to avoid tens of thousands of kernel launches, and fixes missing split-k buckets.
#48804 supersedes the closed #48393: the previous version embedded warmup logic inside the proposer class and stacked on #41481. This version uses a standalone module called from the central warmup dispatcher, with no dependency on #41481. Covers 4 Eagle Triton kernels including the slot-mapping metadata kernel.
#48805 is the largest — it exhaustively enumerates the full Triton specialization space for _combine_topk_swa_indices_kernel (1728 entries) and _build_c128a_topk_metadata_kernel (24576 entries), including pointer-alignment divisibility, with a _is_safe() pre-check that skips combos that would OOB the warmup buffers. Replaces the hardcoded warmup case tuple with dynamic mode derivation from the model's kv_cache_spec.
#48806 is an enabler: InductorAdaptor.initialize_cache() unconditionally overwrites TRITON_CACHE_DIR, discarding any pre-warmed cache (e.g. baked into a Docker image). Switched to setdefault so the user's value wins. One line, but it's what makes a pre-populated cache actually usable.
#48807 wraps every warmup call in kernel_warmup() in its own try/except. Previously a single exception (missing import, shape mismatch, OOM) would abort the entire warmup sequence, leaving all subsequent kernels un-warmed. Now each failure is logged as a warning and the remaining warmups still run.
Happy to merge in any order. #48806 and #48807 are standalone; #47807 / #48804 / #48805 all touch kernel_warmup.py but only in additive hunks (different sections each), so there are no merge conflicts between them.