执行摘要

• 一句话：V4.1 索引器接入稀疏 MQA logits，长上下文提速
• 推荐动作：值得精读，尤其适合负责长上下文推理性能、KV cache 布局和内核接入的同学。重点关注四处设计决策：一是「一个独立的 SparseMQAIndexer 模块 + 独立 backend，而不是在 dense indexer 里加分支」，候选源层与消费者层的语义被清晰切开；二是「所有校验前移到引擎启动、路径上不留任何 PyTorch fallback」，这是评审反复推动的结果，也是这类新内核接入的可复用范式；三是「weights dtype 由融合 Q 核直接产出 bf16」，把 dtype 转换从热路径消灭掉；四是 KV 块的 block_stride_alignment 对齐方案，用 0.8% 显存换取 DeepGEMM 分页核的 512B 步长约束。阅读时建议配合 tests/model_executor/layers/test_mla_short_prefill_indexer.py 的参考实现一起看，并确认 _INT32_BIG 的模块级 tl.constexpr 是否已在合入版本修正。

功能与动机

DeepSeek V4.1 的注意力选 token 分两级：一个候选源索引器发布上下文的 top 候选块，其后的消费者索引器在这些块内挑自己的 top-k。PR body 指出「Today the consumer layers still score the whole context with dense MQA logits and then mask everything outside the candidates」——即消费者层仍按整段上下文算 dense logits 再 mask。DeepGEMM 2.8 随 deepseek-ai/DeepGEMM#432（Public Release 26/09 中列出的 “Sparse Indexer: hierarchical sparse indexing for V4.1”）提供了可按块打分的稀疏 MQA-logits 内核，于是本 PR 让消费者层只对候选块打分，top-k 在压缩空间运行并映射回上下文位置，把长上下文的索引器开销变成常数级。

实现拆解

1. 开关与选路：vllm/config/attention.py 新增 indexer_sparse_logits；vllm/models/deepseek_v41/attention.py 中 DeepseekV4IndexerCache.get_attn_backend() 据此返回 DeepseekV41SparseIndexerBackend，DeepseekV4Indexer.__init__ 用 use_sparse_logits = indexer_sparse_logits and candidate_block_buffer is not None and not candidate_write 决定实例化 SparseMQAIndexer 还是旧的 SparseAttnIndexer（候选源层及其之前的层保持 dense），并据 SparseMQAIndexer.weights_dtype 设置 self.indexer_weights_dtype。
2. 独立后端与元数据（启动期 fail-closed）：新增 vllm/v1/attention/backends/mla/sparse_indexer.py，包含 SparseMQARowsMetadata（行级 scratch）、DeepseekV41SparseIndexerMetadata（继承 DeepseekV32IndexerMetadata，追加 sparse_block_kv / sparse_decode / sparse_prefill）、DeepseekV41SparseIndexerBackend 与 DeepseekV41SparseIndexerMetadataBuilder。builder 在引擎启动时一次性校验模型两级索引、MXFP4 indexer cache、SM100、DeepGEMM >= 2.8、DeepSelect 扩展、top-k backend 取值、page 可被 sparse_block_kv 整除，并把 sparse_indices_buffer / end_buffer / col_indices_buffer / zero_row_ks 作为自有缓冲；build() 在 varlen decode 与每个 prefill chunk 上切出这些缓冲的视图，同一 step 的所有消费者层共享一次候选展平与 DeepGEMM schedule。
3. 内核算子层：新增 vllm/model_executor/kernels/attention/dsa/sparse_mqa_logits.py。_expand_candidates_kernel 逐行完成候选块展平、区间过滤、升序排序与重复末尾补齐，candidate_blocks_to_sparse_indices 是它的 Python 封装并支持调用方传入输出缓冲；sparse_topk_remap 用 DeepSelect 在 bf16 logits 上做 top-k，再由 _sparse_topk_remap_kernel 把稀疏列映射回 request-local 位置；sparse_mqa_logits_prefill_chunk / sparse_mqa_logits_paged_decode 负责端到端串联（含 DeepGEMM schedule 的复用）。
4. DeepGEMM 桥接与 KV 布局：vllm/utils/deep_gemm.py 增加四个懒加载符号与 has_deep_gemm_sparse_mqa()，metadata 包装函数固定 use_unaligned_ks=True 以避免 host 端同步选择内核变体；vllm/v1/kv_cache_interface.py 给 MLAAttentionSpec 加 block_stride_alignment（并纳入 merge() 的一致性断言），vllm/v1/core/kv_cache_utils.py:_get_kv_cache_bytes_per_block 用 round_up 把块字节数对齐到 lcm(512, page_alignment)，满足 DeepGEMM 分页核 base + page * stride 的 512B 步长要求（开启后 KV 多占 0.8%）。
5. 配套改动：vllm/v1/attention/backends/mla/indexer.py 抽出 _prefill_split_seq_lens 钩子，sparse builder 覆写 min_split_seq_len = num_sparse_cols * 2 // 4，让按 fp32 [rows, seq_len] 预算的 prefill 分块器按 bf16 稀疏行宽切块；fused_indexer_q_rope_quant（Triton 与 CuTeDSL MXFP4 变体）新增 weights_out_dtype，直接写 bf16 per-head weights，热路径不再 cast。测试侧新增 tests/model_executor/layers/test_mla_short_prefill_indexer.py 的候选展平数学/随机对拍/稀疏与 dense-masked top-k 选中值一致性（含 CUDA graph capture 与 schedule 复用）用例，以及 tests/kernels/test_fused_indexer_q_rope_quant.py 的 bf16 权重用例。

