执行摘要

• 一句话：DSv4 MLA 元数据内核迁移至共享 JIT warmup 契约
• 推荐动作：值得精读。核心看三点：VllmTritonJitKernel + kernel_launcher 的「预热即假启动」设计（运行时 launch 规格是唯一事实来源，避免 launch 与 warmup 参数漂移）；triton_scalar_specialization_rep 对 Triton 标量特化类（exact-one / divisible-by-16 / generic）的处理；zip_inputs + _trace_dispatch 对编译键空间的声明式枚举。评审对话本身也是很好的教材：三处「预热变体不可达」说明编译键必须覆盖运行时真实特化类。若团队要迁移其他 Triton 内核，建议先读 jit_warmup_triton_helper.py 与本 PR 的 review 记录。

功能与动机

目标来自 Issue #49349「Zero JIT compilation during runtime」：把内核编译全部前移到启动期 warmup，使模型能在 --jit-monitor-mode error 下无任何运行时 JIT。本 PR 是 DSv4 模型 de-JIT 化的 1/N 拆分之一，负责迁移「generic MLA metadata and indexing kernels」。旧实现（vllm/model_executor/warmup/sparse_mla_triton_warmup.py）在 worker 上按 attention backend 名称探测后手工调用各内核 warmup，代表性形状写死在外部函数里，内核自身不掌握自己的编译键空间。PR body 明确列出新方案要点：迁移 sparse-MLA 元数据与索引转换内核、迁移 prefill chunk 元数据与 top-k/SWA 组合内核、用 kernel-owned traced dispatch spaces 取代外部代表性形状循环、只注册激活后端选中的 warmup provider。

实现拆解

整个迁移分 5 步推进：

1. 扩展共享预热基建。vllm/model_executor/warmup/jit_warmup_triton_helper.py 新增 VllmTritonJitKernel 基类（约 +90 行）：抽象方法 warmup_inputs(compile_key) 返回能复现该编译键的输入；compile() 置 _warming 标志后走真实 call，launch() 依据标志路由到 kernel.warmup(grid=(1,)) 做纯编译期代码生成。配套新增 kernel_launcher 装饰器与 LaunchSpec 类型，让 call 只声明 grid 与 launch kwargs；TritonWarmupTensor 增加 strides 字段以模拟非紧凑/转置张量布局。vllm/model_executor/warmup/jit_warmup.py 的 registry.warmup() 改为用 inspect.signature 探测 get_warmup_keys 是否需要参数、按需注入 vllm_config，从而支持 config-free 内核；kernel_warmup.py 移除对已删除文件的 import。

2. 迁移四个 MLA 内核模块为自描述内核。sparse_utils.py（+370/-173）：ConvertReqIndexToGlobalIndexKernel 的 max_num_blocks_per_req 由 tl.constexpr 降为 do_not_specialize 运行时参数，避免按 block-table 宽度无限重编译；get_warmup_keys 用 zip_inputs 枚举 5 个变体（prefill workspace、count_valid、DCP size/rank/interleave、compact_to_front）；warmup_inputs 为 token_indices/out 显式提供 (16, 1) 步长以匹配分页缓存布局；valid_counts 分配恢复 single-tile 用 torch.empty、multi-tile 用 torch.zeros 的分支。compressor_utils.py（+125/-40）：CompressedSlotMappingKernel 的 block_table_stride 进 do_not_specialize，block_size 用 triton_scalar_specialization_rep 代表值，压缩比列表合并 hf_config.index_kpool，无大于 1 的压缩比时返回空 key 列表跳过预热。indexer.py（+163/-101）：PrepareUniformDecodeKernel 增加显式 num_decode_tokens 驱动 grid（修复按 buffer 容量 over-launch 的越界读），BuildPrefillChunkMetadataKernel 编译键只保留算法维度（block size、压缩比、指针对齐），query_slice 与 DCP 参数走特化类代表值。sparse_swa.py（+304/-155）：ComputePrefillMetadataKernel 与 ComputeSWAIndicesAndLensKernel 迁移，window_size/block_size 编译键从 vllm_config 取值，调用点改用命名参数 num_tokens=/token_offset=，并新增 _COMPUTE_DSPARK_NONCAUSAL_SWA_INDICES_KERNEL 包装。

