[Test] Complete Kimi-K3 Prefill MoE backend report

2026-08-19 16:54:45 +08:00 · 2026-08-19 16:54:45 +08:00 · 08a35066d7
commit 08a35066d7
parent 74ec19dd48
4 changed files with 351 additions and 693 deletions
--- a/README.md
+++ b/README.md
@ -1,5 +1,7 @@
 # sskj — 多平台大模型推理性能基准测试项目

+**更新（2026-08-19 16:51:43 CST）**：完成 Kimi-K3 / RTX 6000D / SGLang Prefill MoE backend 交付文档收敛与 Chunk=4K 补测。四节点 TP32×EP4、16K→1、C=8/16 下 12/12 重复、480/480 请求成功；FlashInfer 相对 Marlin 提升约 9.8% Input TPS、降低约 9.0% TTFT。结合 4K/8K/16K 全部结果，最终推荐保持 `flashinfer_mxfp4 + chunked_prefill_size=8192`。
+
 **更新（2026-08-19 14:23:52 CST）**：601 的 GitLab SSH key 已生效，内部 Kimi-K3 SM120 审计分支已推送至 `qqtang/sskj:hzy-kimi-k3-sm120-flashinfer-mxfp4`。远端已有普通分支 `hzy`，因此不能创建 `hzy/...` ref，最终采用无层级冲突的分支名；GitHub SGLang feature 分支保持不变。

 **更新（2026-08-19 13:58:50 CST）**：用户同步 GitHub fork 后，将 Kimi-K3 SM120 SGLang 的 7 个提交无冲突重放到最新 `main` `593b1a9`，`range-diff` 逐提交确认补丁等价；最终 Draft HEAD `b95b534` 已推送至 `ACEEE-1222/sglang:hzy/kimi-k3-sm120-flashinfer-mxfp4`。在 601 统一验证镜像中重跑最终 HEAD 的语法与定向回归，结果为 `10 passed, 3 deselected`；代表性四机性能数据保留，不重复消耗 GPU。
--- a/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/README.md
+++ b/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/README.md
@ -1,725 +1,182 @@
 # Kimi-K3 SM120 FlashInfer MXFP4 MoE

-> **2026-08-18 去重审计更正**：下文 Phase 0-4 保留当时的调查与实验
-> 时间线，但“没有开放 PR、需要独立实现 FlashInfer SiTU kernel”的判断已失效。
-> 当前 kernel 依赖方向是尚未合并的 FlashInfer PR #4460；我们不提交
-> FlashInfer PR，SGLang Draft 只保留
-> Kimi-K3 SGLang 接入、精确 MXFP8×MXFP4+SiTU 测试和 6000D 实测证据。
-> 权威边界见 [`UPSTREAM_DUPLICATION_AUDIT.md`](UPSTREAM_DUPLICATION_AUDIT.md)。
-
-当前 SGLang Draft 基于 `main` `593b1a9b8a`，HEAD 为 `b95b5341ab95`，已推送到
-`ACEEE-1222/sglang:hzy/kimi-k3-sm120-flashinfer-mxfp4`。公开 PR 正文见
-[`SGLANG_DRAFT_PR.md`](SGLANG_DRAFT_PR.md)，可审查补丁见
-[`patches/sglang-kimi-k3-sm120-flashinfer-mxfp4-draft.patch`](patches/sglang-kimi-k3-sm120-flashinfer-mxfp4-draft.patch)。
-五个变更文件已通过完整 pre-commit；同步 fork 后，7 个提交的
-`range-diff` 全部为补丁等价，最终非 JIT 回归为
-`10 passed, 3 deselected, 17 warnings in 33.30s`。
-逐项验收和唯一剩余交付项见
-[`COMPLETION_AUDIT.md`](COMPLETION_AUDIT.md)。
-
-最终 PR 代表点已经在同步前的等价 patch `300c87a` 统一镜像中完成：四节点
-TP32/EP4、16K→1、C=8、Chunk=8K，Marlin 与 FlashInfer 各三次重复，
-240/240 请求成功。FlashInfer 相对 Marlin 的 Input TPS 中位数为
-`+10.96%`，TTFT P50/P95 为 `-9.91%/-9.92%`。轻量结果见
-`results/pr_representative_16k_c8_chunk8k_summary.json` 和
-`results/pr_representative_16k_c8_chunk8k_results.csv`。
-同步后的 `b95b534` 与该 7 提交系列逐提交等价，因此保留原始性能证据，
-不把容器中的旧 commit object 错写成最终 HEAD，也不重复运行四机实验。
-
 ## 目标

-为 Kimi-K3 在 NVIDIA SM120（RTX PRO 6000D）上的 SGLang 部署补齐 FlashInfer MXFP8 activation × MXFP4 weight fused MoE，并形成可回归、可性能比较、可提交上游的最小改动。
+为 Kimi-K3 在 NVIDIA SM120（RTX 6000D）上的 SGLang 部署补齐
+FlashInfer MXFP8 activation × MXFP4 weight fused MoE 支持，并形成：

-本目录不修改已有 Kimi-K3 基线实验。功能实现必须先通过算子正确性，再进入模型加载和性能测试。
+- 可上游审查的最小 SGLang 改动；
+- 精确覆盖 Kimi-K3 的正确性测试；
+- RTX 6000D 单卡与四节点真实服务证据；
+- Marlin 与 FlashInfer 的同口径性能结论。

-## Phase 0 审计结论
+本目录不修改已有 Kimi-K3 基线实验，也不提交与 FlashInfer 上游重复的 kernel。

-状态：**完成**。
+## 交付状态

-审计环境：
+状态：**完成，SGLang Draft 代码与审计分支已推送。**

-| 对象 | 版本或提交 |
+| 项目 | 结果 |
 |---|---|
-| GPU | NVIDIA RTX PRO 6000 Blackwell Server Edition，SM120 |
-| SGLang | `827552bc1dda8507f2a1b11c7274340a83da49ce` |
-| FlashInfer | `a9e03bf2f5fe`，Python package `0.6.18` |
-| PyTorch / CUDA | PyTorch 2.11，CUDA 13 |
-| 构建镜像 | `lmsysorg/sglang:kimi-k3-fiv617situ-warm` |
-| Kimi-K3 | hidden 7168，expert hidden 3584，896 experts，top-k 16，SiTU `(alpha=4.0, beta=25.0)` |
+| GPU | NVIDIA RTX 6000D，SM120 |
+| 模型 | Kimi-K3，MXFP4，SiTU `(4.0, 25.0)` |
+| SGLang base | `593b1a9b8aeedb908ee21870c02a0fd1a3dc30a5` |
+| SGLang Draft HEAD | `b95b5341ab9555ba5ea53f5828f6de203d749cce` |
+| GitHub 分支 | `ACEEE-1222/sglang:hzy/kimi-k3-sm120-flashinfer-mxfp4` |
+| FlashInfer 依赖 | PR #4460，implementation `b525c513fef7` |
+| 默认行为 | 不改变默认 MoE backend，Marlin 仍可用 |

-已确认：
+公开 PR 正文：[`SGLANG_DRAFT_PR.md`](SGLANG_DRAFT_PR.md)

-1. FlashInfer 当前 main 的 SM120 CUTLASS fused MoE 能完成基础 MXFP8×MXFP4 运算；官方最小测试通过。
-2. Kimi-K3 所需的 `ActivationType.Situ` 尚未接入该 SM120 CUTLASS C++ runner。
-3. Python `ActivationType.Situ` 的值为 10，而 CUTLASS C++ `ActivationType` 的值 10 当前是 `InvalidType`。SiTU 调用因此被当作非门控激活，FC1 门控维度判断错误。
-4. SGLang 当前 SM120 MXFP4 路径仍按 GPT-OSS 的 SwiGLU 权重布局和参数组织，未透传 Kimi SiTU 参数。
-5. 旧目录 `kimi3_pro6000_sglang_tp32ep32_deepep_flashinfer_mxfp4` 实际使用 `flashinfer_cutedsl`，并在模型加载前因量化配置识别失败；它不能作为 FlashInfer MXFP4 SM120 的成功或失败证据。
+可应用补丁：
+[`patches/sglang-kimi-k3-sm120-flashinfer-mxfp4-draft.patch`](patches/sglang-kimi-k3-sm120-flashinfer-mxfp4-draft.patch)