关键文件：

• vllm/model_executor/kernels/attention/dsa/sparse_mqa_logits.py（模块 索引器内核；类别 source；类型 data-contract；符号 _expand_candidates_kernel, _sparse_topk_remap_kernel, candidate_blocks_to_sparse_indices, sparse_topk_remap）: 新增的稀疏 MQA logits 内核层：候选块展平/过滤/排序、DeepSelect top-k 调用与稀疏列到 request-local 位置的重映射，以及 prefill/decode 两条端到端串联函数，是全 PR 的算法核心。
• vllm/model_executor/layers/sparse_mqa_indexer.py（模块 索引器层；类别 source；类型 data-contract；符号 SparseMQAIndexer, init, forward, _reserve_workspaces）: 新增的消费者侧索引器模块，替代 SparseAttnIndexer 走稀疏打分；forward 串起 prefill 的 K gather 与 paged decode 两条路径，并定义 weights_dtype = torch.bfloat16 作为与融合 Q 核之间的 dtype 契约。
• vllm/v1/attention/backends/mla/sparse_indexer.py（模块 注意力后端；类别 source；类型 dependency-wiring；符号 SparseMQARowsMetadata, DeepseekV41SparseIndexerMetadata, DeepseekV41SparseIndexerBackend, get_name）: 新增后端与元数据：定义 SparseMQARowsMetadata/DeepseekV41SparseIndexerMetadata，并在 builder 中完成所有前置条件的启动期校验与行级 scratch 的持有，是多层共享一次候选展平与 schedule 的关键。
• vllm/utils/deep_gemm.py（模块 内核桥接；类别 source；类型 core-logic；符号 has_deep_gemm_sparse_mqa, get_sparse_mqa_logits_metadata, get_paged_sparse_mqa_logits_metadata, fp8_fp4_sparse_mqa_logits）: DeepGEMM 懒加载桥接：新增 4 个稀疏符号与 has_deep_gemm_sparse_mqa()，并把 unaligned-ks 变体固定在包装函数里，避免每一步做 host 端同步决策。
• vllm/models/deepseek_v41/attention.py（模块 模型接入；类别 source；类型 data-contract；符号 DeepseekV4IndexerCache.get_kv_cache_spec, DeepseekV4IndexerCache.get_attn_backend, DeepseekV4Indexer.init, indexer_weights_dtype）: 模型侧接入点：DeepseekV4IndexerCache 依据开关选择后端并设置 block_stride_alignment，DeepseekV4Indexer 在消费者层实例化 SparseMQAIndexer 并据其 weights_dtype 决定融合 Q 核的输出精度。
• tests/model_executor/layers/test_mla_short_prefill_indexer.py（模块 索引器测试；类别 test；类型 test-coverage；符号 _make_candidates, _candidate_token_mask, _reference_topk, _assert_same_selection）: 测试主体：候选展平的定点用例、与 PyTorch 参考实现的随机对拍、稀疏路径与 dense-masked torch.topk 的选中值一致性（含 CUDA graph capture 与后续层复用 schedule）。
• vllm/v1/core/kv_cache_utils.py（模块 缓存分配；类别 source；类型 core-logic；符号 _get_kv_cache_bytes_per_block）: KV 分配改动：_get_kv_cache_bytes_per_block 收集所有 MLA spec 的 block_stride_alignment 并对齐块字节数，是 DeepGEMM 分页核 512B 步长要求的落地点。
• vllm/v1/kv_cache_interface.py（模块 缓存规格；类别 source；类型 data-contract；符号 MLAAttentionSpec.block_stride_alignment, MLAAttentionSpec.merge）: 给 MLAAttentionSpec 新增 block_stride_alignment 字段并纳入 merge() 的一致性断言，是跨层 KV 布局契约的一部分。
• vllm/v1/attention/backends/mla/indexer.py（模块 分块预算；类别 source；类型 core-logic；符号 DeepseekV32IndexerMetadataBuilder._prefill_split_seq_lens）: 抽出 _prefill_split_seq_lens 钩子，让稀疏路径可以按 bf16 稀疏行宽而非上下文长度给 prefill 分块器做预算。
• vllm/models/deepseek_v4/common/ops/fused_indexer_q.py（模块 融合算子；类别 source；类型 core-logic；符号 fused_indexer_q_rope_quant, _indexer_weights_out_dtypes）: 融合 Q RoPE-quant 核新增 weights_out_dtype，让 per-head weights 直接以 bf16 产出（Triton 与 CuTeDSL 两条实现同步改动），消除热路径 cast。
• vllm/config/attention.py（模块 配置开关；类别 source；类型 configuration；符号 indexer_sparse_logits）: 新增用户可见开关 indexer_sparse_logits，是整条稀疏路径的总闸门。
• tests/kernels/test_fused_indexer_q_rope_quant.py（模块 算子测试；类别 test；类型 test-coverage；符号 test_fused_indexer_q_rope_quant_writes_bf16_weights）: 验证融合 Q 核的 bf16 权重输出等于 fp32 结果四舍五入一次，覆盖 Triton 与 CuTeDSL 两条实现。
• vllm/models/deepseek_v4/nvidia/ops/fused_indexer_q_cutedsl.py（模块 算子后端；类别 infra；类型 infrastructure；符号 fused_indexer_q_rope_quant）: CuTeDSL 版本的融合 Q 核同步支持 weights 输出 dtype，保证两条后端行为一致。