3. 注册点迁移到模型构造处并删除外部编排。删除 vllm/model_executor/warmup/sparse_mla_triton_warmup.py（-111 行）。vllm/models/deepseek_v4/attention.py 的 DeepseekV4Attention.init 与 vllm/model_executor/layers/attention/mla_attention.py 的 MLAAttention.init/bind_kv_cache 在 enable_jit_warmup 下按激活后端注册对应 provider；block_stride_rows 依赖真实 KV-cache stride，故放到 bind_kv_cache 中注册。review 中补回了 _COMBINE_TOPK_SWA_INDICES_KERNEL 的注册。

4. runner 集成。vllm/v1/worker/gpu/model_runner.py 与 vllm/v1/worker/gpu_model_runner.py 把 load_model 与 initialize_kv_cache 包进 jit_warmup_registry.activate() 上下文，让模型构造与 KV-cache 分配期间产生的注册被 registry 收集。

5. 测试与文档配套。tests/v1/attention/test_indexer_deepseek_v4_slot_mapping.py 新增 test_compressed_slot_mapping_warmup_includes_index_kpool 与 test_index_conversion_warmup_uses_physical_block_stride；tests/model_executor/test_jit_warmup.py 新增 test_registry_default_request_supports_config_free_kernel；test_gpu_model_runner_v2_eplb.py、test_fused_shared_expert.py 做 enable_jit_warmup 相关适配；docs/contributing/jit_kernel_warmup.md 更新注册示例（内部单例统一下划线命名）。

关键文件：