-当时的上游重叠审计遗漏了后续出现的 FlashInfer #4460。最新复核结论是：
+逐项验收：[`COMPLETION_AUDIT.md`](COMPLETION_AUDIT.md)

- FlashInfer #4460 已实现 CUTLASS SiTU kernel 和独立 `situ_beta` /
-  `situ_linear_beta` ABI，当前开放、未合并。
- SGLang #32668 已合并 GPT-OSS SM120 接入，但不覆盖 Kimi 的连续 gate/up
-  半区布局和 SiTU 参数。
- 605 的可运行镜像也是移植 #4460，不是 0.6.17 原版直接支持。
+## 交付内容

-因此独立增量仅位于 Kimi SGLang 接入、精确测试和 6000D 服务证据；不再
-包含平行 FlashInfer kernel。完整文件级审计见 `UPSTREAM_DUPLICATION_AUDIT.md`。
+### SGLang 实现

-## 真机证据
+1. 为 Kimi-K3 在 SM120 上接入 FlashInfer CUTLASS MXFP4 MoE；
+2. 将 Kimi 连续 `[gate; up]` 权重与 scale 转换为 CUTLASS 所需布局；
+3. 独立透传 `situ_beta=4.0` 与 `situ_linear_beta=25.0`；
+4. 在 MXFP8 量化前处理非连续 hidden states；
+5. 排除 SM120 误入 tcgen05 attention-residual 路径；
+6. 保持 GPT-OSS 路径、旧 FlashInfer import 和默认 Marlin 行为不变。

-FlashInfer 源码与 wheel：
+变更文件：

 ```text
-/data/hzy/src/flashinfer-kimi-sm120
-/data/hzy/artifacts/flashinfer-main-a9e03bf/flashinfer_python-0.6.18-py3-none-any.whl
-/data/hzy/cache/flashinfer-sm120
+python/sglang/srt/layers/attn_residual.py
+python/sglang/srt/layers/moe/moe_runner/flashinfer_cutlass.py
+python/sglang/srt/layers/quantization/mxfp4.py
+test/registered/unit/layers/quantization/test_mxfp4_sm120_cutlass.py
+test/registered/unit/layers/test_attn_residual_arch_gate.py
 ```

-官方 SM120 MXFP8×MXFP4 基线：
+### 测试覆盖

-```bash
-FLASHINFER_DISABLE_VERSION_CHECK=1 pytest -q \
-  tests/moe/test_trtllm_cutlass_fused_moe.py::test_moe_mxfp8_mxfp4\[False-None-None-None-otype1-128-2-2-128-1\] -s
-```
+- Kimi gate/up 权重和 scale 布局；
+- SiTU `(4.0, 25.0)` 参数；
+- MXFP8 activation × MXFP4 weight；
+- 非连续输入；
+- standard routing 与 EP expert IDs；
+- eager、CUDA Graph capture/replay；
+- SM120 attention-residual architecture gate；
+- 缺少新 API 时的兼容性与明确报错。

-实际结果：
+## 上游边界

-```text
-1 passed, 3 warnings in 540.58s
-```
+FlashInfer PR #4460 负责公共 CUTLASS SiTU kernel 和参数 API；本工作不提交
+平行的 FlashInfer kernel PR。

-SiTU 最小负例沿用官方 `test_moe` 数据构造，只将激活改为 `ActivationType.Situ`：
+SGLang Draft 只保留 Kimi-K3 的真实框架增量：权重/scale 布局、SiTU 参数、
+非连续输入、SM120 guard 和测试。它不 vendor kernel、不 pin 未合并私人 commit、
+不改变默认 backend。

-```python
-test_moe(1, 128, 2, 2, 128, ActivationType.Situ)
-```
+完整去重审计见：
+[`UPSTREAM_DUPLICATION_AUDIT.md`](UPSTREAM_DUPLICATION_AUDIT.md)

-实际首个错误：
+## 正确性结果