关键符号：SparseMQAIndexer.init, SparseMQAIndexer.forward, SparseMQAIndexer._reserve_workspaces, _prefill_k_workspaces, _gather_prefill_chunk_k, _expand_candidates_kernel, _sparse_topk_remap_kernel, candidate_blocks_to_sparse_indices, sparse_topk_remap, sparse_mqa_logits_prefill_chunk, sparse_mqa_logits_paged_decode, pick_sparse_block_kv, has_deep_select, check_deep_select_layout, DeepseekV41SparseIndexerMetadataBuilder.init, DeepseekV41SparseIndexerMetadataBuilder.build, DeepseekV41SparseIndexerMetadataBuilder._prefill_split_seq_lens, has_deep_gemm_sparse_mqa, get_sparse_mqa_logits_metadata, get_paged_sparse_mqa_logits_metadata, DeepseekV32IndexerMetadataBuilder._prefill_split_seq_lens, _get_kv_cache_bytes_per_block, MLAAttentionSpec.merge, fused_indexer_q_rope_quant

关键源码片段

vllm/model_executor/kernels/attention/dsa/sparse_mqa_logits.py

新增的稀疏 MQA logits 内核层：候选块展平/过滤/排序、DeepSelect top-k 调用与稀疏列到 request-local 位置的重映射，以及 prefill/decode 两条端到端串联函数，是全 PR 的算法核心。