• vllm/v1/attention/backends/mla/sparse_swa.py（模块 SWA 内核；类别 source；类型 core-logic；符号 compile, warmup_inputs, _compute_swa_indices_and_lens_kernel, ComputeSWAIndicesAndLensKernel）: 核心迁移文件：ComputePrefillMetadataKernel 与 ComputeSWAIndicesAndLensKernel 迁入 VllmTritonJitKernel 契约，新增 DSpark 非因果包装；update_draft_decode_metadata 从 shape[-1] 改 window_size 的行为变化存疑，风险最集中。
• vllm/v1/attention/backends/mla/sparse_utils.py（模块 索引内核；类别 source；类型 core-logic；符号 _convert_req_index_to_global_index_kernel, ConvertReqIndexToGlobalIndexKernel, CompileKey, kernel）: 最大改动文件（+370/-173）：ConvertReqIndexToGlobalIndexKernel 迁移，max_num_blocks_per_req 降为 do_not_specialize 运行时参数，get_warmup_keys 以 zip_inputs 枚举 5 个 DCP/prefill 变体，warmup_inputs 显式模拟分页步长。
• vllm/model_executor/warmup/jit_warmup_triton_helper.py（模块 预热基建；类别 source；类型 data-contract；符号 stride, VllmTritonJitKernel, warmup_inputs, compile）: 整个迁移的契约核心：新增 VllmTritonJitKernel 基类、kernel_launcher 装饰器、LaunchSpec 类型，TritonWarmupTensor 增加 strides 支持；后续所有 Triton 内核迁移都依赖此文件。
• vllm/v1/attention/backends/mla/compressor_utils.py（模块 压缩映射；类别 source；类型 core-logic；符号 _compressed_slot_mapping_kernel, CompressedSlotMappingKernel, CompileKey, kernel）: CompressedSlotMappingKernel 迁移的代表性完整示例（+125/-40）：压缩比并入 index_kpool，block_table_stride 进 do_not_specialize，block_size 用特化类代表值；也是 review 中「预热变体不可达」问题的修复样板。
• vllm/v1/attention/backends/mla/indexer.py（模块 索引器；类别 source；类型 core-logic；符号 _prepare_uniform_decode_kernel, PrepareUniformDecodeKernel, CompileKey, kernel）: PrepareUniformDecodeKernel 与 BuildPrefillChunkMetadataKernel 迁移；修复 grid 按 buffer 容量 over-launch 的越界读问题；malaiwah 指出的首轮 prefill 残留编译缺口也集中在此文件。
• vllm/model_executor/warmup/sparse_mla_triton_warmup.py（模块 预热编排；类别 source；类型 deletion；符号 _attention_backend_name, _has_attention_backend, _compile_sparse_swa_prefill_metadata_kernel, _compile_prefill_chunk_metadata_kernel）: 整个旧式外部预热编排被删除（-111 行）：按 attention backend 名称探测并手工调用各内核 warmup 的逻辑，被模型构造处的就地注册取代，是本次迁移的方向性标志。
• vllm/models/deepseek_v4/attention.py（模块 DSv4 模型；类别 source；类型 data-contract）: DeepseekV4Attention.init 按激活后端就地注册 7 个 warmup provider；review 中发现 _COMBINE_TOPK_SWA_INDICES_KERNEL 覆盖丢失后在此补回。
• vllm/model_executor/layers/attention/mla_attention.py（模块 MLA 层；类别 source；类型 data-contract）: 通用稀疏 MLA 后端（FLASHMLA_SPARSE、FLASHINFER_MLA_SPARSE、DEEPSEEK_V32_INDEXER 等）在 MLAAttention 构造与 bind_kv_cache 处注册 provider；block_stride_rows 依赖真实 KV-cache stride，必须在缓存绑定后注册。
• vllm/model_executor/warmup/jit_warmup.py（模块 预热注册表；类别 source；类型 data-contract）: registry.warmup() 改为 signature sniffing 按需注入 vllm_config，支持 config-free 内核；Lucas 曾建议统一契约，最终同意维持此设计。
• vllm/v1/worker/gpu/model_runner.py（模块 运行器；类别 source；类型 data-contract；符号 load_model, initialize_kv_cache）: load_model 与 initialize_kv_cache 包进 jit_warmup_registry.activate() 上下文，让依赖 KV-cache stride 的 provider 在分配后注册。
• vllm/v1/worker/gpu_model_runner.py（模块 运行器；类别 source；类型 data-contract）: 旧版 runner 同步包裹 registry.activate()，保持两条 runner 路径行为一致。
• vllm/model_executor/warmup/kernel_warmup.py（模块 预热基建；类别 source；类型 data-contract）: 移除对已删除 sparse_mla_triton_warmup 模块的 import，保持 warmup 入口干净。
• tests/v1/attention/test_indexer_deepseek_v4_slot_mapping.py（模块 索引测试；类别 test；类型 test-coverage；符号 test_compressed_slot_mapping_warmup_includes_index_kpool, test_index_conversion_warmup_uses_physical_block_stride）: 新增 test_compressed_slot_mapping_warmup_includes_index_kpool 与 test_index_conversion_warmup_uses_physical_block_stride，直接覆盖 review 中暴露的两类编译键问题。
• tests/model_executor/test_jit_warmup.py（模块 预热测试；类别 test；类型 test-coverage；符号 test_registry_default_request_supports_config_free_kernel, ConfigFreeKernel, CompileKey, init）: 新增 test_registry_default_request_supports_config_free_kernel，验证 signature sniffing 对无配置内核的注册路径。
• tests/v1/worker/test_gpu_model_runner_v2_eplb.py（模块 运行器测试；类别 test；类型 test-coverage）: 适配 enable_jit_warmup 配置，保证 EP LB 路径在新注册机制下测试通过。
• tests/model_executor/layers/test_fused_shared_expert.py（模块 融合专家测试；类别 test；类型 test-coverage）: 小幅适配 enable_jit_warmup，保持融合共享专家层测试与新预热注册兼容。
• docs/contributing/jit_kernel_warmup.md（模块 文档；类别 docs；类型 documentation）: 更新注册示例命名约定：内部单例统一用下划线前缀，避免被误当作外部 API。

关键符号：VllmTritonJitKernel.compile, VllmTritonJitKernel.launch, kernel_launcher, ConvertReqIndexToGlobalIndexKernel.dispatch, ConvertReqIndexToGlobalIndexKernel.get_warmup_keys, ConvertReqIndexToGlobalIndexKernel.warmup_inputs, CompressedSlotMappingKernel.dispatch, CompressedSlotMappingKernel.get_warmup_keys, CompressedSlotMappingKernel.warmup_inputs, PrepareUniformDecodeKernel.get_warmup_keys, BuildPrefillChunkMetadataKernel.get_warmup_keys, ComputePrefillMetadataKernel.warmup_inputs, ComputeSWAIndicesAndLensKernel, register_warmup, JitWarmupRegistry.warmup

