Compare commits

..

10 Commits

10 changed files with 659 additions and 38 deletions

View File

@ -1,5 +1,11 @@
# sskj — 多平台大模型推理性能基准测试项目 # sskj — 多平台大模型推理性能基准测试项目
**更新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。
**更新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/B240/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 wheelSGLang 定向测试 8/8 通过。Draft 标题、依赖、兼容策略、测试和四机 EP4 数据见 `experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/SGLANG_DRAFT_PR.md` **更新2026-08-18 23:00:42 CST**:完成依赖 FlashInfer #4460 的 SGLang Draft PR 收敛。确认不提交任何 FlashInfer PR也不 vendor 或 pin 未合并 kernel在 601 GPU6 上从 #4460 源码构建 FlashInfer 0.6.18 wheelSGLang 定向测试 8/8 通过。Draft 标题、依赖、兼容策略、测试和四机 EP4 数据见 `experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/SGLANG_DRAFT_PR.md`
**更新2026-08-18 22:21:40 CST**:完成 Kimi-K3 SM120 FlashInfer MXFP4 去重审计。确认 CUTLASS SiTU kernel 已由未合并的 FlashInfer PR #4460 实现,不再提交平行 kernelSGLang 贡献收敛为 Kimi gate/up 与 scale 布局、SiTU 4.0/25.0 参数映射、非连续输入和 SM120 attention-residual guard并保留 601-604 EP4 的全部服务级结果。详见 `experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/UPSTREAM_DUPLICATION_AUDIT.md` **更新2026-08-18 22:21:40 CST**:完成 Kimi-K3 SM120 FlashInfer MXFP4 去重审计。确认 CUTLASS SiTU kernel 已由未合并的 FlashInfer PR #4460 实现,不再提交平行 kernelSGLang 贡献收敛为 Kimi gate/up 与 scale 布局、SiTU 4.0/25.0 参数映射、非连续输入和 SM120 attention-residual guard并保留 601-604 EP4 的全部服务级结果。详见 `experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/UPSTREAM_DUPLICATION_AUDIT.md`

View File