# Sparse-block 大小（token 数）的支持集合：DeepGEMM 稀疏核只接受 16 或 8
SPARSE_BLOCK_KV_CHOICES = (16, 8)
​
# 排序哨兵。注意：review 指出模块级调用 tl.constexpr(...) 在无 Triton 的平台上
# 会导入即崩（vllm.triton_utils 的占位对象里 constexpr 为 None），普通整数即可
_INT32_BIG = tl.constexpr(2**31 - 1)
​
​
@triton.jit
def _expand_candidates_kernel(
    cand_ptr, cand_stride,
    ks_ptr, ks_stride,
    ke_ptr, ke_stride,
    si_ptr, si_stride,
    end_ptr, end_stride,
    K: tl.constexpr, S: tl.constexpr,
    RATIO: tl.constexpr, SBK: tl.constexpr,
):
    """每行一个 program：展平候选块 -> 过滤越界 -> 升序排序 -> 重复末尾补齐。
​
    生产环境的候选块本身唯一，所以这里只排序不去重；``end`` 统计有效稀疏列数，
    由于块 id 已排序，有效列一定构成前缀。
    """
    row = tl.program_id(0)
    ks = tl.load(ks_ptr + row * ks_stride)
    ke = tl.load(ke_ptr + row * ke_stride)
    cols = tl.arange(0, K)
    cand = tl.load(cand_ptr + row * cand_stride + cols) # K 是 2 的幂，无需 mask
    if RATIO == 2:
        # 候选块 16 个位置、稀疏块 8 个位置时，一个候选块拆成两个稀疏块
        cand = tl.interleave(cand * RATIO, cand * RATIO + 1)
​
    # 只保留落在本行 KV 区间 [0, ke - ks) 内的块；候选块非负，下界无需检查
    kv_len = tl.maximum(ke - ks, 0)
    max_block = (kv_len + SBK - 1) // SBK
    valid = (cand >= 0) & (cand < max_block)
    vals = tl.sort(tl.where(valid, cand, _INT32_BIG))
​
    # 锚定到本行 ks 向下取整后的稀疏块边界（paged decode 路径 ks == 0，锚点恒为 0）
    base = ks // SBK
    si = tl.where(vals != _INT32_BIG, vals, 0) + base
    n_valid = tl.sum(valid.to(tl.int32))
    last_rel = tl.max(tl.where(valid, cand, -1))
    pad_val = tl.where(n_valid > 0, last_rel + base, base)
    tl.store(
        si_ptr + row * si_stride + tl.arange(0, S),
        tl.where(tl.arange(0, S) < n_valid, si, pad_val),
    )
​
    # 有效稀疏 token 列数：只有最后一个块可能是半块，其余都是整块有效
    last_start = last_rel * SBK + ks
    fill = tl.minimum(tl.maximum(ke - last_start, 1), SBK)
    end = tl.where(n_valid > 0, (n_valid - 1) * SBK + fill, 0)
    tl.store(end_ptr + row * end_stride, end)
