执行摘要

• 一句话：修复 SM120 FP8 KV 下 HiSparse 启动被 DSA 后端校验误拒
• 推荐动作：建议精读。该 PR 展示了参数解析与校验之间的协同设计，暴露了分散验证导致启动失败的典型问题。修复思路（让校验认可解析结果）简单有效，值得在类似多关卡配置校验中借鉴。另外，PR 中附带的性能对比数据为 HiSparse 在 SM120 上的价值提供了证据。

功能与动机

PR body 明确指出 --enable-hisparse 在 SM120/SM121 上完全无法启动，因为两条参数管线对合法 DSA 后端判断不一致：解析层已选择 flashinfer_sparse_mla（SM120 上唯一可用的 DSA 内核），而校验层只允许 flashmla_kv。错误信息甚至提出了不可满足的建议（SM120 上无 flashmla_kv），导致 omit 也会再次失败。该修复只是让启动校验认可解析器自己选出的后端。

实现拆解

1. 修改 python/sglang/srt/arg_groups/hisparse_hook.py：在 HISPARSE_CUDA_DSA_BACKENDS_BY_DTYPE["fp8_e4m3"] 集合中加入 flashinfer_sparse_mla，并在 _hisparse_allowed_backends 的未知 dtype fallback 集合中也加上该后端，避免未来新 dtype 再次出现校验与解析脱节。
2. 重写 validate_hisparse_dsa_backend 的错误提示，从"建议单一默认后端"改为"列出全部允许后端"，避免给出不可满足的建议。
3. 在 test/registered/unit/server_args/test_server_args.py 新增 test_hisparse_accepts_flashinfer_sparse_mla_on_cuda_fp8，验证在 CUDA FP8 KV 下 prefill/decode validator 均接受该后端。
4. 更新 docs/docs/advanced_features/hisparse_guide.mdx，在 DSA 后端说明和 Key Notes 中补充 SM120/SM121 GLM + FP8 的特例。
5. 后置校验 _validate_flashinfer_sparse_mla_backend 保持不变，仍只允许 GLM DSA + FP8 KV + SM120/SM121 组合。

关键文件：

• python/sglang/srt/arg_groups/hisparse_hook.py（模块 参数校验；类别 source；类型 core-logic）: 核心校验逻辑所在文件，修改了 fp8_e4m3 后端的允许集合，并优化错误提示，是修复启动失败的关键。
• test/registered/unit/server_args/test_server_args.py（模块 参数测试；类别 test；类型 test-coverage；符号 test_hisparse_accepts_flashinfer_sparse_mla_on_cuda_fp8）: 新增单元测试验证 CUDA FP8 下 HiSparse 接受 flashinfer_sparse_mla 后端，防止回归。
• docs/docs/advanced_features/hisparse_guide.mdx（模块 用户文档；类别 docs；类型 documentation）: 文档更新，明确 SM120/SM121 GLM+FP8 下的后端特例，降低用户困惑。

关键符号：validate_hisparse_dsa_backend, _hisparse_allowed_backends, test_hisparse_accepts_flashinfer_sparse_mla_on_cuda_fp8

关键源码片段

python/sglang/srt/arg_groups/hisparse_hook.py

核心校验逻辑所在文件，修改了 fp8_e4m3 后端的允许集合，并优化错误提示，是修复启动失败的关键。

# python/sglang/srt/arg_groups/hisparse_hook.py
​
# 按 KV cache dtype 划分的 CUDA 合法 DSA 后端集合。
# fp8_e4m3 在 SM120 上只有 flashinfer_sparse_mla 可用，因此必须同时允许。
HISPARSE_CUDA_DSA_BACKENDS_BY_DTYPE = {
    "bfloat16": {"flashmla_sparse"}, # BF16 保持原有唯一后端
    "fp8_e4m3": {"flashmla_kv", "flashinfer_sparse_mla"}, # 新增 SM120 专属后端
}
​
HISPARSE_ROCM_DSA_BACKENDS = {"tilelang", "aiter"} # ROCm 平台不变
HISPARSE_KV_CACHE_DTYPES = ("bfloat16", "fp8_e4m3") # HiSparse 仅支持这两类 KV dtype
​
​
def _hisparse_default_backend(kv_cache_dtype: str) -> str:
    """返回该 dtype 下的默认后端（用于错误提示，不再强制单一建议）。"""
    if _is_hip():
        return "tilelang"
    return "flashmla_kv" if kv_cache_dtype == "fp8_e4m3" else "flashmla_sparse"