@ -1,6 +1,6 @@
# Kimi-K3 SM120 FlashInfer MXFP4 完成审计 # Kimi-K3 SM120 FlashInfer MXFP4 完成审计
审计时间2026-08-18。本文只记录 goal 的逐项验收状态;技术结论见 审计时间2026-08-19。本文只记录 goal 的逐项验收状态;技术结论见
`UPSTREAM_DUPLICATION_AUDIT.md`,公开 PR 正文见 `SGLANG_DRAFT_PR.md` `UPSTREAM_DUPLICATION_AUDIT.md`,公开 PR 正文见 `SGLANG_DRAFT_PR.md`
## 验收矩阵 ## 验收矩阵
@ -11,18 +11,21 @@
| 区分官方 0.6.17 与实际能力 | 已证明 | 605 profile 设置 `FLASHINFER_DISABLE_VERSION_CHECK=1` 并启动前执行两个补丁;报告未再声称 0.6.17 原版直接支持完整组合 | | 区分官方 0.6.17 与实际能力 | 已证明 | 605 profile 设置 `FLASHINFER_DISABLE_VERSION_CHECK=1` 并启动前执行两个补丁;报告未再声称 0.6.17 原版直接支持完整组合 |
| 逐文件比较 605 补丁、FlashInfer #4460、SGLang 上游 | 已证明 | `UPSTREAM_DUPLICATION_AUDIT.md` 的“上游已有/605 已有/我们新增/明确排除”四部分605 补丁目标已从原文件重新核验 | | 逐文件比较 605 补丁、FlashInfer #4460、SGLang 上游 | 已证明 | `UPSTREAM_DUPLICATION_AUDIT.md` 的“上游已有/605 已有/我们新增/明确排除”四部分605 补丁目标已从原文件重新核验 |
| 不提交重复 FlashInfer kernel | 已满足 | FlashInfer 历史实现与测试仅作为本地研究记录SGLang Draft 不含 FlashInfer C++、依赖 pin 或 private capability marker | | 不提交重复 FlashInfer kernel | 已满足 | FlashInfer 历史实现与测试仅作为本地研究记录SGLang Draft 不含 FlashInfer C++、依赖 pin 或 private capability marker |
| SGLang 仅保留真实增量 | 已满足 | SGLang base `63d783bbe095` 到 head `300c87a431ac` 仅修改 5 个文件;补丁位于 `patches/sglang-kimi-k3-sm120-flashinfer-mxfp4-draft.patch` | | SGLang 仅保留真实增量 | 已满足 | SGLang base `593b1a9b8aee` 到 head `b95b5341ab95` 仅修改 5 个文件;补丁位于 `patches/sglang-kimi-k3-sm120-flashinfer-mxfp4-draft.patch` |
| 可应用 patch 与 Draft HEAD 一致 | 已证明 | 在全新 detached base `593b1a9b8aee` 应用仓库 patch 并 `git add -A` 后,暂存树与 `b95b5341ab95` 的 diff 为空 |
| Kimi gate/up 权重和 scale 布局 | 已证明 | 注册测试覆盖连续 `[gate; up]` 到 CUTLASS `[up; gate]`,并保留 GPT-OSS pair-interleaved 分支 | | Kimi gate/up 权重和 scale 布局 | 已证明 | 注册测试覆盖连续 `[gate; up]` 到 CUTLASS `[up; gate]`,并保留 GPT-OSS pair-interleaved 分支 |
| SiTU 4.0/25.0 参数透传 | 已证明 | #4460 ABI 的 SM120 GPU correctness 测试覆盖独立 `situ_beta``situ_linear_beta` | | SiTU 4.0/25.0 参数透传 | 已证明 | #4460 ABI 的 SM120 GPU correctness 测试覆盖独立 `situ_beta``situ_linear_beta` |
| 非连续 hidden states | 已证明 | GPU correctness 测试输入为非连续 viewrunner 在 MXFP8 quantize 边界 materialize contiguous tensor | | 非连续 hidden states | 已证明 | GPU correctness 测试输入为非连续 viewrunner 在 MXFP8 quantize 边界 materialize contiguous tensor |
| attn_residual SM120 guard | 已证明 | 注册架构测试覆盖 SM100/103/110 为 true、SM120/SM90 为 false | | attn_residual SM120 guard | 已证明 | 注册架构测试覆盖 SM100/103/110 为 true、SM120/SM90 为 false |
| autotune 稳定性边界 | 已证明且不纳入 PR | 605 的 skip workaround 未进入 Draft正式 24 点候选矩阵未复现 autotune 崩溃,因此没有用历史 workaround 扩大 PR | | autotune 稳定性边界 | 已证明且不纳入 PR | 605 的 skip workaround 未进入 Draft正式 24 点候选矩阵未复现 autotune 崩溃,因此没有用历史 workaround 扩大 PR |
| 最终定向 correctness | 已通过 | RTX 6000D GPU6、#4460 snapshot`8 passed, 17 warnings in 559.87s`wheel SHA256 `deafdac2dc44b331441daba64e5b5dd28f591fe44f087af0aac3521c6ddbfa1f` | | 最终定向 correctness | 已通过 | RTX 6000D GPU6、#4460 snapshot`8 passed, 17 warnings in 559.87s`wheel SHA256 `deafdac2dc44b331441daba64e5b5dd28f591fe44f087af0aac3521c6ddbfa1f` |
| 最终兼容回归 | 已通过 | Draft HEAD `300c87a``8 passed, 5 deselected, 17 warnings in 22.23s` | | fork 同步与补丁等价性 | 已通过 | 旧 base/head `63d783b..300c87a` 重放到 `593b1a9..b95b534`7/7 commits 在 `git range-diff` 中均为 `=` |
| 最终兼容回归 | 已通过 | Draft HEAD `b95b534`compileall 通过,`10 passed, 3 deselected, 17 warnings in 33.30s` |
| 格式和 CI 注册 | 已通过 | 五个变更文件完整 pre-commit 全部通过;两个测试文件注册到 CUDA CI `base-b/1-gpu-small` | | 格式和 CI 注册 | 已通过 | 五个变更文件完整 pre-commit 全部通过;两个测试文件注册到 CUDA CI `base-b/1-gpu-small` |
| 保留 EP4 服务级结果 | 已证明 | 601 原始目录有 24 个正式 JSONL 点;每组 3 repeats、每点 40/40 成功;错误扫描为空;未重跑 EP32 | | 保留 EP4 服务级结果 | 已证明 | 601 原始目录有 24 个正式 JSONL 点;每组 3 repeats、每点 40/40 成功;错误扫描为空;未重跑 EP32 |
| 公开 PR 标题、正文、依赖和已知限制 | 已完成 | `SGLANG_DRAFT_PR.md` 遵循 SGLang 官方模板;标题为 `[Feature] Support Kimi-K3 FlashInfer MXFP4 MoE on SM120` | | 当前 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% |
| GitHub fork 交付 | 已完成 | `ACEEE-1222/sglang:hzy/kimi-k3-sm120-flashinfer-mxfp4`Draft HEAD `300c87a` | | 公开 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 `b95b534`,远端 ref 已复核 |
| 内部仓库独立分支 | 本地和 601 已完成,远端待推送 | 601 分支 `hzy/feature-kimi-k3-sm120-flashinfer-mxfp4` 已到审计提交;`git.meta-stone.net` 尚未为本机或 601 key 授权HTTPS 也无 Token | | 内部仓库独立分支 | 本地和 601 已完成,远端待推送 | 601 分支 `hzy/feature-kimi-k3-sm120-flashinfer-mxfp4` 已到审计提交;`git.meta-stone.net` 尚未为本机或 601 key 授权HTTPS 也无 Token |
## 原始证据位置 ## 原始证据位置
@ -38,6 +41,10 @@
# 当前 SGLang Draft worktree601 # 当前 SGLang Draft worktree601
/data/hzy/src/sglang-kimi-sm120-draft-wt /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
``` ```
## 唯一未关闭项 ## 唯一未关闭项

View File