​
​
def candidate_blocks_to_sparse_indices(
    candidate_blocks, row_ks, row_ke,
    candidate_block_size, sparse_block_kv, out=None,
):
    """把 request-local 候选块 id 展平成 DeepGEMM 需要的稀疏块 id（排序 + 补齐 + 计数）。
​
    调用方可传入 ``(sparse_indices, end)`` 复用缓冲，热路径因此零分配；
    行 stride 通过 ``.stride(0)`` 读取，因此支持列切片得到的非连续输入。
    """
    rows, num_candidates = candidate_blocks.shape
    ratio = candidate_block_size // sparse_block_kv
    assert ratio * sparse_block_kv == candidate_block_size and ratio in (1, 2)
    assert num_candidates & (num_candidates - 1) == 0, num_candidates
    num_sparse = num_candidates * ratio
    if out is None: # 未复用缓冲时退化为分配，主要供测试与首次调用使用
        int32 = dict(dtype=torch.int32, device=candidate_blocks.device)
        out = (
            torch.empty(rows, num_sparse, **int32),
            torch.empty(rows, **int32),
        )
    sparse_indices, end = out
    _expand_candidates_kernel[(rows,)](
        candidate_blocks, candidate_blocks.stride(0),
        row_ks, row_ks.stride(0),
        row_ke, row_ke.stride(0),
        sparse_indices, sparse_indices.stride(0),
        end, end.stride(0),
        K=num_candidates, S=num_sparse, RATIO=ratio, SBK=sparse_block_kv,
        num_warps=16,
    )
    return sparse_indices, end

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

新增后端与元数据：定义 SparseMQARowsMetadata/DeepseekV41SparseIndexerMetadata，并在 builder 中完成所有前置条件的启动期校验与行级 scratch 的持有，是多层共享一次候选展平与 schedule 的关键。

@dataclass
class SparseMQARowsMetadata:
    """单次稀疏 logits 调用的行级状态：decode 的各行，或一个 prefill chunk。
​
    这些缓冲都是 builder 持有的视图；``kernel_metadata``（DeepGEMM schedule）
    与本行的 ``sparse_indices`` / ``end`` 一起由本 step 的第一个消费者层写入，
    后续消费者层直接复用，因此同一 step 只做一次候选展平。
    """
​
    row_ks: torch.Tensor # [rows] int32 K 区间起点；paged decode 全 0
    row_ke: torch.Tensor # [rows] int32 K 区间终点，或压缩后的上下文长度
    sparse_indices: torch.Tensor # [rows, num_sparse_blocks] int32 稀疏块 id
    end: torch.Tensor # [rows] int32 每行有效稀疏列数
    col_indices: torch.Tensor # [rows, topk] int32 稀疏列 top-k 暂存
    block_table: torch.Tensor | None = None # decode 专用页表
    row_indices: torch.Tensor | None = None # decode 专用 行 -> 请求 映射
    kernel_metadata: torch.Tensor | None = None
​
​
class DeepseekV41SparseIndexerMetadataBuilder(DeepseekV32IndexerMetadataBuilder):
    """dense 索引器元数据 + 稀疏 logits 的行级 scratch。
​
    稀疏核的每一条前置条件都在这里、也就是引擎启动时校验，不满足直接报错，
    而不是在每步推理时退回 dense 路径。
    """
​
    def __init__(self, kv_cache_spec, layer_names, vllm_config, device, **kwargs):
        super().__init__(kv_cache_spec, layer_names, vllm_config, device, **kwargs)
        hf_config = vllm_config.model_config.hf_text_config
        candidate_block_size = getattr(hf_config, "candidate_block_size", 0)
        candidate_topk_blocks = getattr(hf_config, "candidate_topk_blocks", 0)
        topk_tokens = getattr(hf_config, "index_topk", 0)
        prefix = "attention_config.indexer_sparse_logits"
        # fail-closed：模型必须两级索引 + MXFP4 indexer cache + SM100 + DeepGEMM >= 2.8
        if candidate_block_size <= 0 or candidate_topk_blocks <= 0:
            raise ValueError(f"{prefix} requires a two-level (candidate block) indexer model ...")
        if not dsa_indexer_uses_fp4(vllm_config):
            raise ValueError(f"{prefix} requires the MXFP4 indexer cache ...")
        if not has_deep_gemm_sparse_mqa():
            raise ValueError(f"{prefix} requires DeepGEMM >= 2.8 ...")
        if self.dcp_world_size > 1 or self.use_pcp:
            raise NotImplementedError(f"{prefix} is not supported with decode/prefill context parallel.")
        # top-k 只能跑在 bf16 logits 上的 DeepSelect，路径上没有 fp32 回退
        if not has_deep_select():
            raise ValueError(f"{prefix} requires the DeepSelect top-k extension ...")
        if vllm_config.kernel_config.sparse_indexer_topk_backend not in ("auto", "deep_select"):
            raise ValueError(f"{prefix} uses DeepSelect for the top-k; ...")