关键源码片段

vllm/v1/attention/backends/mla/sparse_utils.py

最大改动文件（+370/-173）：ConvertReqIndexToGlobalIndexKernel 迁移，max_num_blocks_per_req 降为 do_not_specialize 运行时参数，get_warmup_keys 以 zip_inputs 枚举 5 个 DCP/prefill 变体，warmup_inputs 显式模拟分页步长。

    def get_warmup_keys(
        self,
        vllm_config: VllmConfig,
        *,
        block_stride_rows: int,
    ) -> list[CompileKey]:
        # 编译键空间由运行时配置推导，取代旧实现里写死的外部“代表性形状”循环：
        # block size、DCP 参数、index_topk、block-table 宽度全部来自 vllm_config。
        block_size = vllm_config.cache_config.block_size
        dcp_size = vllm_config.parallel_config.decode_context_parallel_size
        dcp_interleave = vllm_config.parallel_config.cp_kv_cache_interleave_size
        dcp_rank = get_dcp_group().rank_in_group if dcp_size > 1 else 0
        num_topk_tokens = vllm_config.model_config.hf_config.index_topk
        max_num_blocks = cdiv(
            vllm_config.model_config.max_model_len,
            block_size * dcp_size,
        )
        max_num_blocks_per_req = get_block_table_width(max_num_blocks, block_size)
        # zip_inputs 对 5 个“语义变体”做笛卡尔积式的组合枚举：
        # prefill workspace / count_valid / DCP 过滤与紧凑化是正交的算法分支。
        return self._trace_dispatch(self.dispatch)(
            zip_inputs(
                dict(HAS_PREFILL_WORKSPACE=False, COUNT_VALID=False,
                     COMPACT_TO_FRONT=False, DCP_SIZE=1, DCP_RANK=0, DCP_INTERLEAVE=1),
                dict(HAS_PREFILL_WORKSPACE=False, COUNT_VALID=True,
                     COMPACT_TO_FRONT=False, DCP_SIZE=1, DCP_RANK=0, DCP_INTERLEAVE=1),
                dict(HAS_PREFILL_WORKSPACE=True, COUNT_VALID=True,
                     COMPACT_TO_FRONT=False, DCP_SIZE=1, DCP_RANK=0, DCP_INTERLEAVE=1),
                dict(HAS_PREFILL_WORKSPACE=False, COUNT_VALID=True,
                     COMPACT_TO_FRONT=False, DCP_SIZE=dcp_size,
                     DCP_RANK=dcp_rank, DCP_INTERLEAVE=dcp_interleave),
                dict(HAS_PREFILL_WORKSPACE=False, COUNT_VALID=True,
                     COMPACT_TO_FRONT=dcp_size > 1, DCP_SIZE=dcp_size,
                     DCP_RANK=dcp_rank, DCP_INTERLEAVE=dcp_interleave),
            ),
            BLOCK_SIZE=block_size,
            BLOCK_STRIDE_ROWS=block_stride_rows,
            BLOCK_N=128,
            NUM_TOPK_TOKENS=num_topk_tokens,
            max_num_blocks_per_req=max_num_blocks_per_req,
        )
​
    def warmup_inputs(self, compile_key: CompileKey) -> dict[str, Any]:
        int32_ptr = TritonWarmupTensor(torch.int32)
        token_shape = (1, compile_key.num_topk_tokens)
        # token_indices / out 显式给出 (16, 1) 步长：运行时张量是分页布局，
        # 若按紧凑布局推断，Triton 指针特化键会与真实启动不一致，预热即失效。
        token_strides = (16, 1)
        return dict(
            req_id=int32_ptr,
            block_table=TritonWarmupTensor(
                torch.int32,
                shape=(1, compile_key.block_table_stride),
            ),
            token_indices=TritonWarmupTensor(
                torch.int32, shape=token_shape, strides=token_strides,
            ),
            out=TritonWarmupTensor(
                torch.int32, shape=token_shape, strides=token_strides,
            ),
            valid_counts=int32_ptr if compile_key.count_valid else None,
            prefill_workspace_request_ids=(
                int32_ptr if compile_key.has_prefill_workspace else None
            ),
            prefill_workspace_starts=(
                int32_ptr if compile_key.has_prefill_workspace else None
            ),
            # max_num_blocks_per_req 已是 do_not_specialize 运行时参数，
            # 值本身不参与编译键，这里给 1 即可。
            max_num_blocks_per_req=1,
            BLOCK_SIZE=compile_key.block_size,
            BLOCK_STRIDE_ROWS=compile_key.block_stride_rows,
        )