-```text
-fc1_expert_weights inter size must be equal to fc2_expert_weights inter size
-(256 vs. 128)
-```
-
-对应源码事实：
-
-```text
-flashinfer/tllm_enums.py:
-  Situ = 10
-  InvalidType = 11
-
-csrc/nv_internal/tensorrt_llm/kernels/cutlass_kernels/include/common.h:
-  Identity = 9
-  InvalidType = 10
-```
-
-## Phase 1 真实调用链审计
-
-状态：**完成**。本阶段只确认数据契约和最小失败边界，未修改 SGLang 或 FlashInfer 实现。
-
-### Serving 调用链
-
-```text
-KimiK3MoE.forward
-  -> TopK(select_experts): sigmoid + correction bias + top-k 16 + renormalize
-  -> KimiK3MoE._forward_routed
-  -> FusedMoE.run_moe_core
-  -> Mxfp4MoEMethod.apply
-  -> fused_experts_none_to_flashinfer_mxfp4
-  -> flashinfer.fused_moe.cutlass_fused_moe
-  -> flashinfer_cutlass_fused_moe_binding.cu
-  -> CUTLASS grouped GEMM1 -> gated activation -> GEMM2 -> finalize
-```
-
-关键源码：
-
-| 层次 | 文件 | 已确认行为 |
-|---|---|---|
-| 模型 | `sglang/srt/models/kimi_k3.py` | 创建 `FusedMoE` 时传 `activation="situ"`、`gemm1_alpha=4.0`、`gemm1_clamp_limit=25.0`、`gate_up_interleaved=False` |
-| 路由 | `sglang/srt/layers/moe/topk.py` | K3 使用 sigmoid 分数；correction bias 只参与选 expert，最终权重从未加 bias 的 sigmoid 分数 gather，并做归一化 |
-| 量化 | `sglang/srt/layers/quantization/mxfp4.py` | checkpoint 的 packed MXFP4 权重与 E8M0 group-32 scale 被装入 `w13/w2` |
-| SGLang runner | `sglang/srt/layers/moe/moe_runner/flashinfer_cutlass.py` | SM120 当前固定传 `ActivationType.Swiglu`，未读取 `runner_config.activation` |
-| FlashInfer binding | `csrc/fused_moe/cutlass_backend/flashinfer_cutlass_fused_moe_binding.cu` | 根据 C++ `isGatedActivation` 判断 FC1 是否为 GEMM2 intermediate 的两倍 |
-| FlashInfer activation | `csrc/fused_moe/cutlass_backend/cutlass_fused_moe_kernels.cuh` | gated activation 是独立 CUDA kernel；已有 per-expert 参数向量，可做最小 SiTU 扩展，无需重写 grouped GEMM |
-
-### 真实 checkpoint 契约
-
-模型配置来自 `/data/hf_models/Kimi-K3/config.json`：
-
-| 项目 | 值 |
-|---|---:|
-| hidden size | 7168 |
-| routed latent width | 3584 |
-| expert intermediate | 3072 |
-| experts / active experts | 896 / 16 |
-| layers | 93 |
-| activation | SiTU |
-| SiTU gate beta / linear beta | 4.0 / 25.0 |
-| quantization | packed MXFP4 weight，group size 32，E8M0 scale |
-
-第 3 层 expert 0 的真实 safetensors 元数据：
-
-```text
-w1.weight_packed [3072, 1792] U8   w1.weight_scale [3072, 112] U8
-w3.weight_packed [3072, 1792] U8   w3.weight_scale [3072, 112] U8
-w2.weight_packed [3584, 1536] U8   w2.weight_scale [3584,   96] U8
-```
-
-`1792=3584/2`、`1536=3072/2` 是每 byte 两个 FP4 元素。SGLang loader 将独立 `w1` 与 `w3` 顺序装入 `w13` 的两个连续半区，因此加载后是 `[gate; up]`，不是 GPT-OSS 的逐行 `[gate0, up0, ...]`。
-
-SiTU 的参考公式为：
-
-```text
-gate = 4 * tanh(gate / 4) * sigmoid(gate)
-up   = 25 * tanh(up / 25)
-out  = gate * up
-```
-
-### TP/EP expert mapping
-
-SGLang 的关系为：
-
-```text
-moe_tp_size = tp_size / moe_ep_size / moe_dp_size
-num_local_experts = num_global_experts / moe_ep_size
-```
-
-当前四节点基线为 TP32 / EP32 / MoE-DP1，因此：
-
-```text
-moe_tp_size = 32 / 32 / 1 = 1
-num_local_experts = 896 / 32 = 28
-```
-
-每个 rank 持有连续的 28 个完整专家，FC1/FC2 不再按 MoE-TP 切 intermediate。候选 kernel 的真实核心 shape 是 `E_local=28, K=3584, N=3072, topK=16`；输入 token 数随 prefill/decode batch 变化。
-
-### 已确认的差异矩阵
-
-| 契约 | Kimi-K3 正确值 | 当前 SM120 路径 | 影响 |
-|---|---|---|---|
-| Python/C++ enum 10 | `Situ` | C++ 为 `InvalidType` | SiTU 被当成非门控，首先触发 FC1 shape 校验失败 |
-| 激活数学 | SiTU `(4, 25)` | 固定 `Swiglu` | 即使绕过 shape 校验也会产生错误输出 |
-| checkpoint `w13` | 连续 `[gate; up]` | SM120 processor 假定 GPT-OSS 交错行 | `0::2/1::2` 会把 gate/up 行混合打乱 |
-| CUTLASS 期望 `w31` | 连续 `[up; gate]` | GPT-OSS 专用转换 | Kimi 需要按半区交换，不能按奇偶行拆分 |
-| activation 参数 | gate beta=4，linear beta=25 | 固定 SwiGLU alpha=1.702、beta=1、limit=7 | 参数语义完全不同 |
-| routing | 已计算的标准 top-k ids/weights | CUTLASS runner 支持标准 routing | 这一层无需重写 |
-| EP mapping | global ids + EP32/rank -> local 28 experts | CUTLASS ABI 已有 TP/EP 参数 | 需要 correctness 覆盖，但不是当前首错 |
-
-因此这不是单一“参数名不匹配”。至少要同时修复 FlashInfer SiTU activation 和 SGLang Kimi SM120 权重/参数适配，缺一项都不能得到正确结果。
-
-### CUDA Graph 与 fallback
-
- `flashinfer_mxfp4` runner 位于普通 model forward 内，eager 与 CUDA Graph 使用同一条 MoE 调用，不存在另一个自动回退到 Marlin 的 graph 专用分支。
- Decode 只在 batch 能被已捕获 bucket 覆盖且满足 graph 条件时 replay；超过 capture max、宽度不匹配或其他 eligibility 失败时，`ModelRunner._forward_raw` 转入 eager runner。
- Graph replay 会把实际 batch pad 到最近的 captured batch。`tune_max_num_tokens=next_power_of_2(x.shape[0])` 因而看到的是静态 padded token 数。
- 后续必须分别验证 eager、capture 和 replay；“eager 正确”不能代表 CUDA Graph replay 正确。
- `--moe-runner-backend flashinfer_mxfp4` 当前没有运行时按 shape 回退 Marlin 的机制。生产 fallback 应在服务初始化选择阶段明确决定，不能等 kernel 运行失败后热切换。
-
-## 最小复现
-
-目标 MXFP8 activation × MXFP4 weight 的复现文件：
-
-```text
-reproduce_kimi_situ_sm120_mxfp4.py
-```
-
-它复用 FlashInfer 官方测试的数据构造，只把 activation 改为 SiTU 并传 Kimi 的 `4/25` 参数。当前 main 的预期首错是：
-
-```text
-fc1_expert_weights inter size must be equal to fc2_expert_weights inter size
-```
-
-601 上的实际复现命令：
-
-```bash
-cd /data/hzy/sskj/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4
-CUDA_VISIBLE_DEVICES=2 \
-PYTHONPATH=/tmp/hzy-fi-main-phase1 \
-FLASHINFER_DISABLE_VERSION_CHECK=1 \
-/data/hzy/envs/sparse-opd-sm120/bin/python \
-  reproduce_kimi_situ_sm120_mxfp4.py \
-  > phase1_reproduce_sm120_situ.log 2>&1
-```
-
-证据日志：
-
-```text
-/data/hzy/sskj/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/phase1_reproduce_sm120_situ.log
-```
-
-2026-08-14 实测进入 `flashinfer.cutlass_fused_moe` 后，在
-`flashinfer_cutlass_fused_moe_binding.cu:369` 以 `256 vs. 128` 命中上述检查。
-这证明失败发生在 FlashInfer SM120 CUTLASS ABI/activation 边界，不依赖完整模型、SGLang scheduler、DeepEP 或多节点网络。
-
-## Phase 1 决策记录
-
-1. 不在 SGLang 中用独立 Triton SiTU 对 CUTLASS 的错误输出做事后修补；CUTLASS 已有独立 activation kernel，正确位置在 FlashInfer。
-2. 不复用 GPT-OSS 的奇偶行权重转换；Kimi loader 已明确给出连续 `[gate; up]` 半区。
-3. 不新增 FlashInfer ABI 参数；现有两组 per-expert float 参数足以承载 SiTU 的 `4/25`，但字段命名需要在代码和测试中解释清楚。
-4. 不把旧 `flashinfer_cutedsl` 实验当作本路径证据。
-5. 不进入整模型性能测试，直到 Phase 2 的 BF16/PyTorch 或 Marlin 参考正确性通过。
-
-## Phase 2 correctness harness
-
-状态：**完成**。本阶段建立验证框架并校准 reference，但 SiTU candidate 按预期尚未通过；未产生性能晋级结论。
-
-唯一入口：
-
-```text
-validate_kimi_sm120_mxfp4.py
-```
-
-### 对称性约束
-
- reference 与 candidate 共用同一 `Case`、随机种子、routing IDs/weights、量化后的 MXFP8 input、packed MXFP4 weights 和 E8M0 scales；reference 不读取量化前的 BF16 权重。
- FC1 在生成时显式采用 CUTLASS oracle 布局 `[up; gate]`。这使 harness 同时约束 SGLang 后续必须把 Kimi checkpoint 的 `[gate; up]` 正确交换。
- SwiGLU 作为当前 kernel 已支持的控制组；SiTU `(4, 25)` 使用完全相同的数据与调用，只改变 activation contract。
- correctness 固定 `use_fused_finalize=False`，避免非结合 atomic reduction 干扰确定性；benchmark 也由同一个 CUDA Event 计时函数驱动，但本阶段设置 iterations=0，不在共享 GPU 负载下生成性能数字。
-
-覆盖矩阵：
-
-| 类别 | 覆盖 |
-|---|---|
-| token shape | `M=1/2/8/32` |
-| routing | uniform、skewed、空 expert、`-1` masked slot |
-| EP | `EP2` global/local ID 过滤；Kimi `EP32/rank7/E_local=28` |
-| 数值 | finite、NaN/Inf row-level propagation、三次 bitwise determinism |
-| execution | eager、CUDA Graph capture、两次 replay |
-| Kimi 真尺寸 | `H=3584, N=3072, E_local=28, topK=16`；decode `M=1`、prefill `M=128` |
-
-### 执行命令
-
-```bash
-cd /data/hzy/sskj/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4
-
-# 已支持的 SwiGLU 控制组；本阶段最终证据统一使用 GPU 6
-CUDA_VISIBLE_DEVICES=6 PYTHONPATH=/tmp/hzy-fi-main-phase1 \
-FLASHINFER_DISABLE_VERSION_CHECK=1 \
-/data/hzy/envs/sparse-opd-sm120/bin/python validate_kimi_sm120_mxfp4.py \
-  --suite smoke --activation swiglu \
-  --output results/phase2_control_swiglu_smoke.json
-
-# Kimi SiTU 目标真尺寸
-CUDA_VISIBLE_DEVICES=6 PYTHONPATH=/tmp/hzy-fi-main-phase1 \
-FLASHINFER_DISABLE_VERSION_CHECK=1 \
-/data/hzy/envs/sparse-opd-sm120/bin/python validate_kimi_sm120_mxfp4.py \
-  --suite real --activation situ \
-  --output results/phase2_target_situ_real_gpu6.json
-```
-
-实现完成后的 CI-style gate 必须增加 `--require-situ`；若仍命中已知缺口，入口返回非零。
-
-### 结果与证据
-
-| 组别 | 结果 | 关键证据 |
-|---|---|---|
-| SwiGLU smoke control | 6/6 PASS | cosine `0.999625–0.999812`，NRMSE `1.50%–2.75%` |
-| SwiGLU Kimi real shape | 2/2 PASS | decode/prefill cosine `0.999597/0.999621`，NRMSE `2.85%/2.75%` |
-| SiTU smoke target | 6/6 EXPECTED_MISSING_SITU | 全部命中 binding `256 vs. 128` 首错 |
-| SiTU Kimi real shape | 2/2 EXPECTED_MISSING_SITU | GPU6 上 decode/prefill 均命中同一首错 |
-| determinism | PASS | non-fused finalize 三次 bitwise equal |
-| CUDA Graph | PASS | control replay 与 eager `max_abs=0`，两次 replay bitwise equal |
-
-原始证据：
-
-```text
-/data/hzy/sskj/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/phase2_control_swiglu_smoke_gpu6.log
-/data/hzy/sskj/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/phase2_target_situ_smoke_gpu6.log
-/data/hzy/sskj/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/phase2_control_swiglu_real_gpu6.log
-/data/hzy/sskj/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/phase2_target_situ_real_gpu6.log
-/data/hzy/sskj/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/results/phase2_*.json
-```
-
-### NaN/Inf 决策记录
-
-2026-08-17 的统一 gate 发现旧 oracle 不成立：原始 BF16 NaN/Inf 经
-`mxfp8_quantize` 后，E4M3 payload 全部有限，scale 最大值为 `0xFE`。
-NVIDIA E8M0 定义中 `0x00..0xFE` 表示有限的二次幂，只有 `0xFF` 表示
-NaN；转换采用 saturating finite 行为。旧 reference 把 `0xFE` scale
-重新展开到 BF16 时发生溢出，再要求 Tensor Core candidate 传播相同
-non-finite row，不是有效的 MXFP8 oracle。
-
-修正后的 nonfinite case 同时要求：
-
-1. 量化后的 E4M3 payload 全部有限；
-2. E8M0 scale 不得出现 `0xFF`；
-3. candidate 输出全部有限且三次 bitwise deterministic；
-4. BF16 dequant reference 的 overflow rows 继续记录为 diagnostic，不参与数值判定。
-
-这项修正没有改变任何 finite case 的 reference、容差或性能计时。有限值仍
-使用同一量化输入、权重、scale 与 routing，并满足 `rtol=0.2, atol=0.2`。
-
-官方格式依据：
-`https://docs.nvidia.com/cuda/cuda-math-api/cuda_math_api/struct____nv__fp8__e8m0.html`。
-
-## Phase 3 最小实现
-
-状态：**完成；有限值正确性与定义明确的 MXFP8 nonfinite 边界通过，可以进入 Phase 4 profiling。尚未形成整模型性能结论。**
-
-### 源码提交
-
-| 仓库 | 分支 | 提交 |
-|---|---|---|
-| FlashInfer | `hzy/kimi-k3-sm120-situ-mxfp4` | `7d1a4ee` `[Feature] Add CUTLASS SiTU fused MoE support on SM120` |
-| SGLang | `hzy/kimi-k3-sm120-flashinfer-mxfp4` | `fb929bb` `[Feature] Support Kimi-K3 FlashInfer MXFP4 MoE on SM120` |
-
-FlashInfer 的最小增量：
-
-1. C++ `ActivationType` 增加与 Python ABI 对齐的 `Situ`；
-2. 把 SiTU 标为 gated activation；
-3. 增加 `4*tanh(gate/4)*sigmoid(gate) * 25*tanh(up/25)` adaptor，复用已有 grouped GEMM 和 per-expert 参数；
-4. Python 暴露显式 capability marker，避免仅凭早已存在的 enum 误判 kernel 可用。
-
-SGLang 的最小增量：
-
-1. 同时支持 GPT-OSS 交错 `[gate_i, up_i]` 与 Kimi 连续 `[gate; up]` checkpoint 布局，统一生成 FlashInfer `[up; gate]`；
-2. 从 `MoeRunnerConfig` 透传 `Situ`、alpha `4`、beta `25`，不再硬编码 SwiGLU；
-3. bias-free Kimi 路径不向 kernel 传零占位 bias；
-4. FlashInfer 缺少 SiTU capability marker 时在 runner 初始化阶段明确报错，并提示以 `--moe-runner-backend marlin` 重启；不做 kernel 失败后的热切换。
-
-### 构建产物
-
-```text
-/data/hzy/artifacts/flashinfer-kimi-sm120-situ-phase3-v3/flashinfer_python-0.6.18-py3-none-any.whl
-SHA256 9298d34b6f9cec9b8bf535d1688fc6e8c73f7db65f64e0c496580e881c3d997b
-```
-
-### 正确性与集成结果
-
-统一 Phase 2/3 gate 命令：
-
-```bash
-cd /data/hzy/sskj/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4
-CUDA_VISIBLE_DEVICES=6 \
-PYTHONPATH=/tmp/hzy-fi-situ-phase3-v3 \
-FLASHINFER_DISABLE_VERSION_CHECK=1 \
-/data/hzy/envs/sparse-opd-sm120/bin/python validate_kimi_sm120_mxfp4.py \
-  --suite all --activation situ --require-situ \
-  --warmup 2 --benchmark-iterations 3 \
-  --output results/phase3_target_situ_all_gpu6_v3.json
-```
-
-结果：`8/8 PASS`。
-
-| 覆盖 | 结果 |
-|---|---|
-| M=1/2/8/32，uniform/skewed/empty/masked | PASS |
-| EP2 global/local IDs | PASS |
-| MXFP8 NaN/Inf saturation boundary | PASS，candidate finite、3-run bitwise |
-| CUDA Graph capture/replay | PASS，replay 与 eager `max_abs=0` |
-| Kimi decode `M=1, H=3584, N=3072, E=28, topK=16` | cosine `0.9996223`，NRMSE `2.748%` |
-| Kimi prefill `M=128` 同真尺寸 | cosine `0.9996152`，NRMSE `2.774%` |
-
-SGLang 集成测试：`6 passed, 18 warnings in 563.41s`。warnings 为 asyncio
-配置、PyTorch deprecation 和只读 pytest cache，不含断言失败。
-
-原始证据：
-
-```text
-/data/hzy/sskj/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/results/phase3_target_situ_all_gpu6_v3.log
-/data/hzy/sskj/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/results/phase3_target_situ_all_gpu6_v3.json
-/data/hzy/sskj/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/results/phase3_sglang_integration_gpu6_final.log
-```
-
-### Phase 3 决策
-
-1. 当前缺口同时位于 FlashInfer SiTU activation 与 SGLang Kimi layout/参数集成；两个补丁都必要。
-2. 没有修改 checkpoint，没有 runtime monkeypatch，没有通过放宽 finite tolerance 获得通过。
-3. microbenchmark latency 仅用于检查 candidate 可执行，不作为生产性能结论。
-4. 下一阶段必须用 Nsight Systems/Compute 和三轮交错 A/B 区分 wrapper、同步、权重搬运、dispatch 与 kernel 本体；在此之前不宣称优于 Marlin。
-
-## Phase 4 算子性能与 Profiler 归因
-
-状态：**完成。** 本阶段只证明单卡本地 MoE runner 的正确性、性能边界和
-kernel 归因，不代表完整 Kimi-K3 serving 性能。真实 Scheduler、Chunked
-Prefill、跨节点 EP 通信和 CUDA Graph 的联合效果留到 Phase 5。
-
-### 实验边界
-
-两后端均通过 SGLang 生产 runner 调用，而不是直接调用底层 kernel：
-
- baseline：`fused_experts_none_to_marlin`；
- candidate：`fused_experts_none_to_flashinfer_mxfp4`；
- 相同 BF16 输入、packed MXFP4 权重、E8M0 scale、Top-K IDs/weights、
-  SiTU `(alpha=4, beta=25)`；
- Kimi 真尺寸 `H=3584, N=3072, topK=16`；
- `EP32/rank7`，每个 rank 驻留 28 个本地专家，总专家数 896；
- `ep_mixed` 路由中每个 token 的 8 个 slot 指向本地专家、8 个指向远端专家；
- 仅使用 601 的物理 GPU6，并用 UUID
-  `GPU-45354dff-e8b3-1c0d-99c9-d3562bcab5fd` 做运行时绑定校验。
-
-这里的 `M` 是单次本地 MoE runner 收到的 token 数，不等于 HTTP 层的
-`ISL * concurrency`。本地平均路由项约为：
-
-```text
-local assignments / expert = M * 8 / 28
-```
-
-因此本阶段不能替代 `ISL=8192, C=16` 的真实服务测试。
-
-### 统一入口
-
-```bash
-python3 profile_kimi_sm120_mxfp4.py \
-  --shapes prefill prefill_m512 prefill_m2048 prefill_m8192 \
-  --rounds 5 --iterations 30 --warmup 5 \
-  --output results/phase4_mscan_ep32_5rounds_gpu6.json
-```
-
-正式运行时使用镜像 `lmsysorg/sglang:kimi-k3-fiv617situ-warm`，挂载：
-
-```text
-/tmp/hzy-fi-situ-phase3-v3             -> patched FlashInfer Python/AOT package
-/data/hzy/src/sglang-kimi-sm120        -> patched SGLang source
-/data/hzy/cache/sglang-jit-sm120       -> persistent SGLang JIT cache
-```
-
-并设置：
-
-```text
-FLASHINFER_DISABLE_JIT=1
-FLASHINFER_DISABLE_VERSION_CHECK=1
-SGLANG_JIT_CACHE_DIR=/root/.cache/sglang/jit
-EXPECTED_GPU_UUID=45354dff-e8b3-1c0d-99c9-d3562bcab5fd
-```
-
-### 五轮交错 A/B
-
-每个 shape 先做 BF16 参考正确性，再 warmup 5 次。五轮按
-`Marlin -> FlashInfer`、`FlashInfer -> Marlin` 交替执行，每轮每后端 30 次。
-下表为 GPU latency 中位数：
-
-| M | 平均本地路由项/专家 | Marlin | FlashInfer | FlashInfer 相对速度 | FI vs reference |
-|---:|---:|---:|---:|---:|---:|
-| 128 | 36.6 | 0.870 ms | 0.926 ms | 0.94x，慢约 6.5% | cosine 0.999615，NRMSE 2.774% |
-| 512 | 146.3 | 2.610 ms | 1.782 ms | 1.46x | cosine 0.999614，NRMSE 2.779% |
-| 2048 | 585.1 | 9.070 ms | 4.355 ms | 2.08x | cosine 0.999623，NRMSE 2.746% |
-| 8192 | 2340.6 | 34.379 ms | 16.167 ms | 2.13x | cosine 0.999623，NRMSE 2.744% |
-
-五轮逐轮比值稳定：M512 为 `1.428x-1.474x`，M2048 为
-`2.065x-2.089x`，M8192 为 `2.126x-2.128x`。这证明 FlashInfer 的优势
-依赖每次 runner 的实际 M；它不是小 batch 的无条件替代品。
-
-M512 的第一次 pilot 只 warmup 1 次，得到 FlashInfer 4.26 ms、Marlin
-1.94 ms，与正式结果相反。该结果保留为 warmup 敏感性证据，不用于性能结论。
-
-### Nsight Systems
-
-捕获文件：
-
-```text
-results/phase4_nsys_ab_gpu6.nsys-rep
-results/phase4_nsys_ab_gpu6.log
-results/phase4_nsys_stats_*.csv
-results/phase4_nsys_*_iteration_5_*.csv
-```
-
-稳定的第 5 次迭代中，主 GEMM kernel 总时间：
-
-| Shape | Marlin | FlashInfer | 观察 |
-|---|---:|---:|---|
-| Decode M1 | 109.4 us | 191.8 us | FlashInfer 主 kernel 约慢 75% |
-| Prefill M128 | 709.9 us | 865.3 us | FlashInfer 主 kernel 约慢 22% |
-
-Decode 的非主 kernel 开销约为 FlashInfer 20.3 us、Marlin 6.4 us；Prefill
-分别约 45.8 us、44.2 us。小 M 差距主要来自 kernel 本体，不能归因成纯
-Python wrapper 或同步开销。
-
-### Nsight Compute
-
-RTX 6000D 需要较新的工具。镜像内 NCU 2025.3.1 返回
-`Profiling is not supported on device 0`；正式报告由 NVIDIA 官方包
-Nsight Compute 2026.2.1 生成：
-
-```text
-/data/hzy/tools/nsight-compute-2026.2.1/
-results/phase4_ncu_2026_2_1_ab_gpu6.ncu-rep
-results/phase4_ncu_2026_2_1_ab_gpu6.log
-results/phase4_ncu_2026_2_1_ab_gpu6_raw.csv
-```
-
-`--set basic` 捕获到的两组主 kernel 指标：
-
-| Shape/backend | 主 kernel 合计 | DRAM 峰值利用率 | SM 峰值利用率 | registers/thread | shared memory/block |
-|---|---:|---:|---:|---:|---:|
-| Decode Marlin | 121.2 us | 77.8%-86.8% | 21.2%-22.5% | 113 | 33.8 KiB |
-| Decode FlashInfer | 199.4 us | 49.9%-51.7% | 47.7%-47.8% | 168 | 84.0 KiB |
-| M128 Marlin | 513.0 us | 73.7%-76.1% | 61.1%-65.9% | 255 | 102.4 KiB |
-| M128 FlashInfer | 535.1 us | 67.2%-70.7% | 61.1%-63.3% | 168 | 84.0 KiB |
-
-NCU 使用 replay，绝对时间不与 Nsight Systems 混用。它支持相同方向的
-结论：Decode 时 FlashInfer 使用更多寄存器和共享内存，DRAM 利用率更低，
-当前 CUTLASS 配置没有在小 M 上胜过 Marlin；M128 的主 kernel 已接近，
-但完整 runner 仍有其他步骤。
-
-### 工程问题与修复记录
-
-1. `--privileged` 会重新暴露全部 `/dev/nvidia*`，曾令容器落到物理 GPU0。
-   UUID guard 立即终止实验。之后固定使用 `--gpus device=6`、
-   `NVIDIA_VISIBLE_DEVICES=6`，不再使用 privileged。
-2. 纯 Python wheel 在每个新进程触发 FlashInfer JIT；NCU 注入子进程后会
-   卡住大量 `nvcc`。正式 profiling 把三个已验证 `.so` 放入
-   `flashinfer/data/aot/`，并设置 `FLASHINFER_DISABLE_JIT=1`。
-3. SGLang 自身 JIT cache 持久化到 `/data/hzy/cache/sglang-jit-sm120`。
-4. `torch.quantile` 对 M8192 的约 2936 万元素输出报输入过大。统计代码改用
-   精确 `torch.kthvalue`；cosine、NRMSE、max/mean 与正确性阈值均未放宽。
-
-详细时间线和失败证据见 `decision_log.md`。
-
-### Phase 4 决策
-
-1. FlashInfer SiTU MXFP4 路径已通过正确性门槛，并在 M512 及以上的本地
-   runner shape 显示明确吞吐潜力。
-2. 它在 Decode/M128 小 M 上仍落后于 Marlin，不能现在就设为无条件默认。
-3. 不再扩展 synthetic EP4/kernel sweep。下一步应进入四节点真实 Kimi-K3
-   serving，观察 Scheduler 实际 M 分布及端到端 TTFT/TPOT/TPS。
-4. 在 Phase 5 完成前，不宣称 FlashInfer 提升生产性能；当前 PR 的可靠表述
-   是补齐 Kimi-K3 + SM120 + SiTU + MXFP4 功能支持，并提供显式 Marlin fallback。
-
-### Phase 4 原始证据
-
-服务器根目录：
-
-```text
-/data/hzy/sskj/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/results/
-```
-
-关键文件：
-
-```text
-phase4_ab_5rounds_gpu6.json
-phase4_mscan_ep32_5rounds_gpu6.json
-phase4_mscan_ep32_5rounds_gpu6.log
-phase4_aot_nojit_smoke_gpu6.json
-phase4_nsys_ab_gpu6.nsys-rep
-phase4_nsys_stats_*.csv
-phase4_ncu_2026_2_1_ab_gpu6.ncu-rep
-phase4_ncu_2026_2_1_ab_gpu6_raw.csv
-```
-
-本地备份：
-
-```text
-/Users/hzy/Desktop/infra/.phase3_work/phase4_results/
-```
-
-## Phase 5 四节点真实 Serving
-
-状态：**完成。** 本阶段在 601-604 共 32 张 RTX PRO 6000D 上验证
-TP32/EP4 Kimi-K3，使用真实 SGLang Scheduler、Chunked Prefill、CUDA
-Graph、四节点 NCCL/RDMA 和 HTTP benchmark。
-
-### 正式矩阵
-
-```text
-ISL=16384, OSL=1, requests=40
-TP32, EP4, DP1
-backend=marlin / flashinfer_mxfp4
-chunked_prefill_size=8192 / 16384
-concurrency=8 / 16
-3 repeats per point
-```
-
-Prefix/Radix cache 已禁用；client 使用 `random-ids` 保证精确输入长度，且
-tokenizer 从本地模型目录读取。24/24 repeat 均完成 40/40 请求，errors 为空。
-
-### 三次重复中位数
-
-| Chunk | C | Backend | Input TPS | TTFT P50 | TTFT P95 |
-|---:|---:|---|---:|---:|---:|
-| 8K | 8 | Marlin | 2935.02 | 43.51 s | 46.22 s |
-| 8K | 8 | FlashInfer | 3257.96 | 39.19 s | 41.62 s |
-| 8K | 16 | Marlin | 2934.90 | 87.09 s | 89.86 s |
-| 8K | 16 | FlashInfer | 3260.14 | 78.39 s | 80.88 s |
-| 16K | 8 | Marlin | 2812.91 | 45.47 s | 45.86 s |
-| 16K | 8 | FlashInfer | 3027.50 | 42.27 s | 42.63 s |
-| 16K | 16 | Marlin | 2812.33 | 90.97 s | 91.36 s |
-| 16K | 16 | FlashInfer | 3027.34 | 84.52 s | 84.85 s |
-
-FlashInfer 在四个场景均胜过 Marlin：8K chunk 的 Input TPS 提升约 11%，
-TTFT P50 降低约 10%；16K chunk 的收益约为 +7.6% / -7.1%。同一后端内，
-8K chunk 也优于 16K。因此本工作负载的推荐组合是：
-
-```text
-EP4 + flashinfer_mxfp4 + chunked_prefill_size=8192
-```
-
-已有 EP32 Marlin/8K/C8 两次结果仅作同口径对照：EP4 相对 EP32 的 Input
-TPS +15.94%，TTFT P50 -13.88%。按要求没有补跑 EP32。
-
-### 生产路径修复
-
-真实服务额外发现并修复：旧 Kimi 镜像缺少通用 `gemm1_beta` 字段、
-Prefill hidden states 可能非连续、benchmark tokenizer 必须固定本地路径。
-这些问题均在正式矩阵前复现，并由最终 24 点结果覆盖。
-
-### 证据
-
-服务器原始目录：
-
-```text
-/data/hzy/sskj/experiments/pro6000/kimi3_pro6000_sglang_tp32ep32_moe_backend_prefill/results/kimi3-ep4-moe-full-20260818-151349/
-/data/hzy/kimi3-ep4-moe-full-20260818-151349.log
-```
-
-仓库内汇总：
-
-```text
-results/phase5_ep4_moe_backend_prefill_summary.json
-results/phase5_ep4_moe_backend_prefill_results.csv
-```
-
-日志扫描未发现 OOM、Traceback、NCCL error 或 EngineDeadError。实验结束后
-601-604 的容器全部删除，32 张 GPU 均回到 0 MiB / 0% utilization。
-
-### PR 级回归
-
-真实 serving 发现 `mxfp8_quantize` 要求连续输入，因此 SGLang runner 在量化
-边界显式 materialize contiguous tensor，并把原 Kimi test 改为构造非连续
-输入。GPU6 定向测试结果：
-
-```text
-1 passed, 5 deselected in 20.22s
-```
-
-FlashInfer 官方 `tests/moe/test_trtllm_cutlass_fused_moe.py` 增加 SiTU 的
-MXFP8×MXFP4 参数组，覆盖 FP16/BF16 与 autotune on/off：
-
-```text
-4 passed, 153 deselected in 8.50s
-```
-
-候选提交：
-
-| 仓库 | 提交 | 内容 |
-|---|---|---|
-| SGLang | `8a2d501` | 非连续输入修复与回归 |
-| FlashInfer | `e0bc653` | 官方 SiTU MXFP8×MXFP4 测试 |
-
-服务器日志：
-
-```text
-/data/hzy/phase5_sglang_noncontiguous_gpu6_v2.log
-/data/hzy/phase5_flashinfer_official_situ_gpu6.log
-```
-
-去重后又在 FlashInfer #4460 的正式参数接口上完成最终组合验证：从 #4460
-源码构建 `flashinfer-python 0.6.18` wheel，加载当前 SGLang Draft，结果为：
+使用 FlashInfer #4460 构建的 wheel，在 RTX 6000D GPU6 上执行精确 Kimi 组合：

 ```text
 8 passed, 17 warnings in 559.87s
-/data/hzy/phase5_sglang_pr4460_draft_gpu6_v4.log
 ```