​
        self.sparse_block_kv = pick_sparse_block_kv(candidate_block_size)
        self.num_sparse_blocks = candidate_topk_blocks * (
            candidate_block_size // self.sparse_block_kv
        )
        num_sparse_cols = self.num_sparse_blocks * self.sparse_block_kv
        check_deep_select_layout(num_sparse_cols, topk_tokens)
        # 稀疏 logits 是 bf16 [rows, num_sparse_cols]；让按 fp32 [rows, seq_len]
        # 做预算的 prefill 分块器按这个宽度切片，避免把行提前切碎
        self.min_split_seq_len = (num_sparse_cols * 2 + 3) // 4
​
        # 行号按 batch token 排布（decode 行在前，prefill chunk 按 token 区间），
        # 行数与 ``arange_buffer`` 一致；scratch 全部由 builder 持有，热路径零分配
        max_rows = self.arange_buffer.shape[0]
        int32 = dict(dtype=torch.int32, device=device)
        self.sparse_indices_buffer = torch.zeros((max_rows, self.num_sparse_blocks), **int32)
        self.end_buffer = torch.zeros(max_rows, **int32)
        self.col_indices_buffer = torch.zeros((max_rows, topk_tokens), **int32)
        self.zero_row_ks = torch.zeros(max_rows, **int32)
​
    def _prefill_split_seq_lens(self, seq_lens_cpu: torch.Tensor) -> torch.Tensor:
        # 覆写基类钩子：分块预算按稀疏列宽而不是真实上下文长度计算
        return seq_lens_cpu.clamp(min=self.min_split_seq_len)
​
    def build(self, common_prefix_len, common_attn_metadata, fast_build=False):
        base = super().build(common_prefix_len, common_attn_metadata, fast_build)
        metadata = DeepseekV41SparseIndexerMetadata(
            **{f.name: getattr(base, f.name) for f in fields(base)},
            sparse_block_kv=self.sparse_block_kv,
        )
        # varlen 布局把每个 query 压成一行，所以 base 的 decode 张量已是逐行语义
        ...
        return metadata

评论区精华

评审集中在「删繁就简」和「不要在热路径上留退路」两条主线：

WoosukKwon（sparse_mqa_indexer.py）：We don't need these as DSV4.1 does not use torch.compile. We also don't use CustomOp now.

作者随即在第二个 commit 把 SparseMQAIndexer 从 CustomOp 改为普通 nn.Module，并说明 v4.1 的 attention 层本就在 breakable-CUDA-graph 的 eager 段里执行 indexer。

WoosukKwon：When do we take the "slow" path? I think we should just raise an error in this case, since we never want to execute such unoptimized PyTorch-heavy code in practice.（对 dtype 转换追问：Why not do this conversion in a previous kernel?）

结果：候选展平与索引重映射只保留 Triton、top-k 只保留 DeepSelect，PyTorch 参考实现搬到测试里当精度基线（作者回复 “moved to test as accuracy baseline”）；bf16 转换前移到融合 Q 核，fused_indexer_q_rope_quant 新增 weights_out_dtype（第三个 commit）。

WoosukKwon：Why not put these methods on vllm/model_executor/layers/sparse_mqa_indexer.py instead?（并建议 We can support non-contiguous cases by adding stride to the Triton kernel's input）

reserve_indexer_workspaces 一类辅助函数被迁入新模块，Triton 内核改为显式 stride，测试也覆盖了 strided 输入与 strided 输出缓冲。

ywang96（reviewed at 16a2fee）：Thanks for the unusually well-documented PR … one must-fix below（import-time crash on Triton-less platforms），the rest is non-blocking. Verdict: approve with comments