vllm/model_executor/warmup/jit_warmup_triton_helper.py

整个迁移的契约核心：新增 VllmTritonJitKernel 基类、kernel_launcher 装饰器、LaunchSpec 类型，TritonWarmupTensor 增加 strides 支持；后续所有 Triton 内核迁移都依赖此文件。

class VllmTritonJitKernel(VllmJitKernel[CompileKeyT], Generic[CompileKeyT]):
    """Triton 内核的共享 JIT 预热契约：用“假启动”代替手工 warmup 调用。"""
​
    kernel: ClassVar[Any]
    _warming = False
​
    @abstractmethod
    def warmup_inputs(self, compile_key: CompileKeyT) -> dict[str, Any]:
        # 子类负责为每个 CompileKey 给出能复现该编译键的输入：
        # 张量用 TritonWarmupTensor 描述符代替，标量用特化类代表值。
        raise NotImplementedError
​
    def compile(self, compile_key: CompileKeyT) -> None:
        # 预热 = 带着 fake 输入走一遍真实的 __call__。_warming 标志让
        # launch() 路由到 kernel.warmup(grid=(1,))：只做编译期类型推导
        # 与代码生成，不真正启动 grid，因此不需要真实显存与真实 batch。
        inputs = self.warmup_inputs(compile_key)
        self._warming = True
        try:
            cast(Callable[..., None], self)(**inputs)
        finally:
            self._warming = False
​
    @cached_property
    def _kernel_param_names(self) -> frozenset[str]:
        # 调用侧惯用不带 _ptr 后缀的名字（如 seq_lens），launch() 据此
        # 把实参自动映射到 Triton 形参（seq_lens_ptr）。
        return frozenset(self.kernel.arg_names)
​
    def launch(
        self,
        grid: tuple[int, ...],
        inputs: Mapping[str, Any],
        /,
        **kwargs: Any,
    ) -> Any:
        for name, value in inputs.items():
            target = name if name in self._kernel_param_names else f"{name}_ptr"
            if target in self._kernel_param_names and target not in kwargs:
                kwargs[target] = value
        if self._warming:
            warmup = getattr(self.kernel, "warmup", None)
            assert warmup is not None
            return warmup(grid=(1,), **kwargs)
        return self.kernel[grid](**kwargs)
​
​
def kernel_launcher(
    call_fn: Callable[..., LaunchSpec],
) -> Callable[..., None]:
    """把声明式的 __call__（返回 grid + launch kwargs）包装成统一启动入口。"""
    signature = inspect.signature(call_fn)
​
    @wraps(call_fn)
    def wrapper(
        self: VllmTritonJitKernel[Any],
        *args: Any,
        **kwargs: Any,
    ) -> None:
        grid, launch_kwargs = call_fn(self, *args, **kwargs)
        bound = signature.bind(self, *args, **kwargs)
        bound.apply_defaults()
        inputs = {
            name: value for name, value in bound.arguments.items() if name != "self"
        }
        self.launch(grid, inputs, **launch_kwargs)
​
    return wrapper

评论区精华

评审由维护者 LucasWilkinson 主导，多线程为 blocking 级，核心交锋如下：

• LucasWilkinson 在 indexer.py 指出 grid 尺寸改变导致越界读：“The call site passes the full decode_seq_lens_buffer (allocated at max_num_batched_tokens), so this grid launches the buffer capacity, not the batch... Padding programs compute req_id = idx // max_decode_len past the real batch and read seq_lens/block_table rows out of bounds.” 作者以显式 num_decode_tokens 参数修正（“Runtime count is explicit now and drives the grid.”）。