-该 FlashInfer 测试提交仅保留为本地验证，不提交 FlashInfer PR。SGLang Draft
-不 vendor kernel、不 pin 未合并 commit、也不改变默认 MoE backend。
+覆盖 `SM120 + MXFP8 + MXFP4 + SiTU(4,25)`、Kimi 布局、非连续输入和
+CUDA Graph。Kimi 真尺寸结果：

-### 最终 PR 代表点
+| Case | cosine | NRMSE |
+|---|---:|---:|
+| Decode M=1 | 0.9996223 | 2.748% |
+| Prefill M=128 | 0.9996152 | 2.774% |

-为避免把旧候选镜像结果直接当作当前 Draft 的性能证据，最终又使用完全相同的
-SGLang `300c87a` Python 源码和 FlashInfer #4460 实现 `b525c513`，对唯一代表点
-重新执行对称 A/B：
+最终 Draft HEAD 的 compileall 和非 JIT 回归：
+
+```text
+10 passed, 3 deselected, 17 warnings in 33.30s
+```
+
+五个变更文件已通过完整 pre-commit。同步 fork 后的 7 个提交与同步前版本在
+`git range-diff` 中全部 patch-equivalent。
+
+原始日志：
+
+```text
+/data/hzy/phase5_sglang_pr4460_draft_gpu6_v4.log
+/data/hzy/phase5_sglang_pr4460_draft_cpu_final.log
+```
+
+## 四节点真实 Serving
+
+### 实验口径
+
+| 维度 | 取值 |
+|---|---|
+| 节点 | 601-604，共 32 张 RTX 6000D |
+| 并行 | TP32 / EP4 / DP1 |
+| 输入/输出 | 16K / 1 token |
+| 并发 | C=8、C=16 |
+| Backend | Marlin、FlashInfer MXFP4 |
+| Chunk | 4K、8K、16K |
+| 请求与重复 | 每点 40 请求，3 次重复 |
+| Cache | 禁用 Radix Cache |
+
+正式 Run：
+
+```text
+8K/16K: kimi3-ep4-moe-full-20260818-151349
+4K:     kimi3-ep4-moe-chunk4k-20260819-150632
+```
+
+36/36 repeat 均完成 40/40 请求，输出长度均为 1，`errors` 为空；日志未发现
+OOM、Traceback、NCCL error 或 EngineDeadError。下表为三次重复的中位数。
+
+### 实验结果
+
+| Chunk | C | Backend | Input TPS | TTFT P50 | TTFT P95 | FlashInfer vs Marlin |
+|---:|---:|---|---:|---:|---:|---:|
+| 4K | 8 | Marlin | 2694.53 | 48.86 s | 49.20 s | baseline |
+| 4K | 8 | FlashInfer | **2957.23** | **44.48 s** | **44.86 s** | +9.75% TPS / -8.96% TTFT |
+| 4K | 16 | Marlin | 2692.60 | 96.25 s | 97.96 s | baseline |
+| 4K | 16 | FlashInfer | **2960.41** | **87.61 s** | **89.04 s** | +9.95% TPS / -8.98% TTFT |
+| 8K | 8 | Marlin | 2935.02 | 43.51 s | 46.22 s | baseline |
+| 8K | 8 | FlashInfer | **3257.96** | **39.19 s** | **41.62 s** | +11.00% TPS / -9.92% TTFT |
+| 8K | 16 | Marlin | 2934.90 | 87.09 s | 89.86 s | baseline |
+| 8K | 16 | FlashInfer | **3260.14** | **78.39 s** | **80.88 s** | +11.08% TPS / -9.99% TTFT |
+| 16K | 8 | Marlin | 2812.91 | 45.47 s | 45.86 s | baseline |
+| 16K | 8 | FlashInfer | **3027.50** | **42.27 s** | **42.63 s** | +7.63% TPS / -7.04% TTFT |
+| 16K | 16 | Marlin | 2812.33 | 90.97 s | 91.36 s | baseline |
+| 16K | 16 | FlashInfer | **3027.34** | **84.52 s** | **84.85 s** | +7.65% TPS / -7.09% TTFT |
+
+### 结果结论
+
+1. FlashInfer 在 6 个正式 shape 上均优于 Marlin；
+2. 4K Chunk 的收益约为 `+9.8% Input TPS / -9.0% TTFT`；
+3. 8K Chunk 的收益约为 `+11% Input TPS / -10% TTFT`；
+4. 16K Chunk 的收益约为 `+7.6% Input TPS / -7.1% TTFT`；
+5. 同一 backend 下，8K Chunk 同时优于 4K 和 16K；4K 切分轮次过多，16K 单步过重；
+6. C 从 8 增至 16 后吞吐基本不变、TTFT 约翻倍，C=8 已接近 Prefill 吞吐饱和；
+7. 当前推荐为 `EP4 + flashinfer_mxfp4 + chunked_prefill_size=8192`。
+
+该优化有效但不是数量级变化。最佳 TTFT P50 仍约 39 秒，后续需要继续优化
+TP/EP 通信与 Prefill 并行结构。
+
+### Chunk=4K 补测结论
+
+4K 补测复用了 8K/16K 矩阵的镜像、模型、服务参数和入口。12/12 repeat
+全部有效，共 480/480 请求成功。
+
+FlashInfer 在 C=8/16 下仍分别提升 `9.75%/9.95%` Input TPS，但 4K 的绝对
+Input TPS 比 8K 低约 9%，TTFT 也更高。该结果补齐了 Chunk 曲线并确认 8K
+不是网格边界造成的偶然最优点。
+
+## 最终 PR 代表点
+
+最终又用等价 SGLang patch `300c87a` 和 FlashInfer #4460 implementation
+`b525c513` 重跑 8K/C8 代表点：

 ```text
 Run ID: kimi3-sm120-pr-representative-20260819-115530
@ -728,20 +185,97 @@ ISL=16384 / OSL=1 / C=8 / requests=40 / chunk=8192
 Marlin 3 repeats + FlashInfer MXFP4 3 repeats
 ```