​
​
def _hisparse_allowed_backends(kv_cache_dtype: str) -> set[str]:
    """返回当前平台与 dtype 下的全部合法后端集合。
​
    未知 dtype 的 fallback 集合也补入 flashinfer_sparse_mla，
    避免将来引入新 dtype 时再次出现校验与解析不一致。
    """
    if _is_hip():
        return HISPARSE_ROCM_DSA_BACKENDS
    return HISPARSE_CUDA_DSA_BACKENDS_BY_DTYPE.get(
        kv_cache_dtype, {"flashmla_sparse", "flashmla_kv", "flashinfer_sparse_mla"}
    )
​
​
def validate_hisparse_dsa_backend(server_args: ServerArgs, attr: str, label: str) -> None:
    """校验 DSA 后端是否在 HiSparse 允许列表中。
​
    通过 resolved_view 读取解析后的后端，避免与 overrides 的解析逻辑脱节。
    """
    from sglang.srt.arg_groups.overrides import resolved_view
​
    view = resolved_view(server_args)
    backend = getattr(view, attr)
    kv_cache_dtype = view.kv_cache_dtype
    allowed_backends = _hisparse_allowed_backends(kv_cache_dtype)
    if backend is not None and backend not in allowed_backends:
        raise ValueError(
            f"HiSparse supports DSA {label} backend(s) {sorted(allowed_backends)} "
            f"on this platform with --kv-cache-dtype={kv_cache_dtype}, "
            f"but got --dsa-{label}-backend={backend}. "
            f"Please use one of {sorted(allowed_backends)}, or omit the option "
            "to let SGLang pick a backend for this platform."
        )

评论区精华

核心讨论围绕 SM120 上 HiSparse 的性能价值。b8zhong 质疑 H2D 带宽较低时 HiSparse 是否还有收益，作者 gongwei1027 用 8×RTX PRO 6000 的实测数据回应：在 BS≥8 时 HiSparse 显著提升吞吐（BS=54 时 3.0×），因为收益来自打破 KV 内存墙而非 kernel 速度，并解释无 HiSparse 时 decode 实例在 running-req=6 就饱和 KV 池。samuellees 评价"A cool job!"。讨论结论：HiSparse 在 SM120 高并发下仍值得开启。

• SM120 上 HiSparse 性能收益是否因 H2D 带宽受限而打折扣 (performance): 作者用实测数据证明 SM120 上 HiSparse 在高并发下仍有明显收益，低 batch 下略降。

风险与影响

• 风险：该改动扩大了 fp8_e4m3 下的后端允许集合，理论上允许用户在非 SM120/GLM 组合下显式指定 flashinfer_sparse_mla 通过 HiSparse 校验；但后续 _validate_flashinfer_sparse_mla_backend 仍会拒绝不支持的架构/模型组合，因此安全性由后置校验兜底。风险点包括：若该后置校验存在绕过路径，可能引入隐性配置错误；单元测试仅覆盖 CPU 上的 allow-set，未覆盖真实 SM120 硬件（CI 无该配置），端到端验证依赖手动硬件测试。整体风险较低，但对平台行为的变更需要留意。
• 影响：对 SM120/SM121 用户：修复了 --enable-hisparse 启动失败的问题，使 GLM DSA + FP8 KV 可以无缝使用 HiSparse。对其他平台：BF16 KV、ROCm 的允许集合不变，行为无影响。对团队：需要维护两处后端列表（解析层与校验层）的一致性；文档新增了硬件特例说明，降低用户困惑。
• 风险标记：允许列表与解析器耦合风险, 缺少 SM120 硬件 CI 覆盖, 后置校验兜底

关联脉络

• PR #34329 HiSparse: shared-index (IndexShare) plan-then-IO swap-in prefetch: 同一 HiSparse 功能线，该 PR 在当前修复之上继续扩展预取能力，互为演进。
• PR #26928 flashinfer_sparse_mla SM120 kernel support (referenced in PR body): PR body 提到该 PR 落地了 flashinfer_sparse_mla 的 SM120 kernel 路径，是本次允许列表修改的前提。