• 最大的一类问题被归纳为「预热变体不可达」：CompressedSlotMappingKernel 传字面量 1，命中的是 Triton 的 exact-one 特化类，而运行时值落在 divisible-by-16 类，等于白编译。同样的问题出现在 ConvertReqIndexToGlobalIndexKernel 的 constexpr max_num_blocks_per_req：“It's only a bounds check — suggest demoting it to a do_not_specialize runtime arg.” 两处均按建议修复。

• _COMBINE_TOPK_SWA_INDICES_KERNEL 覆盖丢失：“The deleted sparse_mla_triton_warmup.py also warmed _COMBINE_TOPK_SWA_INDICES_KERNEL; nothing registers it after this PR, so it JITs at first use.” Lucas 建议统一 register_dsv4_sparse_warmups() 辅助函数防两处注册漂移；作者补回覆盖但明确不采纳统一 helper：“The migrated kernels are not exclusively DSv4-owned, and different backends require different subsets.”

• update_draft_decode_metadata 的语义变化：用 self.window_size 取代 decode_swa_indices.shape[-1]，Lucas 提醒 DSpark 非因果缓冲宽度是对齐后的 window+spec，行尾可能残留 stale indices；此线程在材料中未见作者明确回复，属遗留确认项。

• malaiwah 在 Issue 评论中给出 fresh-cache 校验结论：BuildPrefillChunkMetadataKernel 仍会逐 worker 编译一次，需把 query_slice_start/stop 与 DCP 参数标记 do_not_specialize，并用一次真实一行执行填充进程内 JITFunction.run cache；修复在独立 review branch，未并入本 PR。

• 设计取舍：Lucas 建议 get_warmup_keys(self, vllm_config) 统一契约以省去 signature sniffing；作者反驳“Some kernels need no configuration... adds unused-argument boilerplate”；Lucas 最终认同“lets ignore this”。另外 lucas 建议把无调用方的 jit_warmup_cutedsl_helper.py 移出本 PR，最终文件清单已无该文件。

• PrepareUniformDecodeKernel 按 buffer 容量 over-launch 导致越界读 (correctness): 作者改为显式 num_decode_tokens 关键字驱动 grid（“Runtime count is explicit now and drives the grid.”），head 版本已修正。

• CompressedSlotMappingKernel 预热变体不可达（Triton 特化类不匹配） (correctness): 改为 do_not_specialize=["block_table_stride"]，block_size 用 triton_scalar_specialization_rep 代表值；indexer.py 同样问题一并修复。
• ConvertReqIndexToGlobalIndexKernel 的 max_num_blocks_per_req 作 constexpr 导致死预热 (correctness): 降为 do_not_specialize 运行时参数，作者确认“It is now a runtime, non-specialized argument.”。
• BuildPrefillChunkMetadataKernel 在 fresh-cache 下仍逐 worker 编译一次 (correctness): 修复方案（do_not_specialize + 一次真实一行执行 + OOT 后端 config-name 回退 + GLM index_kpool）在独立 review branch，未并入本 PR，需后续跟进。
• _COMBINE_TOPK_SWA_INDICES_KERNEL 预热覆盖丢失 (correctness): 作者补回注册，但未采纳统一 helper：“The migrated kernels are not exclusively DSv4-owned, and different backends require different subsets.”。
• update_draft_decode_metadata 的 window_size 替代 shape[-1] 行为变化 (correctness): 材料中未见作者明确回复，head 版本仍用 self.window_size，需回归确认 DSpark 非因果 draft 路径。
• valid_counts 无条件 torch.zeros 丢掉 single-tile 优化 (performance): 作者恢复分支：single-tile 用 torch.empty，multi-tile 用 torch.zeros。
• 硬编码 block size 64 与 NUM_TOPK_TOKENS=2048 的编译键一致性 (correctness): 作者确认编译键取值改从 vllm_config 获取（“Values now come from vllm_config.”）。
• get_warmup_keys 是否统一注入 vllm_config 的契约争论 (design): Lucas 最终认同：“agreed @LopezCastroRoberto lets ignore this”。
• jit_warmup_cutedsl_helper.py 无调用方，建议移出本 PR (design): 最终合并文件清单已无该文件，按建议移出。

风险与影响