-| Backend | Input TPS 中位数 | TTFT P50 中位数 | TTFT P95 中位数 |
+| Backend | Input TPS | TTFT P50 | TTFT P95 |
 |---|---:|---:|---:|
 | Marlin | 2930.03 | 43.587 s | 46.287 s |
 | FlashInfer MXFP4 | 3251.25 | 39.267 s | 41.694 s |
 | 相对变化 | +10.96% | -9.91% | -9.92% |

-六轮均为 40/40 成功，`errors` 中没有非空项，所有 `output_lens` 均为 1；服务
-和 benchmark 日志未命中 OOM、Traceback、NCCL error 或 EngineDeadError。
-FlashInfer 在 32 个 rank 上完成 autotune，未复现 605 补丁中的 skip-autotune
-需求。运行结束后 601-604 的 GPU compute process 均为 0。
+六轮均为 40/40 成功，所有 `output_lens=1`，无 OOM、Traceback、NCCL error
+或 EngineDeadError。

-原始证据位于 601：
+## 复现命令
+
+四节点矩阵只在 601 执行：
+
+```bash
+cd /data/hzy/sskj/experiments/pro6000/kimi3_pro6000_sglang_tp32ep32_moe_backend_prefill
+
+read -rsp 'sudo password: ' SUDO_PASSWORD; echo
+export SUDO_PASSWORD
+export RUN_ID=kimi3-moe-prefill-$(date +%Y%m%d-%H%M%S)
+
+bash run_moe_backend_prefill.sh run
+```
+
+可通过以下环境变量缩小矩阵：
+
+```text
+EP_SIZE
+BACKENDS_CSV
+CHUNKED_PREFILL_SIZES_CSV
+CONCURRENCIES_CSV
+REPEATS
+```
+
+4K 补测使用：
+
+```bash
+EP_SIZE=4 \
+BACKENDS_CSV=marlin,flashinfer_mxfp4 \
+CHUNKED_PREFILL_SIZES_CSV=4096 \
+CONCURRENCIES_CSV=8,16 \
+REPEATS=3 \
+bash run_moe_backend_prefill.sh run
+```
+
+## 证据路径
+
+### 8K/16K 正式矩阵
+
+```text
+/data/hzy/sskj/experiments/pro6000/kimi3_pro6000_sglang_tp32ep32_moe_backend_prefill/results/kimi3-ep4-moe-full-20260818-151349/
+/data/hzy/kimi3-ep4-moe-full-20260818-151349.log
+
+results/phase5_ep4_moe_backend_prefill_summary.json
+results/phase5_ep4_moe_backend_prefill_results.csv
+```
+
+### 4K 补测
+
+```text
+/data/hzy/sskj/experiments/pro6000/kimi3_pro6000_sglang_tp32ep32_moe_backend_prefill/results/kimi3-ep4-moe-chunk4k-20260819-150632/
+/data/hzy/kimi3-ep4-moe-chunk4k-20260819-150632.log
+
+results/phase5_ep4_moe_backend_chunk4k_summary.json
+results/phase5_ep4_moe_backend_chunk4k_results.csv
+```
+
+### 最终 PR 代表点

 ```text
 /data/hzy/sskj/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/results/kimi3-sm120-pr-representative-20260819-115530/
 /data/hzy/kimi3-sm120-pr-representative-20260819-115530.log
+
+results/pr_representative_16k_c8_chunk8k_summary.json
+results/pr_representative_16k_c8_chunk8k_results.csv
 ```