他同时指出 _INT32_BIG = tl.constexpr(2**31 - 1) 在模块级调用，Triton 缺失时 vllm.triton_utils 的占位对象 constexpr = None 不可调用，而 sparse_attn_indexer.py 无条件 import 该模块，建议一行改为普通整数。Copilot 则提了两点合约问题：sparse_topk_remap 每次 decode 新分配 workspace 且 top-k 硬编码 stride (width, 1)；docstring 声称候选块 “unique and strictly increasing” 但实现只排序不去重——最终合约改为 “sorted ascending … unique in production”。

• 无 Triton 平台上的模块级 tl.constexpr 导入崩溃 (correctness): 被标记为 must-fix 并在合入前要求修复；审阅同时确认 candidate_blocks.py 只在类型注解里用 tl.constexpr，没有同类问题。
• 生产路径是否保留 PyTorch fallback / “慢路径” (design): 作者删除候选展平/重映射的 PyTorch fallback（只保留 Triton）与 fp32 转换分支，PyTorch 参考实现移到测试里作为精度基线（回复 “moved to test as accuracy baseline”），所有不满足条件的情况在元数据构建器里 raise。
• per-head weights 的 bf16 转换放在哪里 (performance): 第三个 commit 让 fused_indexer_q_rope_quant 新增 weights_out_dtype，Triton 与 CuTeDSL MXFP4 变体直接写 bf16；SparseMQAIndexer.weights_dtype 与 DeepseekV4Indexer.indexer_weights_dtype 串联起这条 dtype 契约，热路径不再有 cast。
• 非连续输入：报错还是让内核支持 stride (design): _expand_candidates_kernel 与 _sparse_topk_remap_kernel 全部改为显式 stride 参数，测试覆盖列切片候选行与 strided 输出缓冲。
• workspace 预留辅助函数的归属 (design): reserve_indexer_workspaces 一类逻辑迁入新模块，成为 SparseMQAIndexer._reserve_workspaces 与 _prefill_k_workspaces，旧文件不再承担稀疏路径职责。
• 候选块 / 稀疏块 id 的单调性与唯一性合约 (documentation): 合约改为 “sorted ascending per row … unique in production, where candidates are unique”，并说明重复会被容忍但会浪费算力；测试里专门保留了重复块与 -1 padding 的用例。
• decode 热路径是否需要每次分配 workspace (performance): 最终形态把所有中间缓冲（sparse_indices / end / col_indices）交给元数据构建器持有并通过参数传入，SparseMQAIndexer.forward 不做任何分配。
• 不要 CustomOp / torch.compile 相关机制 (design): SparseMQAIndexer 改为普通 nn.Module（模型侧也把类型注解改为 SparseAttnIndexer | SparseMQAIndexer），理由是 v4.1 的 attention 层本就在 breakable-CUDA-graph 的 eager 段里运行 indexer。

风险与影响