• 风险：
  1. 残留首轮运行时编译：malaiwah 的 fresh-cache 校验发现 BuildPrefillChunkMetadataKernel 在部分配置下仍会逐 worker 编译一次（indexer.py）。query_slice_start/stop、DCP_RANK/WORLD/INTERLEAVE 若留在 CompileKey 会生成 callsite 变体，且 compile-only 的 kernel.warmup() 不填充进程内 JITFunction.run cache；OOT 后端的 config-name 回退与 GLM index_kpool 也需纳入。修复在独立 review branch，未并入本 PR，在 --jit-monitor-mode error 下会直接暴露。
  2. DSpark 非因果 draft 路径行为变化：sparse_swa.py 的 update_draft_decode_metadata 用 window_size 取代 decode_swa_indices.shape[-1]，非因果缓冲宽度为 64 对齐后的 window+spec；若 DSpark 可到达该路径，有效窗口收窄且清行循环只覆盖 window_size，行尾会残留 stale indices。
  3. 注册点分散易漂移：同一内核族在 deepseek_v4/attention.py 与 mla_attention.py 多处注册，作者未采纳统一 helper，_COMBINE_TOPK_SWA_INDICES_KERNEL 已被遗漏过一次，长期靠 review 兜底。
  4. 编译键与运行时特化类强耦合：triton_scalar_specialization_rep 的 1、16 整除、其他三类必须与运行时真实值一致，任何新调用点若改变 shape/stride/标量分布，预热即静默失效；本次 review 抓到的三类问题均属此。
  5. 性能回归风险：valid_counts 分配与 grid 尺寸若回归到统一 torch.zeros 或 buffer-capacity 启动，会重新引入 decode 热路径 memset 或越界读。
     - 影响：运行时影响：对启用 enable_jit_warmup 的 DSv4 与通用稀疏 MLA 后端（FLASHMLA_SPARSE、FLASHINFER_MLA_SPARSE、FLASHINFER_MLA_SPARSE_SM120、DEEPSEEK_V32_INDEXER 等），内核编译前移到启动期，覆盖完整时首次推理不再出现 JIT 停顿，直接服务于 Issue #49349 的「Zero JIT」目标。架构影响：确立「内核自描述 + 注册就近」的迁移范式，VllmTritonJitKernel 与 kernel_launcher 成为后续 Triton 内核迁移的模板，删除外部编排文件减少了 backend 名称字符串与内核符号的隐式耦合。团队影响：为系列后续 PR（#50176-#50178、#53564-#53567 等）提供可复制模式，review 暴露的特化类问题将成为后续迁移的检查清单。范围限制：仅覆盖 Triton 元数据/索引内核，DeepGEMM、CuTeDSL、TileLang 等其他后端的 de-JIT 仍在后续工作中。
     - 风险标记：残留首轮运行时编译, DSpark 非因果 draft 路径行为改变, 注册点分散易遗漏内核, 预热键与运行时特化类强耦合, 依赖后续 PR 系列闭环

关联脉络

• PR #49315 共享 JIT warmup 基础设施: PR body 声明 Depends on #49315，本 PR 所有内核迁移都建立在该共享 JIT warmup 注册表基础设施之上；标题为按上下文概括，材料中未直接提供。
• PR #49627 DSv4 模型级 de-JIT 迁移（父 PR）: PR body 声明为 parent PR，本 PR 是其在 MLA 元数据与索引内核维度上的 1/N 拆分解；标题为按上下文概括。
• PR #53576 新增 FlashInfer PCIe IPC all-reduce 后端: 历史 PR 中与本 PR 同改 vllm/model_executor/warmup/kernel_warmup.py，说明 warmup 基建正被分布式通信等多方复用，后续改动需避免相互冲突。
• PR #46167 统一运行时监控（jit-monitor-mode）: Issue #49349 指出 --jit-monitor-mode error 由 #46167 引入，是验证本 PR 是否完全消除运行时 JIT 的判定手段；标题为按上下文概括。
• PR #50176 DSv4 de-JIT 后续迁移步骤: Issue #49349 将其列为紧随本 PR 的 DSv4 迁移任务（与 #53564、#53565、#50177、#50178 等同系列），malaiwah 发现的残留缺口可能由系列后续 PR 收尾；标题为按上下文概括。