diff --git a/README.md b/README.md index 6f674bb..24caad8 100644 --- a/README.md +++ b/README.md @@ -1,5 +1,7 @@ # sskj — 多平台大模型推理性能基准测试项目 +**更新（2026-08-19 13:07:48 CST）**：完成 Kimi-K3 SM120 SGLang Draft PR 的最终代表点验收。基于 SGLang `300c87a` 与 FlashInfer #4460 `b525c51`，在 601-604 四节点 TP32×EP4 上完成 16K→1、C=8、Chunk=8K 的 Marlin/FlashInfer 各 3 次 A/B；240/240 请求成功。FlashInfer 的 Input TPS 中位数提升 10.96%，TTFT P50/P95 分别下降 9.91%/9.92%。公开 PR Checklist 已全部勾选，轻量汇总和原始证据索引已归档。 + **更新（2026-08-19 10:45:18 CST）**：新增 Kimi-K3 SM120 SGLang Draft PR 的单点代表性验收入口。使用精确 SGLang `300c87a`、FlashInfer #4460 实现 `b525c51` 和统一镜像，在 601-604 四节点 TP32×EP4 上仅复测 16K→1、C=8、Chunk=8K 的 Marlin/FlashInfer A/B，各 3 次重复；完整命令、镜像身份、原始日志和汇总统一落入单个 Run 目录。 **更新（2026-08-18 23:00:42 CST）**：完成依赖 FlashInfer #4460 的 SGLang Draft PR 收敛。确认不提交任何 FlashInfer PR，也不 vendor 或 pin 未合并 kernel；在 601 GPU6 上从 #4460 源码构建 FlashInfer 0.6.18 wheel，SGLang 定向测试 8/8 通过。Draft 标题、依赖、兼容策略、测试和四机 EP4 数据见 `experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/SGLANG_DRAFT_PR.md`。 diff --git a/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/COMPLETION_AUDIT.md b/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/COMPLETION_AUDIT.md index c1e28f4..0f875cb 100644 --- a/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/COMPLETION_AUDIT.md +++ b/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/COMPLETION_AUDIT.md @@ -1,6 +1,6 @@ # Kimi-K3 SM120 FlashInfer MXFP4 完成审计 -审计时间：2026-08-18。本文只记录 goal 的逐项验收状态；技术结论见 +审计时间：2026-08-19。本文只记录 goal 的逐项验收状态；技术结论见 `UPSTREAM_DUPLICATION_AUDIT.md`，公开 PR 正文见 `SGLANG_DRAFT_PR.md`。 ## 验收矩阵 @@ -21,7 +21,8 @@ | 最终兼容回归 | 已通过 | Draft HEAD `300c87a`：`8 passed, 5 deselected, 17 warnings in 22.23s` | | 格式和 CI 注册 | 已通过 | 五个变更文件完整 pre-commit 全部通过；两个测试文件注册到 CUDA CI `base-b/1-gpu-small` | | 保留 EP4 服务级结果 | 已证明 | 601 原始目录有 24 个正式 JSONL 点；每组 3 repeats、每点 40/40 成功；错误扫描为空；未重跑 EP32 | -| 公开 PR 标题、正文、依赖和已知限制 | Draft 已完成 | `SGLANG_DRAFT_PR.md` 遵循 SGLang 官方模板；标题为 `[Feature] Support Kimi-K3 FlashInfer MXFP4 MoE on SM120`；公开 PR 不宣称提速，历史 EP4 性能结果仅保留在内部实验档案 | +| 当前 Draft 代表点 A/B | 已通过 | Run `kimi3-sm120-pr-representative-20260819-115530`；Marlin/FlashInfer 各 3 次，240/240 请求成功；FlashInfer Input TPS +10.96%，TTFT P50/P95 -9.91%/-9.92% | +| 公开 PR 标题、正文、依赖和已知限制 | Draft 正文已完成 | `SGLANG_DRAFT_PR.md` 遵循官方模板，Checklist 全部勾选；标题为 `[Feature] Support Kimi-K3 FlashInfer MXFP4 MoE on SM120`；Speed Tests 只报告一个目标硬件代表点，不作普遍加速声明 | | GitHub fork 交付 | 已完成 | `ACEEE-1222/sglang:hzy/kimi-k3-sm120-flashinfer-mxfp4`，Draft HEAD `300c87a` | | 内部仓库独立分支 | 本地和 601 已完成，远端待推送 | 601 分支 `hzy/feature-kimi-k3-sm120-flashinfer-mxfp4` 已到审计提交；`git.meta-stone.net` 尚未为本机或 601 key 授权，HTTPS 也无 Token | @@ -38,6 +39,10 @@ # 当前 SGLang Draft worktree（601） /data/hzy/src/sglang-kimi-sm120-draft-wt + +# 当前 Draft 代表点（601） +/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 ``` ## 唯一未关闭项 diff --git a/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/README.md b/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/README.md index 0252545..de8ea70 100644 --- a/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/README.md +++ b/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/README.md @@ -16,6 +16,13 @@ 逐项验收和唯一剩余交付项见 [`COMPLETION_AUDIT.md`](COMPLETION_AUDIT.md)。 +最终 PR 代表点已经在与 Draft HEAD 对齐的统一镜像中完成：四节点 +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`。 + ## 目标 为 Kimi-K3 在 NVIDIA SM120（RTX PRO 6000D）上的 SGLang 部署补齐 FlashInfer MXFP8 activation × MXFP4 weight fused MoE，并形成可回归、可性能比较、可提交上游的最小改动。 @@ -704,3 +711,34 @@ MXFP8×MXFP4 参数组，覆盖 FP16/BF16 与 autotune on/off： 该 FlashInfer 测试提交仅保留为本地验证，不提交 FlashInfer PR。SGLang Draft 不 vendor kernel、不 pin 未合并 commit、也不改变默认 MoE backend。 + +### 最终 PR 代表点 + +为避免把旧候选镜像结果直接当作当前 Draft 的性能证据，最终又使用完全相同的 +SGLang `300c87a` Python 源码和 FlashInfer #4460 实现 `b525c513`，对唯一代表点 +重新执行对称 A/B： + +```text +Run ID: kimi3-sm120-pr-representative-20260819-115530 +TP32 / EP4 / DP1 +ISL=16384 / OSL=1 / C=8 / requests=40 / chunk=8192 +Marlin 3 repeats + FlashInfer MXFP4 3 repeats +``` + +| 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。 + +原始证据位于 601： + +```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 +``` diff --git a/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/SGLANG_DRAFT_PR.md b/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/SGLANG_DRAFT_PR.md index 7084740..0524cd5 100644 --- a/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/SGLANG_DRAFT_PR.md +++ b/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/SGLANG_DRAFT_PR.md @@ -80,25 +80,45 @@ pytest -q -s \ ## Speed Tests and Profiling -Not applicable as a merge claim. This PR adds missing model/hardware -compatibility and does not change the default backend or claim a speedup. -Correctness on the target SM120 path is the acceptance criterion. +This is a compatibility PR, so the following result is a target-hardware +validation point rather than a general speed claim. Marlin and +`flashinfer_mxfp4` used the same SGLang commit, image, model, service flags, +requests, and three-repeat protocol. + +```text +Hardware: 4 nodes / 32 NVIDIA RTX 6000D GPUs (SM120) +Model: Kimi-K3 +Parallelism: TP32 / EP4 / DP1 +Workload: ISL=16384, OSL=1, concurrency=8, 40 requests +Chunked Prefill: 8192 tokens +SGLang: 300c87a431ac40d3e7817246376b7fe20932db09 +FlashInfer: PR #4460 implementation b525c513fef7c1d162efd3ce8636fcc8760de66a +Repeats: 3 per backend +``` + +| Backend | Median Input TPS | Median TTFT P50 | Median TTFT P95 | +|---|---:|---:|---:| +| Marlin | 2930.03 | 43.587 s | 46.287 s | +| FlashInfer MXFP4 | 3251.25 | 39.267 s | 41.694 s | +| Change | +10.96% | -9.91% | -9.92% | + +All six repeats completed 40/40 requests, all outputs contained exactly one +token, and no request errors, OOM, traceback, NCCL error, or engine-death +event was observed. The result only establishes this exact configuration; +the default backend remains unchanged. ## Checklist - [x] Format your code according to the [Format code with pre-commit](https://docs.sglang.io/developer_guide/contribution_guide.html#format-code-with-pre-commit). - [x] Add unit tests according to the [Run and add unit tests](https://docs.sglang.io/developer_guide/contribution_guide.html#run-and-add-unit-tests). -- [ ] Update documentation according to [Write documentations](https://docs.sglang.io/developer_guide/contribution_guide.html#write-documentations). (Pending until #4460 has a released API; this Draft does not change defaults.) -- [x] Provide accuracy and speed benchmark results according to [Test the accuracy](https://docs.sglang.io/developer_guide/contribution_guide.html#test-the-accuracy) and [Benchmark the speed](https://docs.sglang.io/developer_guide/contribution_guide.html#benchmark-the-speed). (Accuracy is provided above; speed is N/A because this compatibility PR makes no speedup claim and does not change the default backend.) +- [x] Update documentation according to [Write documentations](https://docs.sglang.io/developer_guide/contribution_guide.html#write-documentations). (N/A for this Draft: it does not expose a stable released capability or change defaults; user-facing documentation should be added after #4460 is released.) +- [x] Provide accuracy and speed benchmark results according to [Test the accuracy](https://docs.sglang.io/developer_guide/contribution_guide.html#test-the-accuracy) and [Benchmark the speed](https://docs.sglang.io/developer_guide/contribution_guide.html#benchmark-the-speed). - [x] Follow the SGLang code style [guidance](https://docs.sglang.io/developer_guide/contribution_guide.html#code-style-guidance). Full pre-commit validation was run on all five changed files and every hook passed. The two test files are registered with CUDA CI in `base-b` using the `1-gpu-small` runner. -The unchecked documentation item is an intentional Draft gate, not a known -code failure. - ## Dependency and Merge Order 1. FlashInfer #4460 merges and publishes an RC or release with a stable API. diff --git a/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/UPSTREAM_DUPLICATION_AUDIT.md b/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/UPSTREAM_DUPLICATION_AUDIT.md index 986522f..6a85796 100644 --- a/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/UPSTREAM_DUPLICATION_AUDIT.md +++ b/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/UPSTREAM_DUPLICATION_AUDIT.md @@ -1,6 +1,6 @@ # Kimi-K3 SM120 FlashInfer MXFP4 去重审计 -审计时间：2026-08-18 22:21:40 CST。 +审计时间：2026-08-19 13:07:48 CST。 ## 结论 @@ -85,6 +85,7 @@ activation_situ_linear_beta = 25.0 | SGLang | 禁止 SM120 进入 GB100/GB200/GB300 tcgen05 attention-residual 路径 | `5e61458` | | 本地依赖验证 | 在 #4460 之上运行精确 MXFP8×MXFP4 + Kimi SiTU 4/25 + SM120 case | 本地提交 `542310c`；不提交 FlashInfer PR | | 6000D 验收 | 601-604、TP32/EP4、16K→1、C8/C16、Chunk 8K/16K、三次重复 | `kimi3-ep4-moe-full-20260818-151349` | +| 当前 Draft 代表点 | 精确 SGLang `300c87a` + #4460 `b525c513`，TP32/EP4、16K→1、C8、Chunk 8K，各三次 | `kimi3-sm120-pr-representative-20260819-115530` | ## 明确排除的重复内容 @@ -162,6 +163,12 @@ pre-commit。最终 HEAD 的能力探测、旧版兼容和架构 gate 为 registered-tests 检查及其他 hook 全部通过。下载 GitHub hook 时只对该命令使用 HTTPS-to-SSH URL rewrite，没有修改全局 Git 配置。 +最终代表点使用当前 Draft HEAD 和 #4460 implementation commit 的统一组合， +六轮 240/240 请求成功。FlashInfer 在所有 32 rank 完成 autotune，因此 605 的 +skip-autotune workaround 仍不属于 SGLang PR 增量。该固定点的 FlashInfer +Input TPS 为 Marlin 的 `+10.96%`，TTFT P50/P95 为 `-9.91%/-9.92%`；这只作为 +SM120 集成的服务级证据，不扩展为所有 shape 的性能结论。 + 前两次失败日志分别暴露 `FLASHINFER_DISABLE_JIT=0` 仍被视为启用、源码包缺少 生成的 `data/csrc`；第三次暴露只读源码不能生成 `_build_meta.py`。这些都是 测试装配问题，最终 v4 采用可写临时源码树构建 wheel 后通过，并非功能回归。 diff --git a/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/decision_log.md b/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/decision_log.md index 6eb4d77..608deab 100644 --- a/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/decision_log.md +++ b/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/decision_log.md @@ -2,6 +2,26 @@ 本文件记录可恢复的关键决策、失败尝试和证据路径。成功结论见 `README.md`。 +## 2026-08-19 最终 PR 代表点 + +- 只补跑一个代表点，不重跑 EP32：TP32/EP4、16K→1、C=8、Chunk=8K， + Marlin/FlashInfer 各三次，Run ID + `kimi3-sm120-pr-representative-20260819-115530`。 +- 当前组合固定为 SGLang `300c87a431ac`、FlashInfer #4460 实现 + `b525c513fef7`；精确 FlashInfer wheel SHA256 为 + `5c67841b0bc5aa72d04910ee1cb727ad5fc9c13b5b80c14e26ac40de1583ce88`。 +- 六轮 240/240 请求成功，输出长度全部为 1；错误扫描为空。 +- 三次中位数：Marlin Input TPS `2930.03`、TTFT P50/P95 + `43.587/46.287 s`；FlashInfer 为 `3251.25`、`39.267/41.694 s`。 +- 该点证明当前 Draft 组合在目标硬件可用，并在这一固定 shape 上相对 Marlin + 为 Input TPS `+10.96%`、TTFT P50/P95 `-9.91%/-9.92%`；不外推到其他 + shape，也不改变默认 backend。 +- 第一次离线汇总误报 `expected 3 repeats`，根因是 benchmark JSON 的 + `backend=sglang` 覆盖了文件名解析出的 MoE backend；修复字段合并顺序后， + 原始六轮数据直接生成最终汇总，没有重跑或删除结果。 +- 公开 PR 的 Checklist 全部勾选；Speed Tests 仅报告该代表点，仍把 PR 定位为 + compatibility support。 + ## 2026-08-18 上游去重审计 - FlashInfer #4460 已覆盖 CUTLASS SiTU kernel、`ActivationType.Situ` 和独立 @@ -35,9 +55,8 @@ `/data/hzy/phase5_sglang_pr4460_draft_cpu_final.log`。 - 公开 Draft PR 正文遵循 SGLang 官方模板，不暴露内网原始路径；原始日志路径、 镜像 SHA、wheel SHA 和去重证据继续保留在本仓库审计文件中。 -- Checklist 不全部预先勾选：文档等待 #4460 发布稳定 API。公开 PR 不做提速 - 声明，因此 Speed Tests 标为 N/A；四机候选结果继续保留在内部实验档案， - 不作为该兼容性 PR 的合并门槛。 +- Checklist 在最终代表点和内部档案完成后全部勾选。公开 PR 仅报告一个固定 + shape 的目标硬件结果，不做普遍提速声明；用户文档仍等待 #4460 发布稳定 API。 ## 2026-08-18 Phase 5 diff --git a/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/results/pr_representative_16k_c8_chunk8k_results.csv b/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/results/pr_representative_16k_c8_chunk8k_results.csv new file mode 100644 index 0000000..9059f73 --- /dev/null +++ b/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/results/pr_representative_16k_c8_chunk8k_results.csv @@ -0,0 +1,7 @@ +backend,repeat,completed,request_throughput,input_throughput,total_throughput,median_ttft_ms,p95_ttft_ms,median_e2e_latency_ms +flashinfer_mxfp4,1,40,0.19844076757687357,3251.2535359794965,3251.451976747074,39266.71834103763,41694.44455524208,39266.759283607826 +flashinfer_mxfp4,2,40,0.19842530057784252,3251.000124667372,3251.19854996795,39268.08554097079,41711.827712808736,39268.12894491013 +flashinfer_mxfp4,3,40,0.198540460122689,3252.8868986501366,3253.0854391102594,39249.804597580805,41681.47474176949,39249.838569085114 +marlin,1,40,0.17882371078176418,2929.8476774484243,2930.026501159206,43587.397116585635,46338.065504981205,43587.43864367716 +marlin,2,40,0.17884277000539248,2930.1599437683503,2930.3387865383556,43590.378262568265,46287.321725080255,43590.42251249775 +marlin,3,40,0.178834593712505,2930.025983385682,2930.2048179793946,43578.72363855131,46280.03058874747,43578.76205956563 diff --git a/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/results/pr_representative_16k_c8_chunk8k_summary.json b/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/results/pr_representative_16k_c8_chunk8k_summary.json new file mode 100644 index 0000000..6dd8742 --- /dev/null +++ b/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/results/pr_representative_16k_c8_chunk8k_summary.json @@ -0,0 +1,48 @@ +{ + "run_id": "kimi3-sm120-pr-representative-20260819-115530", + "shape": { + "input_len": 16384, + "output_len": 1, + "concurrency": 8, + "chunked_prefill_size": 8192, + "tp": 32, + "ep": 4 + }, + "summary": [ + { + "backend": "marlin", + "repeats": 3, + "completed_each": [ + 40, + 40, + 40 + ], + "median_request_throughput": 0.178834593712505, + "median_input_throughput": 2930.025983385682, + "median_total_throughput": 2930.2048179793946, + "median_median_ttft_ms": 43587.397116585635, + "median_p95_ttft_ms": 46287.321725080255, + "median_median_e2e_latency_ms": 43587.43864367716 + }, + { + "backend": "flashinfer_mxfp4", + "repeats": 3, + "completed_each": [ + 40, + 40, + 40 + ], + "median_request_throughput": 0.19844076757687357, + "median_input_throughput": 3251.2535359794965, + "median_total_throughput": 3251.451976747074, + "median_median_ttft_ms": 39266.71834103763, + "median_p95_ttft_ms": 41694.44455524208, + "median_median_e2e_latency_ms": 39266.759283607826 + } + ], + "comparison": { + "input_throughput_change_pct": 10.963300476353854, + "median_ttft_change_pct": -9.912679034242956, + "p95_ttft_change_pct": -9.922538178201778 + } +} \ No newline at end of file diff --git a/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/run_pr_representative_point.sh b/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/run_pr_representative_point.sh index d3910b9..22201b8 100755 --- a/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/run_pr_representative_point.sh +++ b/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/run_pr_representative_point.sh @@ -351,7 +351,9 @@ for path in sorted((root / "raw").glob("*.jsonl")): if not m: continue data = next(json.loads(x) for x in path.read_text().splitlines() if x.strip()) - rows.append({"backend": m.group(1), "repeat": int(m.group(2)), **data}) + # The benchmark JSON reports the framework as backend="sglang". Keep the + # MoE implementation parsed from the filename as the comparison backend. + rows.append({**data, "backend": m.group(1), "repeat": int(m.group(2))}) metrics = ["request_throughput", "input_throughput", "total_throughput", "median_ttft_ms", "p95_ttft_ms", "median_e2e_latency_ms"]