@ -0,0 +1,29 @@
ARG BASE_IMAGE=lmsysorg/sglang:kimi-k3-fiv617situ-warm
FROM ${BASE_IMAGE}
ARG SGLANG_COMMIT
ARG FLASHINFER_COMMIT
ENV FLASHINFER_DISABLE_VERSION_CHECK=1 \
PYTHONPATH=/sgl-workspace/sglang/python \
PYTHONUNBUFFERED=1
COPY flashinfer_python-0.6.18-py3-none-any.whl /tmp/
RUN python3 -m pip install --no-deps --force-reinstall \
/tmp/flashinfer_python-0.6.18-py3-none-any.whl && \
rm -f /tmp/*.whl
# Replace the image's older Python package with the exact Draft tree. Keep the
# base image's ABI-matched torch and sglang-kernel binaries: the Draft changes
# only the Python Kimi/FlashInfer integration and does not require a new kernel.
RUN rm -rf /sgl-workspace/sglang/python/sglang
COPY sglang/ /sgl-workspace/sglang/python/sglang/
COPY test_mxfp4_sm120_cutlass.py /opt/pr-tests/test_mxfp4_sm120_cutlass.py
COPY source_identity.txt /opt/pr-build/source_identity.txt
RUN python3 -m compileall -q /sgl-workspace/sglang/python/sglang && \
python3 -c "import inspect; from flashinfer.fused_moe import cutlass_fused_moe; from flashinfer.fused_moe.core import ActivationType; assert hasattr(ActivationType, 'Situ'); assert 'situ_beta' in inspect.signature(cutlass_fused_moe).parameters; import sglang.srt.layers.quantization.mxfp4"
LABEL ai.meta-stone.purpose="Kimi-K3 SM120 SGLang Draft representative validation" \
ai.meta-stone.sglang.commit="${SGLANG_COMMIT}" \
ai.meta-stone.flashinfer.commit="${FLASHINFER_COMMIT}"

View File

@ -7,15 +7,25 @@
> Kimi-K3 SGLang 接入、精确 MXFP8×MXFP4+SiTU 测试和 6000D 实测证据。 > Kimi-K3 SGLang 接入、精确 MXFP8×MXFP4+SiTU 测试和 6000D 实测证据。
> 权威边界见 [`UPSTREAM_DUPLICATION_AUDIT.md`](UPSTREAM_DUPLICATION_AUDIT.md)。 > 权威边界见 [`UPSTREAM_DUPLICATION_AUDIT.md`](UPSTREAM_DUPLICATION_AUDIT.md)。
当前 SGLang Draft HEAD 为 `300c87a431ac`,已推送到 当前 SGLang Draft 基于 `main` `593b1a9b8a`HEAD 为 `b95b5341ab95`,已推送到
`ACEEE-1222/sglang:hzy/kimi-k3-sm120-flashinfer-mxfp4`。公开 PR 正文见 `ACEEE-1222/sglang:hzy/kimi-k3-sm120-flashinfer-mxfp4`。公开 PR 正文见
[`SGLANG_DRAFT_PR.md`](SGLANG_DRAFT_PR.md),可审查补丁见 [`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)。 [`patches/sglang-kimi-k3-sm120-flashinfer-mxfp4-draft.patch`](patches/sglang-kimi-k3-sm120-flashinfer-mxfp4-draft.patch)。
五个变更文件已通过完整 pre-commit最终非 JIT 回归为 五个变更文件已通过完整 pre-commit同步 fork 后7 个提交的
`8 passed, 5 deselected, 17 warnings in 22.23s` `range-diff` 全部为补丁等价,最终非 JIT 回归为
`10 passed, 3 deselected, 17 warnings in 33.30s`
逐项验收和唯一剩余交付项见 逐项验收和唯一剩余交付项见
[`COMPLETION_AUDIT.md`](COMPLETION_AUDIT.md)。 [`COMPLETION_AUDIT.md`](COMPLETION_AUDIT.md)。
最终 PR 代表点已经在同步前的等价 patch `300c87a` 统一镜像中完成:四节点
TP32/EP4、16K→1、C=8、Chunk=8KMarlin 与 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 SM120RTX PRO 6000D上的 SGLang 部署补齐 FlashInfer MXFP8 activation × MXFP4 weight fused MoE并形成可回归、可性能比较、可提交上游的最小改动。 为 Kimi-K3 在 NVIDIA SM120RTX PRO 6000D上的 SGLang 部署补齐 FlashInfer MXFP8 activation × MXFP4 weight fused MoE并形成可回归、可性能比较、可提交上游的最小改动。
@ -704,3 +714,34 @@ MXFP8×MXFP4 参数组,覆盖 FP16/BF16 与 autotune on/off
该 FlashInfer 测试提交仅保留为本地验证,不提交 FlashInfer PR。SGLang Draft 该 FlashInfer 测试提交仅保留为本地验证,不提交 FlashInfer PR。SGLang Draft
不 vendor kernel、不 pin 未合并 commit、也不改变默认 MoE backend。 不 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
```

View File

@ -64,8 +64,9 @@ MXFP4 weights, SiTU `(4.0, 25.0)`, Kimi gate/up layout, and non-contiguous
input. It compares the SGLang adapter output directly with the public input. It compares the SGLang adapter output directly with the public
FlashInfer CUTLASS API. FlashInfer CUTLASS API.
After the final compatibility and CI-registration changes, the non-JIT subset After syncing the fork and rebasing the seven-patch series onto SGLang main
was rerun from the final commit: `593b1a9b8aeedb908ee21870c02a0fd1a3dc30a5`, the non-JIT subset was rerun
from the final Draft head `b95b5341ab9555ba5ea53f5828f6de203d749cce`:
```bash ```bash
pytest -q -s \ pytest -q -s \
@ -75,38 +76,51 @@ pytest -q -s \
``` ```
```text ```text
8 passed, 5 deselected, 17 warnings in 22.23s 10 passed, 3 deselected, 17 warnings in 33.30s
``` ```
## Speed Tests and Profiling ## Speed Tests and Profiling
Four-node serving validation used 32 RTX 6000D GPUs with TP32/EP4/DP1, This is a compatibility PR, so the following result is a target-hardware
16K input tokens, one output token, concurrency 8/16, and chunked Prefill validation point rather than a general speed claim. Marlin and
sizes 8K/16K. Each point was repeated three times; all 24 runs completed with `flashinfer_mxfp4` used the same SGLang commit, image, model, service flags,
40/40 successful requests and no OOM, Traceback, NCCL error, or requests, and three-repeat protocol.
`EngineDeadError`.
| Chunk | C | Backend | Input TPS | TTFT P50 (s) | FlashInfer vs Marlin | ```text
|---:|---:|---|---:|---:|---:| Hardware: 4 nodes / 32 NVIDIA RTX 6000D GPUs (SM120)
| 8K | 8 | Marlin | 2935.02 | 43.51 | baseline | Model: Kimi-K3
| 8K | 8 | FlashInfer | 3257.96 | 39.19 | +11.00% TPS, -9.92% TTFT | Parallelism: TP32 / EP4 / DP1
| 8K | 16 | Marlin | 2934.90 | 87.09 | baseline | Workload: ISL=16384, OSL=1, concurrency=8, 40 requests
| 8K | 16 | FlashInfer | 3260.14 | 78.39 | +11.08% TPS, -9.99% TTFT | Chunked Prefill: 8192 tokens
| 16K | 8 | Marlin | 2812.91 | 45.47 | baseline | SGLang patch series used for the run: 300c87a431ac40d3e7817246376b7fe20932db09
| 16K | 8 | FlashInfer | 3027.50 | 42.27 | +7.63% TPS, -7.04% TTFT | Final rebased Draft head: b95b5341ab9555ba5ea53f5828f6de203d749cce
| 16K | 16 | Marlin | 2812.33 | 90.97 | baseline | Final SGLang main base: 593b1a9b8aeedb908ee21870c02a0fd1a3dc30a5
| 16K | 16 | FlashInfer | 3027.34 | 84.52 | +7.65% TPS, -7.09% TTFT | FlashInfer: PR #4460 implementation b525c513fef7c1d162efd3ce8636fcc8760de66a
Repeats: 3 per backend
```
The service benchmark used the previously validated candidate integration The service benchmark was collected before the fork sync. `git range-diff`
with the same Kimi layout, SiTU, non-contiguous-input, and SM120 guard fixes. reports all seven commits as patch-equivalent after the rebase, and the
The final public #4460 API is covered by the accuracy test above; no claim is directed tests above were rerun from the final Draft head. The performance
made that an unreleased #4460 wheel was used for this four-node benchmark. result is therefore retained without claiming that the benchmark container
contained the new commit object.
| 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 ## 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] 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). - [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).
- [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 docs should be added after #4460 is released.) - [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] 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). - [x] Follow the SGLang code style [guidance](https://docs.sglang.io/developer_guide/contribution_guide.html#code-style-guidance).

View File

@ -1,6 +1,6 @@
# Kimi-K3 SM120 FlashInfer MXFP4 去重审计 # 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` | | SGLang | 禁止 SM120 进入 GB100/GB200/GB300 tcgen05 attention-residual 路径 | `5e61458` |
| 本地依赖验证 | 在 #4460 之上运行精确 MXFP8×MXFP4 + Kimi SiTU 4/25 + SM120 case | 本地提交 `542310c`;不提交 FlashInfer PR | | 本地依赖验证 | 在 #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` | | 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` |
## 明确排除的重复内容 ## 明确排除的重复内容
@ -102,11 +103,15 @@ activation_situ_linear_beta = 25.0
SGLang Draft 的完整可审查 diff SGLang Draft 的完整可审查 diff
```text ```text
base: 63d783bbe0955237ec41f9ddabf7235ddf04673c base: 593b1a9b8aeedb908ee21870c02a0fd1a3dc30a5
head: 300c87a431ac40d3e7817246376b7fe20932db09 head: b95b5341ab9555ba5ea53f5828f6de203d749cce
patch: patches/sglang-kimi-k3-sm120-flashinfer-mxfp4-draft.patch patch: patches/sglang-kimi-k3-sm120-flashinfer-mxfp4-draft.patch
``` ```
该 7 提交系列由旧范围 `63d783b..300c87a` 重放到用户新同步的 main
`git range-diff` 逐项均为 `=`。在新 base 的独立 worktree 应用上述 patch、
暂存全部文件后,与 `b95b534` 比较的 diff 为空。
GitHub 交付分支: GitHub 交付分支:
```text ```text
@ -153,15 +158,21 @@ GPU: NVIDIA RTX 6000D, SM120
最终审查发现普通 SwiGLU 路径也曾无条件访问 `ActivationType.Situ`,可能误伤 最终审查发现普通 SwiGLU 路径也曾无条件访问 `ActivationType.Situ`,可能误伤
旧 FlashInfer。提交 `775aa076` 改为动态加入 SiTU enum并新增旧版兼容回归。 旧 FlashInfer。提交 `775aa076` 改为动态加入 SiTU enum并新增旧版兼容回归。
最终提交 `300c87a` 为两个测试文件增加正式 CUDA CI 注册,并通过完整 同步前提交 `300c87a` 为两个测试文件增加正式 CUDA CI 注册,并通过完整
pre-commit。最终 HEAD 的能力探测、旧版兼容和架构 gate 为 pre-commit。同步后的最终 HEAD `b95b534` 在 601 通过 compileall能力探测、
`8 passed, 5 deselected, 17 warnings in 22.23s`GPU kernel 路径继续由前述 旧版兼容和架构 gate 为 `10 passed, 3 deselected, 17 warnings in 33.30s`
8/8 测试覆盖。 GPU kernel 路径继续由前述 8/8 测试覆盖。
完整 pre-commit 命令覆盖全部五个变更文件isort、ruff、black、codespell、 完整 pre-commit 命令覆盖全部五个变更文件isort、ruff、black、codespell、
registered-tests 检查及其他 hook 全部通过。下载 GitHub hook 时只对该命令使用 registered-tests 检查及其他 hook 全部通过。下载 GitHub hook 时只对该命令使用
HTTPS-to-SSH URL rewrite没有修改全局 Git 配置。 HTTPS-to-SSH URL rewrite没有修改全局 Git 配置。
最终代表点使用同步前等价 patch `300c87a`#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` 仍被视为启用、源码包缺少 前两次失败日志分别暴露 `FLASHINFER_DISABLE_JIT=0` 仍被视为启用、源码包缺少
生成的 `data/csrc`;第三次暴露只读源码不能生成 `_build_meta.py`。这些都是 生成的 `data/csrc`;第三次暴露只读源码不能生成 `_build_meta.py`。这些都是
测试装配问题,最终 v4 采用可写临时源码树构建 wheel 后通过,并非功能回归。 测试装配问题,最终 v4 采用可写临时源码树构建 wheel 后通过,并非功能回归。

View File

@ -2,6 +2,38 @@
本文件记录可恢复的关键决策、失败尝试和证据路径。成功结论见 `README.md` 本文件记录可恢复的关键决策、失败尝试和证据路径。成功结论见 `README.md`
## 2026-08-19 fork 同步与最终交付
- 用户重新同步 `ACEEE-1222/sglang:main` 后,最终 base 为 `593b1a9b8aee`
- 原 7 提交系列从 `63d783b..300c87a` 无冲突重放为
`593b1a9..b95b534``git range-diff` 的 7 项全部为 `=`,没有语义改动。
- 最终 HEAD 在 601 统一验证镜像中通过 compileall 和定向非 JIT 回归:
`10 passed, 3 deselected, 17 warnings in 33.30s`
- GitHub feature 分支通过精确 `force-with-lease``300c87a` 更新到
`b95b534`,随后用 `ls-remote` 复核远端 ref。
- 四机代表点不重跑。其容器记录的是同步前 `300c87a`,但对应 patch 与最终
HEAD 逐提交等价PR 正文同时列出实验 patch SHA 和最终 HEAD不混淆对象。
## 2026-08-19 最终 PR 代表点
- 只补跑一个代表点,不重跑 EP32TP32/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 上游去重审计 ## 2026-08-18 上游去重审计
- FlashInfer #4460 已覆盖 CUTLASS SiTU kernel、`ActivationType.Situ` 和独立 - FlashInfer #4460 已覆盖 CUTLASS SiTU kernel、`ActivationType.Situ` 和独立
@ -35,6 +67,8 @@
`/data/hzy/phase5_sglang_pr4460_draft_cpu_final.log` `/data/hzy/phase5_sglang_pr4460_draft_cpu_final.log`
- 公开 Draft PR 正文遵循 SGLang 官方模板,不暴露内网原始路径;原始日志路径、 - 公开 Draft PR 正文遵循 SGLang 官方模板,不暴露内网原始路径;原始日志路径、
镜像 SHA、wheel SHA 和去重证据继续保留在本仓库审计文件中。 镜像 SHA、wheel SHA 和去重证据继续保留在本仓库审计文件中。
- Checklist 在最终代表点和内部档案完成后全部勾选。公开 PR 仅报告一个固定
shape 的目标硬件结果,不做普遍提速声明;用户文档仍等待 #4460 发布稳定 API。
## 2026-08-18 Phase 5 ## 2026-08-18 Phase 5

View File

@ -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
1 backend repeat completed request_throughput input_throughput total_throughput median_ttft_ms p95_ttft_ms median_e2e_latency_ms
2 flashinfer_mxfp4 1 40 0.19844076757687357 3251.2535359794965 3251.451976747074 39266.71834103763 41694.44455524208 39266.759283607826
3 flashinfer_mxfp4 2 40 0.19842530057784252 3251.000124667372 3251.19854996795 39268.08554097079 41711.827712808736 39268.12894491013
4 flashinfer_mxfp4 3 40 0.198540460122689 3252.8868986501366 3253.0854391102594 39249.804597580805 41681.47474176949 39249.838569085114
5 marlin 1 40 0.17882371078176418 2929.8476774484243 2930.026501159206 43587.397116585635 46338.065504981205 43587.43864367716
6 marlin 2 40 0.17884277000539248 2930.1599437683503 2930.3387865383556 43590.378262568265 46287.321725080255 43590.42251249775
7 marlin 3 40 0.178834593712505 2930.025983385682 2930.2048179793946 43578.72363855131 46280.03058874747 43578.76205956563

View File

@ -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
}
}

View File

@ -0,0 +1,424 @@
#!/usr/bin/env bash
# Validate the exact SGLang Draft with one four-node Kimi-K3 serving point.
set -Eeuo pipefail
SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
REPO_ROOT="$(cd "${SCRIPT_DIR}/../../.." && pwd)"
# shellcheck source=/dev/null
source "${REPO_ROOT}/scripts/common/lib.sh"
ACTION="${1:-all}"
RUN_ID="${RUN_ID:-kimi3-sm120-pr-representative-$(date '+%Y%m%d-%H%M%S')}"
RESULT_ROOT="${RESULT_ROOT:-${SCRIPT_DIR}/results/${RUN_ID}}"
MODEL_PATH="${MODEL_PATH:-/data/hf_models/Kimi-K3}"
SERVED_MODEL_NAME="${SERVED_MODEL_NAME:-kimi-k3}"
BASE_IMAGE="${BASE_IMAGE:-lmsysorg/sglang:kimi-k3-fiv617situ-warm}"
PR_IMAGE="${PR_IMAGE:-local/sglang:kimi-k3-sm120-pr-300c87a-fi-b525c51}"
SGLANG_SOURCE="${SGLANG_SOURCE:-/data/hzy/src/sglang-kimi-sm120-draft-wt}"
FLASHINFER_SOURCE="${FLASHINFER_SOURCE:-/data/hzy/src/flashinfer-pr4460-kimi-test-wt}"
ARTIFACT_DIR="${ARTIFACT_DIR:-/data/hzy/artifacts/sglang-pr-kimi-sm120-300c87a}"
FLASHINFER_WHEEL="${FLASHINFER_WHEEL:-${ARTIFACT_DIR}/flashinfer_python-0.6.18-py3-none-any.whl}"
BUILD_CONTEXT="${BUILD_CONTEXT:-/tmp/kimi3_sm120_pr_validation_context}"
FLASHINFER_CACHE="${FLASHINFER_CACHE:-/data/hzy/cache/flashinfer-pr-b525c51}"
HEAD_HOST="${HEAD_HOST:-174.1.60.1}"
NODE_SSH_USER="${NODE_SSH_USER:-user}"
NODE_HOSTS=(174.1.60.1 174.1.60.2 174.1.60.3 174.1.60.4)
SSH_OPTS=(-o BatchMode=yes -o StrictHostKeyChecking=no -o ConnectTimeout=10)
DIST_PORT="${DIST_PORT:-20000}"
PORT="${PORT:-30000}"
CONTAINER_PREFIX="kimi3_sm120_pr_validation"
INPUT_LEN=16384
OUTPUT_LEN=1
CONCURRENCY=8
CHUNKED_PREFILL_SIZE=8192
NUM_PROMPTS="${NUM_PROMPTS:-40}"
REPEATS="${REPEATS:-3}"
WARMUP_REQUESTS="${WARMUP_REQUESTS:-2}"
HEALTH_WAIT_S="${HEALTH_WAIT_S:-2400}"
BACKENDS=(marlin flashinfer_mxfp4)
mkdir -p "${RESULT_ROOT}"/{build,service,raw,bench,gpu}
log_init "${RESULT_ROOT}/orchestrator.log"
require_password() {
if [[ -z "${SUDO_PASSWORD:-}" && -n "${SUDO_PASSWORD_FILE:-}" ]]; then
[[ -r "${SUDO_PASSWORD_FILE}" ]] || {
echo "ERROR: cannot read SUDO_PASSWORD_FILE=${SUDO_PASSWORD_FILE}" >&2
exit 2
}
IFS= read -r SUDO_PASSWORD <"${SUDO_PASSWORD_FILE}"
fi
[[ -n "${SUDO_PASSWORD:-}" ]] || {
echo "ERROR: set SUDO_PASSWORD or SUDO_PASSWORD_FILE" >&2
exit 2
}
}
is_head() { [[ "$1" == "${HEAD_HOST}" ]]; }
sudo_host() {
local host="$1"
shift
require_password
if is_head "$host"; then
printf '%s\n' "${SUDO_PASSWORD}" | sudo -S -p '' -- "$@"
else
local remote_cmd
printf -v remote_cmd '%q ' "$@"
printf '%s\n' "${SUDO_PASSWORD}" | \
ssh "${SSH_OPTS[@]}" "${NODE_SSH_USER}@${host}" \
"sudo -S -p '' -- ${remote_cmd}"
fi
}
check_inputs() {
local path
for path in "${MODEL_PATH}" "${SGLANG_SOURCE}/python/sglang" \
"${FLASHINFER_SOURCE}" "${FLASHINFER_WHEEL}" \
"${SCRIPT_DIR}/Dockerfile.pr_validation"; do
[[ -e "$path" ]] || { echo "ERROR: missing ${path}" >&2; exit 2; }
done
SGLANG_COMMIT="$(git -C "${SGLANG_SOURCE}" rev-parse HEAD)"
FLASHINFER_COMMIT="$(git -C "${FLASHINFER_SOURCE}" rev-parse HEAD~1)"
[[ "${SGLANG_COMMIT}" == 300c87a431ac40d3e7817246376b7fe20932db09 ]] || {
echo "ERROR: unexpected SGLang commit ${SGLANG_COMMIT}" >&2
exit 2
}
[[ "${FLASHINFER_COMMIT}" == b525c51* ]] || {
echo "ERROR: unexpected FlashInfer implementation commit ${FLASHINFER_COMMIT}" >&2
exit 2
}
export SGLANG_COMMIT FLASHINFER_COMMIT
}
prepare_build_context() {
log "preparing exact Draft build context"
rm -rf "${BUILD_CONTEXT}"
mkdir -p "${BUILD_CONTEXT}/sglang"
cp "${SCRIPT_DIR}/Dockerfile.pr_validation" "${BUILD_CONTEXT}/Dockerfile"
cp "${FLASHINFER_WHEEL}" "${BUILD_CONTEXT}/"
cp "${SGLANG_SOURCE}/test/registered/unit/layers/quantization/test_mxfp4_sm120_cutlass.py" \
"${BUILD_CONTEXT}/test_mxfp4_sm120_cutlass.py"
tar -C "${SGLANG_SOURCE}/python/sglang" -cf - . | \
tar -C "${BUILD_CONTEXT}/sglang" -xf -
{
printf 'sglang=%s\n' "${SGLANG_COMMIT}"
printf 'flashinfer_implementation=%s\n' "${FLASHINFER_COMMIT}"
sha256sum "${FLASHINFER_WHEEL}"
} >"${BUILD_CONTEXT}/source_identity.txt"
cp "${BUILD_CONTEXT}/source_identity.txt" "${RESULT_ROOT}/build/"
du -sh "${BUILD_CONTEXT}" | tee "${RESULT_ROOT}/build/context_size.txt"
}
stage_context() {
local host="$1"
is_head "$host" && return
ssh "${SSH_OPTS[@]}" "${NODE_SSH_USER}@${host}" \
"rm -rf '${BUILD_CONTEXT}' && mkdir -p '${BUILD_CONTEXT}'"
tar -C "${BUILD_CONTEXT}" -cf - . | \
ssh "${SSH_OPTS[@]}" "${NODE_SSH_USER}@${host}" \
"tar -C '${BUILD_CONTEXT}' -xf -"
}
build_one_node() {
local host="$1"
sudo_host "$host" docker build \
--build-arg "BASE_IMAGE=${BASE_IMAGE}" \
--build-arg "SGLANG_COMMIT=${SGLANG_COMMIT}" \
--build-arg "FLASHINFER_COMMIT=${FLASHINFER_COMMIT}" \
--tag "${PR_IMAGE}" "${BUILD_CONTEXT}" \
>"${RESULT_ROOT}/build/${host}.log" 2>&1
sudo_host "$host" docker image inspect "${PR_IMAGE}" \
--format '{{.Id}} {{.Size}} {{json .Config.Labels}}' \
>"${RESULT_ROOT}/build/${host}.image.txt"
sudo_host "$host" docker run --rm --entrypoint python3 "${PR_IMAGE}" -c \
"import importlib.metadata as m; print('sglang=' + m.version('sglang')); print('sglang-kernel=' + m.version('sglang-kernel')); print('flashinfer-python=' + m.version('flashinfer-python')); print(open('/opt/pr-build/source_identity.txt').read(), end='')" \
>"${RESULT_ROOT}/build/${host}.packages.txt"
}
build_all_nodes() {
check_inputs
prepare_build_context
local host pid rc=0
for host in "${NODE_HOSTS[@]}"; do stage_context "$host"; done
local -a pids=()
for host in "${NODE_HOSTS[@]}"; do
build_one_node "$host" &
pids+=("$!")
done
for pid in "${pids[@]}"; do wait "$pid" || rc=1; done
(( rc == 0 )) || { log "ERROR: image build failed"; return 1; }
log "exact Draft image built on all nodes"
}
prewarm_one_node() {
local host="$1"
sudo_host "$host" mkdir -p "${FLASHINFER_CACHE}"
sudo_host "$host" docker run --rm --gpus device=0 \
-v "${FLASHINFER_CACHE}:/root/.cache/flashinfer" \
-e FLASHINFER_WORKSPACE_BASE=/root \
-e FLASHINFER_DISABLE_VERSION_CHECK=1 \
--entrypoint python3 "${PR_IMAGE}" -m pytest -q -s \
/opt/pr-tests/test_mxfp4_sm120_cutlass.py \
-k kimi_k3_sm120_situ_layout_and_noncontiguous_input \
>"${RESULT_ROOT}/build/${host}.prewarm.log" 2>&1
}
prewarm_all_nodes() {
local host pid rc=0
local -a pids=()
for host in "${NODE_HOSTS[@]}"; do
prewarm_one_node "$host" &
pids+=("$!")
done
for pid in "${pids[@]}"; do wait "$pid" || rc=1; done
(( rc == 0 )) || { log "ERROR: FlashInfer prewarm failed"; return 1; }
log "exact #4460 kernel prewarmed on all nodes"
}
container_name() { printf '%s_node%s' "${CONTAINER_PREFIX}" "$1"; }
stop_service() {
local rank host
for rank in 0 1 2 3; do
host="${NODE_HOSTS[$rank]}"
sudo_host "$host" docker rm -f "$(container_name "$rank")" \
>/dev/null 2>&1 || true
done
}
collect_service_logs() {
local label="$1" rank host
for rank in 0 1 2 3; do
host="${NODE_HOSTS[$rank]}"
sudo_host "$host" docker logs "$(container_name "$rank")" \
>"${RESULT_ROOT}/service/${label}_node${rank}.log" 2>&1 || true
done
}
collect_gpu() {
local label="$1" rank host
for rank in 0 1 2 3; do
host="${NODE_HOSTS[$rank]}"
sudo_host "$host" nvidia-smi \
--query-gpu=timestamp,index,memory.used,memory.total,utilization.gpu,power.draw \
--format=csv,noheader,nounits \
>"${RESULT_ROOT}/gpu/${label}_node${rank}.csv" 2>&1 || true
done
}
verify_service_logs() {
local label="$1"
local pattern='CUDA out of memory|torch\.OutOfMemoryError|Traceback|EngineDeadError|NCCL[^[:cntrl:]]*(error|failed)|connection refused|Terminated'
if grep -Ein "${pattern}" "${RESULT_ROOT}/service/${label}_node"*.log \
>"${RESULT_ROOT}/service/${label}_fatal_scan.txt"; then
log "ERROR: fatal pattern found in service logs label=${label}"
return 1
fi
: >"${RESULT_ROOT}/service/${label}_fatal_scan.txt"
}
start_node() {
local rank="$1" backend="$2"
local host="${NODE_HOSTS[$rank]}" name bootstrap
name="$(container_name "$rank")"
bootstrap="export SGLANG_HOST_IP=174.1.60.$((rank + 1)); exec python3 -m sglang.launch_server --model-path ${MODEL_PATH} --served-model-name ${SERVED_MODEL_NAME} --tp-size 32 --ep-size 4 --nnodes 4 --node-rank ${rank} --dist-init-addr ${HEAD_HOST}:${DIST_PORT} --trust-remote-code --moe-runner-backend ${backend} --chunked-prefill-size ${CHUNKED_PREFILL_SIZE} --mem-fraction-static 0.88 --cuda-graph-max-bs-decode 16 --mamba-radix-cache-strategy extra_buffer_lazy --disable-radix-cache --dist-timeout 3600 --mamba-full-memory-ratio 0.36 --host 0.0.0.0 --port ${PORT}"
local -a cmd=(
docker run -d --name "$name"
--gpus all --network host --ipc=host --ulimit memlock=-1
--device /dev/infiniband --shm-size 32g --entrypoint bash
-v "${MODEL_PATH}:${MODEL_PATH}:ro"
-v "${FLASHINFER_CACHE}:/root/.cache/flashinfer"
-e CUDA_VISIBLE_DEVICES=0,1,2,3,4,5,6,7
-e NCCL_SOCKET_IFNAME=bond0 -e GLOO_SOCKET_IFNAME=bond0
-e NCCL_IB_HCA=mlx5_0,mlx5_1,mlx5_2,mlx5_3
-e NCCL_IB_GID_INDEX=3 -e NCCL_IB_TIMEOUT=22 -e NCCL_IB_RETRY_CNT=7
-e NCCL_CUMEM_ENABLE=1 -e PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True
-e SGLANG_ENABLE_TP_MEMORY_INBALANCE_CHECK=0 -e SGLANG_MOE_FUSED_GATE_RADIX=1
-e SGLANG_SKIP_SGL_KERNEL_VERSION_CHECK=1
-e FLASHINFER_WORKSPACE_BASE=/root
-e FLASHINFER_DISABLE_VERSION_CHECK=1
"${PR_IMAGE}" -lc "$bootstrap"
)
printf '%q ' "${cmd[@]}" \
>"${RESULT_ROOT}/service/${backend}_node${rank}.cmd.txt"
printf '\n' >>"${RESULT_ROOT}/service/${backend}_node${rank}.cmd.txt"
sudo_host "$host" "${cmd[@]}" >/dev/null
}
wait_health() {
local backend="$1" i
for ((i = 1; i <= HEALTH_WAIT_S; i++)); do
if curl --fail --silent --max-time 5 \
"http://${HEAD_HOST}:${PORT}/health" >/dev/null 2>&1; then
log "service healthy backend=${backend} wait_s=${i}"
return 0
fi
if (( i % 30 == 0 )); then
log "waiting backend=${backend} elapsed_s=${i}"
collect_service_logs "${backend}_starting"
if grep -Eiq 'Traceback|CUDA out of memory|NCCL.*(error|failed)|EngineDeadError' \
"${RESULT_ROOT}/service/${backend}_starting_node"*.log; then
return 1
fi
fi
sleep 1
done
return 1
}
start_service() {
local backend="$1"
stop_service
log "starting TP32/EP4 backend=${backend} chunk=8192"
start_node 1 "$backend"
start_node 2 "$backend"
start_node 3 "$backend"
sleep 5
start_node 0 "$backend"
wait_health "$backend" || {
collect_service_logs "${backend}_startup_failed"
return 1
}
collect_service_logs "${backend}_healthy"
collect_gpu "${backend}_healthy"
}
run_bench() {
local backend="$1" repeat="$2"
local stem="${backend}_chunk8192_c8_r${repeat}"
local output="${RESULT_ROOT}/raw/${stem}.jsonl"
rm -f "$output"
log "bench backend=${backend} repeat=${repeat}/${REPEATS}"
sudo_host "${HEAD_HOST}" docker run --rm --network host \
-v "${MODEL_PATH}:${MODEL_PATH}:ro" \
-v "${RESULT_ROOT}:${RESULT_ROOT}" \
-e PYTHONUNBUFFERED=1 --entrypoint python3 "${PR_IMAGE}" \
-m sglang.benchmark.serving \
--backend sglang --host "${HEAD_HOST}" --port "${PORT}" \
--tokenizer "${MODEL_PATH}" --dataset-name random-ids \
--random-input-len "${INPUT_LEN}" --random-output-len "${OUTPUT_LEN}" \
--random-range-ratio 1.0 --num-prompts "${NUM_PROMPTS}" \
--max-concurrency "${CONCURRENCY}" --request-rate 10000 \
--warmup-requests "${WARMUP_REQUESTS}" --output-file "$output" \
--output-details --disable-tqdm \
>"${RESULT_ROOT}/bench/${stem}.log" 2>&1
python3 - "$output" "${NUM_PROMPTS}" "${OUTPUT_LEN}" <<'PY'
import json, sys
path, expected, expected_output_len = sys.argv[1], int(sys.argv[2]), int(sys.argv[3])
rows = [json.loads(x) for x in open(path, encoding="utf-8") if x.strip()]
assert len(rows) == 1, (path, len(rows))
assert rows[0].get("completed") == expected, rows[0].get("completed")
errors = rows[0].get("errors") or []
assert all(not error for error in errors), errors
output_lens = rows[0].get("output_lens") or []
assert len(output_lens) == expected, (len(output_lens), expected)
assert all(length == expected_output_len for length in output_lens), output_lens
PY
}
run_ab() {
local backend repeat
collect_gpu before
for backend in "${BACKENDS[@]}"; do
start_service "$backend"
for ((repeat = 1; repeat <= REPEATS; repeat++)); do
run_bench "$backend" "$repeat"
done
collect_service_logs "${backend}_completed"
verify_service_logs "${backend}_completed"
collect_gpu "${backend}_completed"
stop_service
sleep 5
done
collect_gpu after
}
summarize() {
python3 - "${RESULT_ROOT}" <<'PY'
import csv, json, re, statistics, sys
from pathlib import Path
root = Path(sys.argv[1])
pat = re.compile(r"(.+)_chunk8192_c8_r(\d+)\.jsonl$")
rows = []
for path in sorted((root / "raw").glob("*.jsonl")):
m = pat.match(path.name)
if not m:
continue
data = next(json.loads(x) for x in path.read_text().splitlines() if x.strip())
# 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"]
summary = []
for backend in ("marlin", "flashinfer_mxfp4"):
group = [x for x in rows if x["backend"] == backend]
if len(group) != 3:
raise SystemExit(f"expected 3 repeats for {backend}, got {len(group)}")
item = {"backend": backend, "repeats": len(group),
"completed_each": [x.get("completed") for x in group]}
for metric in metrics:
item[f"median_{metric}"] = statistics.median(float(x[metric]) for x in group)
summary.append(item)
idx = {x["backend"]: x for x in summary}
base, cand = idx["marlin"], idx["flashinfer_mxfp4"]
comparison = {
"input_throughput_change_pct":
(cand["median_input_throughput"] / base["median_input_throughput"] - 1) * 100,
"median_ttft_change_pct":
(cand["median_median_ttft_ms"] / base["median_median_ttft_ms"] - 1) * 100,
"p95_ttft_change_pct":
(cand["median_p95_ttft_ms"] / base["median_p95_ttft_ms"] - 1) * 100,
}
payload = {
"run_id": root.name,
"shape": {"input_len": 16384, "output_len": 1,
"concurrency": 8, "chunked_prefill_size": 8192,
"tp": 32, "ep": 4},
"summary": summary,
"comparison": comparison,
}
(root / "summary.json").write_text(json.dumps(payload, indent=2))
with (root / "results.csv").open("w", newline="") as f:
fields = ["backend", "repeat", "completed", *metrics]
w = csv.DictWriter(f, fieldnames=fields, extrasaction="ignore")
w.writeheader(); w.writerows(rows)
print(json.dumps(payload, indent=2))
PY
}
cleanup() {
collect_service_logs cleanup 2>/dev/null || true
stop_service 2>/dev/null || true
}
trap cleanup EXIT INT TERM
case "${ACTION}" in
build)
require_password; check_inputs; build_all_nodes; prewarm_all_nodes
;;
run)
require_password; check_inputs; run_ab; summarize
;;
all)
require_password; check_inputs; build_all_nodes; prewarm_all_nodes
run_ab; summarize
;;
summarize)
summarize
;;
stop)
require_password; stop_service
;;
*)
echo "Usage: $0 {all|build|run|summarize|stop}" >&2
exit 2
;;
esac