+
+### 单卡测试与 Profiler
+
+```text
+results/phase3_target_situ_all_gpu6_v3.json
+results/phase4_mscan_ep32_5rounds_gpu6.json
+results/phase4_nsys_ab_gpu6.nsys-rep
+results/phase4_ncu_2026_2_1_ab_gpu6.ncu-rep
+```
+
+## 后续方向
+
+MoE backend 选型完成后，不再扩大同类 backend 网格。后续 Prefill 优化按独立
+任务推进：
+
+1. TP Reduce Scatter / Attention TP input scattered；
+2. MoE A2A；
+3. Prefill Pipeline Parallelism；
+4. 对最优新配置重新做 Nsight 时间线验证。
--- a/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/results/phase5_ep4_moe_backend_chunk4k_results.csv
+++ b/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/results/phase5_ep4_moe_backend_chunk4k_results.csv
@ -0,0 +1,13 @@
+backend,chunked_prefill_size,concurrency,repeat,completed,duration,request_throughput,input_throughput,output_throughput,total_throughput,mean_ttft_ms,median_ttft_ms,p95_ttft_ms,p99_ttft_ms,mean_e2e_latency_ms,median_e2e_latency_ms,p95_e2e_latency_ms,p99_e2e_latency_ms
+flashinfer_mxfp4,4096,16,1,40,221.47698600799777,0.18060567249436724,2959.043338147713,0.18060567249436724,2959.2239438202073,72603.08281157631,87635.37671801168,89080.70434568217,89951.34060245706,72603.12579314923,87635.42024500202,89080.75052604545,89951.384754295
+flashinfer_mxfp4,4096,16,2,40,221.06144716404378,0.1809451648541732,2964.6055809707736,0.1809451648541732,2964.786526135628,72440.78481566976,87606.67828551959,88998.53118621977,89664.52184239868,72440.8288814302,87606.71954799909,88998.57463886729,89664.56696634414
+flashinfer_mxfp4,4096,16,3,40,221.3745526280254,0.18068924149205084,2960.412532605761,0.18068924149205084,2960.593221847253,72566.08219580376,87594.9859394459,89038.73863747576,89889.50585184386,72566.12714612857,87595.03182361368,89038.78330618609,89889.55271399347
+flashinfer_mxfp4,4096,8,1,40,221.64001683914103,0.1804728251262978,2956.8667668692633,0.1804728251262978,2957.0472396943896,40754.60876591387,44480.76946451329,44945.41641417891,45654.79073660448,40754.653389629675,44480.81373295281,44945.46504281461,45654.83455775771
+flashinfer_mxfp4,4096,8,2,40,221.6125555918552,0.18049518852022162,2957.233168715311,0.18049518852022162,2957.4136639038315,40748.24627056951,44506.1850650236,44862.61514539365,45574.393972705584,40748.29224284622,44506.22835196555,44862.659902649466,45574.439538051374
+flashinfer_mxfp4,4096,8,3,40,221.44526491500437,0.1806315434893263,2959.467208529122,0.1806315434893263,2959.647840072612,40712.37499640556,44453.15464294981,44811.94024076685,45526.09015815426,40712.42369815009,44453.20301700849,44811.989083641674,45526.13830527756
+marlin,4096,16,1,40,243.22425961703993,0.16445727931490292,2694.4680642953695,0.16445727931490292,2694.6325215746842,79749.63234093157,96212.83750247676,97888.6785025592,98836.45602966193,79749.6743179392,96212.87480182946,97888.72231170535,98836.50125178043
+marlin,4096,16,2,40,243.39303705701604,0.1643432387534973,2692.5996237372997,0.1643432387534973,2692.763966976053,79811.2635479949,96251.52806309052,97955.36642154912,98938.37275844067,79811.30552649265,96251.56629201956,97955.41314255679,98938.41525124619
+marlin,4096,16,3,40,243.62333092605695,0.16418788729286585,2690.054345406314,0.16418788729286585,2690.218533293607,79882.47500148718,96372.09082697518,98021.89738572342,99003.82948873332,79882.51752386568,96372.13580147363,98021.94035288412,99003.87401685584
+marlin,4096,8,1,40,243.21893303701654,0.16446088098705797,2694.5270740919577,0.16446088098705797,2694.6915349729447,44725.21210900741,48850.54129792843,49209.063365869224,50020.41739182314,44725.258720549755,48850.58849258348,49209.11144018173,50020.46818676172
+marlin,4096,8,2,40,243.2142291299533,0.16446406175778208,2694.5791878395016,0.16446406175778208,2694.743651901259,44724.67264174484,48868.09575161897,49182.01875761151,50006.30366436206,44724.714723881334,48868.13376657665,49182.06452928716,50006.34664747398
+marlin,4096,8,3,40,243.25217532995157,0.1644384061344705,2694.1588461071647,0.1644384061344705,2694.323284513299,44730.88195841992,48857.85997705534,49201.6794851399,50016.33869816782,44730.92947611003,48857.90269903373,49201.72802127199,50016.38562648789
--- a/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/results/phase5_ep4_moe_backend_chunk4k_summary.json
+++ b/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/results/phase5_ep4_moe_backend_chunk4k_summary.json
@ -0,0 +1,109 @@
+{
+  "run_id": "kimi3-ep4-moe-chunk4k-20260819-150632",
+  "summary": [
+    {
+      "backend": "flashinfer_mxfp4",
+      "chunked_prefill_size": 4096,
+      "concurrency": 8,
+      "repeats": 3,
+      "completed_each": [
+        40,
+        40,
+        40
+      ],
+      "median_request_throughput": 0.18049518852022162,
+      "median_input_throughput": 2957.233168715311,
+      "median_total_throughput": 2957.4136639038315,
+      "median_mean_ttft_ms": 40748.24627056951,
+      "median_median_ttft_ms": 44480.76946451329,
+      "median_p95_ttft_ms": 44862.61514539365,
+      "median_p99_ttft_ms": 45574.393972705584,
+      "median_mean_e2e_latency_ms": 40748.29224284622,
+      "median_median_e2e_latency_ms": 44480.81373295281,
+      "median_p95_e2e_latency_ms": 44862.659902649466
+    },
+    {
+      "backend": "flashinfer_mxfp4",
+      "chunked_prefill_size": 4096,
+      "concurrency": 16,
+      "repeats": 3,
+      "completed_each": [
+        40,
+        40,
+        40
+      ],
+      "median_request_throughput": 0.18068924149205084,
+      "median_input_throughput": 2960.412532605761,
+      "median_total_throughput": 2960.593221847253,
+      "median_mean_ttft_ms": 72566.08219580376,
+      "median_median_ttft_ms": 87606.67828551959,
+      "median_p95_ttft_ms": 89038.73863747576,
+      "median_p99_ttft_ms": 89889.50585184386,
+      "median_mean_e2e_latency_ms": 72566.12714612857,
+      "median_median_e2e_latency_ms": 87606.71954799909,
+      "median_p95_e2e_latency_ms": 89038.78330618609
+    },
+    {
+      "backend": "marlin",
+      "chunked_prefill_size": 4096,
+      "concurrency": 8,
+      "repeats": 3,
+      "completed_each": [
+        40,
+        40,
+        40
+      ],
+      "median_request_throughput": 0.16446088098705797,
+      "median_input_throughput": 2694.5270740919577,
+      "median_total_throughput": 2694.6915349729447,
+      "median_mean_ttft_ms": 44725.21210900741,
+      "median_median_ttft_ms": 48857.85997705534,
+      "median_p95_ttft_ms": 49201.6794851399,
+      "median_p99_ttft_ms": 50016.33869816782,
+      "median_mean_e2e_latency_ms": 44725.258720549755,
+      "median_median_e2e_latency_ms": 48857.90269903373,
+      "median_p95_e2e_latency_ms": 49201.72802127199
+    },
+    {
+      "backend": "marlin",
+      "chunked_prefill_size": 4096,
+      "concurrency": 16,
+      "repeats": 3,
+      "completed_each": [
+        40,
+        40,
+        40
+      ],
+      "median_request_throughput": 0.1643432387534973,
+      "median_input_throughput": 2692.5996237372997,
+      "median_total_throughput": 2692.763966976053,
+      "median_mean_ttft_ms": 79811.2635479949,
+      "median_median_ttft_ms": 96251.52806309052,
+      "median_p95_ttft_ms": 97955.36642154912,
+      "median_p99_ttft_ms": 98938.37275844067,
+      "median_mean_e2e_latency_ms": 79811.30552649265,
+      "median_median_e2e_latency_ms": 96251.56629201956,
+      "median_p95_e2e_latency_ms": 97955.41314255679
+    }
+  ],
+  "comparisons": [
+    {
+      "chunked_prefill_size": 4096,
+      "concurrency": 8,
+      "flashinfer_input_throughput_change_pct": 9.749617925508659,
+      "flashinfer_total_throughput_change_pct": 9.749617925508659,
+      "flashinfer_median_ttft_ms_change_pct": -8.958825692729121,
+      "flashinfer_p95_ttft_ms_change_pct": -8.818935420805618,
+      "flashinfer_median_e2e_latency_ms_change_pct": -8.958814693794636
+    },
+    {
+      "chunked_prefill_size": 4096,
+      "concurrency": 16,
+      "flashinfer_input_throughput_change_pct": 9.946258125697117,
+      "flashinfer_total_throughput_change_pct": 9.946258125697138,
+      "flashinfer_median_ttft_ms_change_pct": -8.981519516141546,
+      "flashinfer_p95_ttft_ms_change_pct": -9.102745576695416,
+      "flashinfer_median_e2e_latency_ms_change_pct": -8.981512797197189
+    }
+  ]
+}