Compare commits
10 Commits
a9206ff105
...
0fdcab9927
| Author | SHA1 | Date | |
|---|---|---|---|
|
|
0fdcab9927 | ||
|
|
a5248ed80e | ||
|
|
fc336a3c7b | ||
|
|
63f2327a90 | ||
|
|
a90c898683 | ||
|
|
ab9a5422f6 | ||
|
|
7f67dfe6b3 | ||
|
|
e8ff3ce1e8 | ||
|
|
d28db48e4b | ||
|
|
39f692caae |
@ -1,5 +1,11 @@
|
||||
# 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/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`。
|
||||
|
||||
**更新(2026-08-18 22:21:40 CST)**:完成 Kimi-K3 SM120 FlashInfer MXFP4 去重审计。确认 CUTLASS SiTU kernel 已由未合并的 FlashInfer PR #4460 实现,不再提交平行 kernel;SGLang 贡献收敛为 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`。
|
||||
|
||||
@ -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`。
|
||||
|
||||
## 验收矩阵
|
||||
@ -11,18 +11,21 @@
|
||||
| 区分官方 0.6.17 与实际能力 | 已证明 | 605 profile 设置 `FLASHINFER_DISABLE_VERSION_CHECK=1` 并启动前执行两个补丁;报告未再声称 0.6.17 原版直接支持完整组合 |
|
||||
| 逐文件比较 605 补丁、FlashInfer #4460、SGLang 上游 | 已证明 | `UPSTREAM_DUPLICATION_AUDIT.md` 的“上游已有/605 已有/我们新增/明确排除”四部分;605 补丁目标已从原文件重新核验 |
|
||||
| 不提交重复 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 分支 |
|
||||
| SiTU 4.0/25.0 参数透传 | 已证明 | #4460 ABI 的 SM120 GPU correctness 测试覆盖独立 `situ_beta` 与 `situ_linear_beta` |
|
||||
| 非连续 hidden states | 已证明 | GPU correctness 测试输入为非连续 view,runner 在 MXFP8 quantize 边界 materialize contiguous tensor |
|
||||
| attn_residual SM120 guard | 已证明 | 注册架构测试覆盖 SM100/103/110 为 true、SM120/SM90 为 false |
|
||||
| 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` |
|
||||
| 最终兼容回归 | 已通过 | 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` |
|
||||
| 保留 EP4 服务级结果 | 已证明 | 601 原始目录有 24 个正式 JSONL 点;每组 3 repeats、每点 40/40 成功;错误扫描为空;未重跑 EP32 |
|
||||
| 公开 PR 标题、正文、依赖和已知限制 | 已完成 | `SGLANG_DRAFT_PR.md` 遵循 SGLang 官方模板;标题为 `[Feature] Support Kimi-K3 FlashInfer MXFP4 MoE on SM120` |
|
||||
| GitHub fork 交付 | 已完成 | `ACEEE-1222/sglang:hzy/kimi-k3-sm120-flashinfer-mxfp4`,Draft HEAD `300c87a` |
|
||||
| 当前 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 `b95b534`,远端 ref 已复核 |
|
||||
| 内部仓库独立分支 | 本地和 601 已完成,远端待推送 | 601 分支 `hzy/feature-kimi-k3-sm120-flashinfer-mxfp4` 已到审计提交;`git.meta-stone.net` 尚未为本机或 601 key 授权,HTTPS 也无 Token |
|
||||
|
||||
## 原始证据位置
|
||||
@ -38,6 +41,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
|
||||
```
|
||||
|
||||
## 唯一未关闭项
|
||||
|
||||
@ -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}"
|
||||
@ -7,15 +7,25 @@
|
||||
> Kimi-K3 SGLang 接入、精确 MXFP8×MXFP4+SiTU 测试和 6000D 实测证据。
|
||||
> 权威边界见 [`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 正文见
|
||||
[`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;最终非 JIT 回归为
|
||||
`8 passed, 5 deselected, 17 warnings in 22.23s`。
|
||||
五个变更文件已通过完整 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,并形成可回归、可性能比较、可提交上游的最小改动。
|
||||
@ -704,3 +714,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
|
||||
```
|
||||
|
||||
@ -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
|
||||
FlashInfer CUTLASS API.
|
||||
|
||||
After the final compatibility and CI-registration changes, the non-JIT subset
|
||||
was rerun from the final commit:
|
||||
After syncing the fork and rebasing the seven-patch series onto SGLang main
|
||||
`593b1a9b8aeedb908ee21870c02a0fd1a3dc30a5`, the non-JIT subset was rerun
|
||||
from the final Draft head `b95b5341ab9555ba5ea53f5828f6de203d749cce`:
|
||||
|
||||
```bash
|
||||
pytest -q -s \
|
||||
@ -75,38 +76,51 @@ pytest -q -s \
|
||||
```
|
||||
|
||||
```text
|
||||
8 passed, 5 deselected, 17 warnings in 22.23s
|
||||
10 passed, 3 deselected, 17 warnings in 33.30s
|
||||
```
|
||||
|
||||
## Speed Tests and Profiling
|
||||
|
||||
Four-node serving validation used 32 RTX 6000D GPUs with TP32/EP4/DP1,
|
||||
16K input tokens, one output token, concurrency 8/16, and chunked Prefill
|
||||
sizes 8K/16K. Each point was repeated three times; all 24 runs completed with
|
||||
40/40 successful requests and no OOM, Traceback, NCCL error, or
|
||||
`EngineDeadError`.
|
||||
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.
|
||||
|
||||
| Chunk | C | Backend | Input TPS | TTFT P50 (s) | FlashInfer vs Marlin |
|
||||
|---:|---:|---|---:|---:|---:|
|
||||
| 8K | 8 | Marlin | 2935.02 | 43.51 | baseline |
|
||||
| 8K | 8 | FlashInfer | 3257.96 | 39.19 | +11.00% TPS, -9.92% TTFT |
|
||||
| 8K | 16 | Marlin | 2934.90 | 87.09 | baseline |
|
||||
| 8K | 16 | FlashInfer | 3260.14 | 78.39 | +11.08% TPS, -9.99% TTFT |
|
||||
| 16K | 8 | Marlin | 2812.91 | 45.47 | baseline |
|
||||
| 16K | 8 | FlashInfer | 3027.50 | 42.27 | +7.63% TPS, -7.04% TTFT |
|
||||
| 16K | 16 | Marlin | 2812.33 | 90.97 | baseline |
|
||||
| 16K | 16 | FlashInfer | 3027.34 | 84.52 | +7.65% TPS, -7.09% TTFT |
|
||||
```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 patch series used for the run: 300c87a431ac40d3e7817246376b7fe20932db09
|
||||
Final rebased Draft head: b95b5341ab9555ba5ea53f5828f6de203d749cce
|
||||
Final SGLang main base: 593b1a9b8aeedb908ee21870c02a0fd1a3dc30a5
|
||||
FlashInfer: PR #4460 implementation b525c513fef7c1d162efd3ce8636fcc8760de66a
|
||||
Repeats: 3 per backend
|
||||
```
|
||||
|
||||
The service benchmark used the previously validated candidate integration
|
||||
with the same Kimi layout, SiTU, non-contiguous-input, and SM120 guard fixes.
|
||||
The final public #4460 API is covered by the accuracy test above; no claim is
|
||||
made that an unreleased #4460 wheel was used for this four-node benchmark.
|
||||
The service benchmark was collected before the fork sync. `git range-diff`
|
||||
reports all seven commits as patch-equivalent after the rebase, and the
|
||||
directed tests above were rerun from the final Draft head. The performance
|
||||
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
|
||||
|
||||
- [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] 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] Follow the SGLang code style [guidance](https://docs.sglang.io/developer_guide/contribution_guide.html#code-style-guidance).
|
||||
|
||||
|
||||
@ -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` |
|
||||
|
||||
## 明确排除的重复内容
|
||||
|
||||
@ -102,11 +103,15 @@ activation_situ_linear_beta = 25.0
|
||||
SGLang Draft 的完整可审查 diff:
|
||||
|
||||
```text
|
||||
base: 63d783bbe0955237ec41f9ddabf7235ddf04673c
|
||||
head: 300c87a431ac40d3e7817246376b7fe20932db09
|
||||
base: 593b1a9b8aeedb908ee21870c02a0fd1a3dc30a5
|
||||
head: b95b5341ab9555ba5ea53f5828f6de203d749cce
|
||||
patch: patches/sglang-kimi-k3-sm120-flashinfer-mxfp4-draft.patch
|
||||
```
|
||||
|
||||
该 7 提交系列由旧范围 `63d783b..300c87a` 重放到用户新同步的 main;
|
||||
`git range-diff` 逐项均为 `=`。在新 base 的独立 worktree 应用上述 patch、
|
||||
暂存全部文件后,与 `b95b534` 比较的 diff 为空。
|
||||
|
||||
GitHub 交付分支:
|
||||
|
||||
```text
|
||||
@ -153,15 +158,21 @@ GPU: NVIDIA RTX 6000D, SM120
|
||||
|
||||
最终审查发现普通 SwiGLU 路径也曾无条件访问 `ActivationType.Situ`,可能误伤
|
||||
旧 FlashInfer。提交 `775aa076` 改为动态加入 SiTU enum,并新增旧版兼容回归。
|
||||
最终提交 `300c87a` 为两个测试文件增加正式 CUDA CI 注册,并通过完整
|
||||
pre-commit。最终 HEAD 的能力探测、旧版兼容和架构 gate 为
|
||||
`8 passed, 5 deselected, 17 warnings in 22.23s`;GPU kernel 路径继续由前述
|
||||
8/8 测试覆盖。
|
||||
同步前提交 `300c87a` 为两个测试文件增加正式 CUDA CI 注册,并通过完整
|
||||
pre-commit。同步后的最终 HEAD `b95b534` 在 601 通过 compileall;能力探测、
|
||||
旧版兼容和架构 gate 为 `10 passed, 3 deselected, 17 warnings in 33.30s`;
|
||||
GPU kernel 路径继续由前述 8/8 测试覆盖。
|
||||
|
||||
完整 pre-commit 命令覆盖全部五个变更文件;isort、ruff、black、codespell、
|
||||
registered-tests 检查及其他 hook 全部通过。下载 GitHub hook 时只对该命令使用
|
||||
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` 仍被视为启用、源码包缺少
|
||||
生成的 `data/csrc`;第三次暴露只读源码不能生成 `_build_meta.py`。这些都是
|
||||
测试装配问题,最终 v4 采用可写临时源码树构建 wheel 后通过,并非功能回归。
|
||||
|
||||
@ -2,6 +2,38 @@
|
||||
|
||||
本文件记录可恢复的关键决策、失败尝试和证据路径。成功结论见 `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 代表点
|
||||
|
||||
- 只补跑一个代表点,不重跑 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,6 +67,8 @@
|
||||
`/data/hzy/phase5_sglang_pr4460_draft_cpu_final.log`。
|
||||
- 公开 Draft PR 正文遵循 SGLang 官方模板,不暴露内网原始路径;原始日志路径、
|
||||
镜像 SHA、wheel SHA 和去重证据继续保留在本仓库审计文件中。
|
||||
- Checklist 在最终代表点和内部档案完成后全部勾选。公开 PR 仅报告一个固定
|
||||
shape 的目标硬件结果,不做普遍提速声明;用户文档仍等待 #4460 发布稳定 API。
|
||||
|
||||
## 2026-08-18 Phase 5
|
||||
|
||||
|
||||
@ -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
|
||||
|
@ -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
|
||||
}
|
||||
}
|
||||
@ -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
|
||||
Loading…
x
Reference in New Issue
Block a user