• 风险：
  1. 无 Triton 平台的导入期崩溃（must-fix，需确认最终形态）：vllm/model_executor/kernels/attention/dsa/sparse_mqa_logits.py 的 _INT32_BIG = tl.constexpr(2**31 - 1) 是模块级调用；本材料给出的 head 片段里它仍是 tl.constexpr(...)，若未按 ywang96 建议改成普通整数，CPU 等无 Triton 构建只要 import 到 sparse_attn_indexer.py 就会 TypeError。
  2. 精度与输出可复现性：稀疏核只支持 bf16 权重与 bf16 logits（无 fp32 选项），与 dense 路径的 fp32 打分不一致。PR body 明确 greedy 输出不再与 dense 逐 token 相同，78K 文档的 NLL 差异约 0.002（与 TP=4 的 run-to-run 噪声同量级）；对精度敏感的评测或回归对比需要重新设定基线。
  3. 启动期硬失败、无回退：DeepseekV41SparseIndexerMetadataBuilder.__init__ 的所有校验都 raise；仓库 pin 的 DeepGEMM 仍早于 2.8，用户在不满足 SM100/MXFP4/DeepSelect/非 DCP-PCP/非 varlen 的环境上打开开关会直接起不来。这符合 fail-closed 设计，但运维上需要清晰的错误信息与回滚路径。
  4. KV 内存与全局分配逻辑：vllm/v1/kv_cache_interface.py 的 block_stride_alignment 会被 MLAAttentionSpec.merge() 的一致性断言检查，vllm/v1/core/kv_cache_utils.py:_get_kv_cache_bytes_per_block 的 round_up 作用在所有 KV cache group 上；虽然关闭开关时该字段为 None，仍属于全局分配路径的改动，开启后固定增加约 0.8% KV 显存。
  5. 测试覆盖的环境依赖：新增用例大量 skipif（需 SM100、deep_gemm、vllm._deepselect_C），在普通 CI 机器上会被跳过，kernel 与端到端性能/精度回归实际靠作者在 4x GB300 上离线跑；缺少 CI 常驻的性能门禁。
  6. 共享 scratch 的隐含假设：同一 step 的所有消费者层复用第一次计算的 sparse_indices / end / kernel_metadata，正确性依赖「各消费者层看到的候选块完全一致」以及层执行顺序；如果模型配置变化或层顺序调整，可能出现静默的错配，目前只靠注释与测试（later-layer reuse 用例）保障。
  7. 融合 Q 核 ABI 变化：weights_out_dtype 引入新参数，Triton 与 CuTeDSL 两条实现都要覆盖，JIT warmup key 必须同时覆盖两种 dtype（commit message 提到这一点），否则首次请求会出现编译/预热抖动。
     - 影响：用户侧：这是一个 opt-in 开关 --attention-config '{"indexer_kv_dtype": "mxfp4", "indexer_sparse_logits": true}'，只影响 DeepSeek V4.1-Flash 类两级索引模型的长上下文服务；1M 上下文预填充吞吐提升到 2.00x、512K 1.43x、256K 1.19x，decode 稳定提升 3-6%，但 8K 左右短上下文会比 dense 路径慢约 5%（PR body 因此把默认值保持为 off）。系统侧：新增一条 attention backend 与 metadata 类型，KV cache 分配新增块步长对齐字段并对齐到 lcm(512, page_alignment)，属于跨模块（config → kv_cache_interface → kv_cache_utils → model → kernel）的联动改动，任何一处不一致都会在启动时报错而不是静默降级。团队侧：这是 vLLM 首次把 DeepGEMM 的 sparse MQA-logits 内核接进 DSA 索引器，建立了「候选块展平 + 压缩空间 top-k + 位置重映射」的实现范式与对应测试基建（PyTorch 参考实现 + 选中值对拍），后续其他稀疏索引模型（如 glm5next 一类）可以复用这套结构；同时把对 DeepGEMM >= 2.8 与 DeepSelect 扩展的依赖显式化，对发布/镜像 pin 版本提出了新要求。
     - 风险标记：无 Triton 平台导入风险, 精度由 fp32 降为 bf16, 启动期硬失败无回退, KV 块步长对齐影响分配路径, CI 覆盖依赖 SM100 会被跳过, 共享 scratch 的层序依赖

关联脉络

• PR #56464 （历史列表未提供标题，按 PR body 引用）DeepSelect top-k 扩展: PR body 明确说明 “DeepSelect (#56464) runs the top-k on the bf16 logits directly”，本 PR 的 sparse_topk_remap 与元数据校验（has_deep_select()、check_deep_select_layout）都依赖该扩展提供的 vllm._deepselect_C。
• PR #56441 [Perf][DSpark] Add KV-only context insertion across V4.1 cache formats: 同属 DeepSeek V4.1（vllm/models/deepseek_v41/）的长上下文性能线，同样改造 V4.1 注意力/KV 写入内核路径；本 PR 是这条脉络上第二级索引器的性能补强。
• PR #55358 [Refactor][GLM-5.3-Flash] Move sparse_attn_indexer_kpool into the model folder and split AMD/NVIDIA: 同样是稀疏索引器（sparse indexer）的模块化与后端拆分工作，与本 PR “新建独立 SparseMQAIndexer 模块 + 独立 backend” 的组织方式属于同一演进方向，可对照阅读两套拆分粒度。