feat(agentx): retune Kimi-K3 FP4 MI355X ATOM DSpark recipe on _0907 - #2852
feat(agentx): retune Kimi-K3 FP4 MI355X ATOM DSpark recipe on _0907#2852zejunchen-zejun wants to merge 6 commits into
Conversation
|
Thanks for the contribution! Please reach out to respective companies' CODEOWNER to fill in the latest PR_REVIEW_CHECKLIST.md before pinging core maintainer on Slack for review. In order for the signoff PR check bot to trigger, you must follow the PR_REVIEW_CHECKLIST.md template correctly, including the phrase For PR verification, add the PR authors are responsible for ensuring that after merging, all GitHub Action jobs fully pass. A lot of the time, failures are just flakes and simply re-running the failed jobs will fix it. See GitHub's docs on re-running failed jobs 感谢你的贡献！请联系相应公司的 CODEOWNER 填写最新的 PR_REVIEW_CHECKLIST.md，然后再在 Slack 上联系核心维护者进行审阅。为了触发 signoff PR 检查机器人，你必须正确遵循 PR_REVIEW_CHECKLIST.md 模板，包括保留英文语句 如需进行 PR 验证，请为此 PR 添加 PR 作者有责任确保合并后所有 GitHub Action 任务完全通过。 很多时候失败只是偶发抖动（flake），重新运行失败的任务即可解决。参见 GitHub 关于重新运行失败任务的文档 |
Track ROCm/ATOM#2131, the checked-in Kimi-K3 AgentX recipe, onto image kimi_k3_agentic_0907 and publish the full validated concurrency set [1, 2, 4, 8, 12, 14, 16, 32, 40, 56, 64], replacing [1, 2, 4, 8, 12, 32, 40, 56]. Concurrency 14, 16, and 64 are new points and every existing point is retuned. Serving is three bands. Concurrency 1-4 is the latency floor: GPU-resident, dcp 1, 7 draft tokens at golden AL 3.84. Concurrency 8 through 16 shard the decode KV read with dcp 8, back the paged KV with the LMCache DRAM tier, run 3 draft tokens at golden AL 3.00, and rebuild the KDA recurrent state in place with ReplaySSM. Concurrency 32 and up serve without a draft model. Both acceptance lengths come from the committed golden curve in golden_al_distribution/kimik3_dspark_probabilistic_sample_method_block_rejection_sample_method.yaml (7 -> 3.84, 3 -> 3.00); evaluations drop the flag and use real acceptance. Full CUDA graphs are captured over the dense range [2 .. window x (1 + draft tokens)] via --level 3, --cudagraph-mode FULL, and an explicit --cudagraph-capture-sizes. The window is 2 x concurrency, and the (1 + draft) factor covers the verify step of a DSpark round, which submits one row per draft token on top of the accepted token. Concurrency 14 runs the concurrency 16 server verbatim, window pinned at 32, and only moves the client: deriving the window from 2 x concurrency there gives 28 and the server dies during graph warmup. The CPU state-offload tier is gone. ReplaySSM rebuilds the KDA state from the in-GPU checkpoint ring on every point that keeps a draft model, so OFFLOAD_STATE and its companions and the --state-checkpoint-slots override all drop out and the whole per-rank CPU budget goes to the paged KV. LMCACHE_MAX_LOCAL_CPU_SIZE is per rank, so the aggregate TOTAL_CPU_DRAM_GB is still divided by TP as the agentic README requires: dram-utilization 0.343 lands 128 GB/rank for concurrency 8 through 40, and 0.513 lands 192 GB/rank for concurrency 56 and 64. The eight ranks are pinned across the two sockets explicitly, and AITER_REUSE_IDENTICAL_COMM_GROUPS is set only at concurrency 56 and 64. Two failure modes the recipe documents are now closed in the launcher rather than left to the image. The Inferact/Kimi-K3-DSpark draft is staged into the shared Hugging Face cache before the server starts, because an uncached repo id makes every rank pull the same 7 GB checkpoint at once and the log simply stops after loading the drafter -- measured at ~0.7 MB/s on one cluster, about three hours, with every GPU at 0%. And the server registers under --served-model-name $MODEL, which is the name the AgentX client asks for on the wire. 将 MI355X 上 Kimi-K3 的 ATOM AgentX 提交对齐到 ROCm/ATOM#2131 中已合入的 recipe，镜像换成 kimi_k3_agentic_0907，并发点从 [1, 2, 4, 8, 12, 32, 40, 56] 换成 recipe 完整验证过的 [1, 2, 4, 8, 12, 14, 16, 32, 40, 56, 64]。14、16、64 为新增点，其余各点逐点重调。 服务分三档。并发 1-4 是时延下界：全部驻留 GPU，dcp 1，草稿 7 token 对应 golden AL 3.84。并发 8 到 16 用 dcp 8 把 decode 的 KV 读取分摊到 8 张卡， 由 LMCache DRAM 层承接分页 KV，草稿 3 token 对应 golden AL 3.00，并用 ReplaySSM 就地重建 KDA 循环状态。并发 32 及以上不加载草稿模型。两个接受长度 均取自仓库内已提交的 golden 曲线（7 -> 3.84，3 -> 3.00）；评测不传该参数， 使用真实接受率。 通过 --level 3、--cudagraph-mode FULL 和显式的 --cudagraph-capture-sizes， 在 [2 .. 窗口 x (1 + 草稿 token 数)] 这个稠密区间上抓取完整 CUDA graph。 窗口取 2 x 并发，(1 + 草稿) 这个系数覆盖 DSpark 一轮中的 verify 步——它在 被接受的 token 之上，每个草稿 token 再提交一行。并发 14 完整复用并发 16 的 server 配置（窗口固定为 32），只改客户端并发：在这里按 2 x 并发推出 28， server 会在 graph warmup 阶段直接挂掉。 CPU state offload 层被移除。凡是保留草稿模型的并发点，KDA 状态都由 ReplaySSM 从 GPU 内的 checkpoint ring 就地重建，因此 OFFLOAD_STATE 及其配套变量、 --state-checkpoint-slots 覆盖项全部去掉，每 rank 的 CPU 预算整块留给分页 KV。 LMCACHE_MAX_LOCAL_CPU_SIZE 仍是每 rank 设置，所以按 agentic README 的要求 把聚合预算 TOTAL_CPU_DRAM_GB 除以 TP：dram-utilization 0.343 在并发 8 到 40 给到每 rank 128 GB，0.513 在并发 56 和 64 给到每 rank 192 GB。八个 rank 按 显式映射绑定到两个 socket；AITER_REUSE_IDENTICAL_COMM_GROUPS 仅在并发 56 和 64 打开。 recipe 里记录的两个失败模式现在由脚本兜住，而不是依赖镜像。草稿模型 Inferact/Kimi-K3-DSpark 在起 server 之前先落进共享 HF 缓存：仓库 id 未命中 缓存时，八个 rank 会同时去拉同一份 7 GB 权重，日志在加载 drafter 之后就停住， 某个集群上实测约 0.7 MB/s、约三小时，期间所有 GPU 都是 0%。另外 server 现在 用 --served-model-name $MODEL 注册，这正是 AgentX 客户端在协议上请求的名字。 Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
84b11ce to
7c3ae97
Compare
seungrokj
added
AMD
full-sweep-enabled
agentx
|
Claude finished @zejunchen-zejun's task in 3m 2s —— View job Review
Summary: One 🟡 WARNING (inline): the PR description claims the launcher "records the triton version and warns when it is not 3.7.x", but no such check exists in the diff — since the silent triton 3.8.0 TTFT regression is the stated reason for the image swap, either add the guard or correct the description. Everything else verified clean: the |
![claude[bot]](https://avatars.githubusercontent.com/in/1236702?s=60&v=4){kind=small}
| # MODEL, MODEL_PATH, TP, DCP_SIZE, CONC, KV_OFFLOADING, KV_OFFLOAD_BACKEND, | ||
| # TOTAL_CPU_DRAM_GB, RESULT_DIR, DURATION, EP_SIZE, DP_ATTENTION | ||
|
|
||
| DRAFT_MODEL="Inferact/Kimi-K3-DSpark" |
There was a problem hiding this comment.
🟡 WARNING: The PR description says "the launcher now records the triton version and warns when it is not 3.7.x", but no such check exists anywhere in this diff.
Why it matters: The triton pin is the stated reason for the image swap — the 3.8.0 regression is silent, costs ~2.1x on prefill TTFT, and leaves nothing in the logs. Without the guard, the next image bump (e.g. an automated repin) can silently reintroduce it, and the recorded version is what would let anyone correlate a TTFT shift with the pin.
Fix: Add the check (or fix the PR description if the guard was intentionally dropped):
| DRAFT_MODEL="Inferact/Kimi-K3-DSpark" | |
| DRAFT_MODEL="Inferact/Kimi-K3-DSpark" | |
| # The 0903 image shipped triton 3.8.0, which silently costs ~2.1x on prefill | |
| # TTFT for this workload (conc 1 p50 TTFT ~0.79 s -> ~1.5 s); 0907 is back on | |
| # 3.7.x. Record the version so a future image bump that drifts off 3.7 shows | |
| # up in the log rather than staying silent. | |
| TRITON_VERSION="$(python -c 'import triton; print(triton.__version__)' 2>/dev/null || echo unknown)" | |
| echo "triton version: $TRITON_VERSION" | |
| if [[ "$TRITON_VERSION" != 3.7.* ]]; then | |
| echo "WARNING: triton $TRITON_VERSION is not 3.7.x; expect ~2.1x prefill TTFT regression on this workload" >&2 | |
| fi |
There was a problem hiding this comment.
I reviewed this PR and didn't find any bugs. Because it's a large, numerically dense retune of a production benchmark recipe (concurrency-band restructuring, CUDA-graph capture-size derivation, LMCache/NUMA/state-offload changes, and the matching config/changelog updates), a human look would still be worthwhile to sanity-check the tuning choices before the results are published.
What was reviewed: the three-band knob table and CUDA-graph capture-size math in kimik3_fp4_mi355x_atom_mtp.sh, the ReplaySSM/state-offload removal and LMCACHE_MAX_LOCAL_CPU_SIZE resizing, and the corresponding dram-utilization/conc-list restructuring in configs/amd-master.yaml. Checked that the perf-changelog.yaml entry is appended at the tail without altering prior bytes, and that the config's image tag matches the recipe's new kimi_k3_agentic_0907 image.
Extended reasoning...
Overview
This PR retunes a single-node MI355X AgentX benchmark recipe (Kimi-K3 FP4 via ATOM/DSpark): a new image tag, an expanded/retuned concurrency matrix, a reorganized three-band per-concurrency knob table in the launcher script, removal of the CPU state-offload tier in favor of ReplaySSM, new NUMA-pinning and AITER comm-group env vars, explicit CUDA-graph capture-size derivation, a retrying model-download step, and matching updates to configs/amd-master.yaml and perf-changelog.yaml.
Security risks
None identified. No new external inputs, no auth/crypto/permission-boundary code. The hf download retry loop and NUMA pinning are benchmark-infra concerns, not security-sensitive paths.
Level of scrutiny
Config/benchmark-tuning changes are normally low-scrutiny, but this one is unusually dense: it touches interacting numeric derivations (window = 2CONC, GRAPH_MAX = window(1+draft), per-rank CPU sizing via TOTAL_CPU_DRAM_GB/TP, dram-utilization aggregates) across two files that must stay consistent, plus a special-cased concurrency-14 point that intentionally overrides the general derivation. That combination of hand-computed constants and one deliberate exception is exactly the kind of change where a human's sanity check adds value even though the automated pass found nothing conclusive.
Other factors
I independently verified the things the bug hunter's ruled-out list flagged: the config's dram-utilization aggregates (1028 GB / 8 ≈ 128 GB/rank, 1538 GB / 8 ≈ 192 GB/rank) match the launcher's per-rank sizing comments, the conc-list splits in amd-master.yaml align with the script's case statement, and the perf-changelog entry is appended at the tail without disturbing prior bytes. No test suite covers these benchmark scripts (they run against real clusters), so review is the primary safety net here.
|
see unofficial run visualizer at https://inferencex.semianalysis.com/inference?unofficialRun=34105610075 |
|
see unofficial run visualizer at https://inferencex.semianalysis.com/inference?unofficialRun=34172490568 |
The kimi_k3_agentic_0907 image was built with podman on a host behind a local squid, and the build environment leaked into the image config. HTTP_PROXY, HTTPS_PROXY, http_proxy, and https_proxy are all baked in as http://127.0.0.1:3128, with NO_PROXY covering only localhost. The 0821 and 0903 images carried none of them. Nothing listens on that port inside the container on a benchmark node, so every outbound request the launcher makes before the server starts -- the uv bootstrap and aiperf install in install_agentic_deps, the trace-corpus fetch, and both hf downloads -- would reach a dead proxy. Clear the four variables and keep the loopback exemption, which leaves the AIPerf client and the /metrics scrape unaffected either way. kimi_k3_agentic_0907 镜像是在一台走本地 squid 的机器上用 podman 构建的， 构建环境泄漏进了镜像配置：HTTP_PROXY、HTTPS_PROXY、http_proxy、https_proxy 四个变量都被写死成 http://127.0.0.1:3128，而 NO_PROXY 只覆盖 localhost。 0821 和 0903 镜像都没有这些变量。 在 benchmark 节点上，容器内没有任何进程监听该端口，因此起 server 之前脚本发出的 所有外部请求——install_agentic_deps 里的 uv 引导与 aiperf 安装、trace 语料拉取、 以及两次 hf download——都会打到一个不存在的代理上。这里清掉这四个变量，并保留 回环地址的豁免，AIPerf 客户端与 /metrics 抓取不受影响。 Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Fold seungrokj's unset in e994743 into the block above. Same four variables, but the launcher's first outbound request is the hf download of the target checkpoint, which runs before rocm-smi and so before that unset; clearing them after it leaves that call pointed at the dead proxy. One unset, placed before every network call, plus the loopback NO_PROXY exemption. 把 e994743（seungrokj）中的 unset 合并进上面的代码块。变量完全相同，但脚本的 第一次外部请求是拉取目标 checkpoint 的 hf download，它在 rocm-smi 之前执行， 也就在原先那处 unset 之前；在其之后再清理，这一次调用仍会打到不存在的代理上。 现在只保留一处 unset，置于所有网络调用之前，并保留回环地址的 NO_PROXY 豁免。 Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
|
see unofficial run visualizer at https://inferencex.semianalysis.com/inference?unofficialRun=34174258644 |
|
see unofficial run visualizer at https://inferencex.semianalysis.com/inference?unofficialRun=34176729530 |
1 similar comment
|
see unofficial run visualizer at https://inferencex.semianalysis.com/inference?unofficialRun=34176729530 |
|
see unofficial run visualizer at https://inferencex.semianalysis.com/inference?unofficialRun=34176729530 |
2 participants
Track ROCm/ATOM#2131, the checked-in Kimi-K3 AgentX recipe, onto image kimi_k3_agentic_0907 and publish the full validated concurrency set [1, 2, 4, 8, 12, 14, 16, 32, 40, 56, 64], replacing [1, 2, 4, 8, 12, 32, 40, 56]. Concurrency 14, 16, and 64 are new points and every existing point is retuned.
The image matters for its triton pin, not only its ATOM revision. The 0903 build shipped triton 3.8.0, which costs roughly 2.1x on prefill TTFT for this workload -- concurrency 1 p50 TTFT goes from ~0.79 s to ~1.5 s -- with decode untouched, and neither the aiter revision nor the ATOM checkout nor the build flavour recovers it. 0907 is back on 3.7.x. The regression is silent, so the launcher now records the triton version and warns when it is not 3.7.x.
Serving is three bands. Concurrency 1-4 is the latency floor: GPU-resident, dcp 1, 7 draft tokens at golden AL 3.84. Concurrency 8 through 16 shard the decode KV read with dcp 8, back the paged KV with the LMCache DRAM tier, run 3 draft tokens at golden AL 3.00, and rebuild the KDA recurrent state in place with ReplaySSM. Concurrency 32 and up serve without a draft model. Both acceptance lengths come from the committed golden curve in golden_al_distribution/kimik3_dspark_probabilistic_sample_method_block_rejection_sample_method.yaml (7 -> 3.84, 3 -> 3.00); evaluations drop the flag and use real acceptance.
Full CUDA graphs are captured over the dense range [2 .. window x (1 + draft tokens)] via --level 3, --cudagraph-mode FULL, and an explicit --cudagraph-capture-sizes. The window is 2 x concurrency, and the (1 + draft) factor covers the verify step of a DSpark round, which submits one row per draft token on top of the accepted token. Concurrency 14 runs the concurrency 16 server verbatim, window pinned at 32, and only moves the client: deriving the window from 2 x concurrency there gives 28 and the server dies during graph warmup.
The CPU state-offload tier is gone. ReplaySSM rebuilds the KDA state from the in-GPU checkpoint ring on every point that keeps a draft model, so OFFLOAD_STATE and its companions and the --state-checkpoint-slots override all drop out and the whole per-rank CPU budget goes to the paged KV. LMCACHE_MAX_LOCAL_CPU_SIZE is per rank, so the aggregate TOTAL_CPU_DRAM_GB is still divided by TP as the agentic README requires: dram-utilization 0.343 lands 128 GB/rank for concurrency 8 through 40, and 0.513 lands 192 GB/rank for concurrency 56 and 64.
Two failure modes the recipe documents are now closed in the launcher rather than left to the image. The Inferact/Kimi-K3-DSpark draft is staged into the shared Hugging Face cache before the server starts, because an uncached repo id makes every rank pull the same 7 GB checkpoint at once and the log simply stops after loading the drafter -- measured at ~0.7 MB/s on one cluster, about three hours, with every GPU at 0%. And the server registers under --served-model-name $MODEL, which is the name the AgentX client asks for on the wire.
将 MI355X 上 Kimi-K3 的 ATOM AgentX 提交对齐到 ROCm/ATOM#2131 中已合入的 recipe，镜像换成 kimi_k3_agentic_0907，并发点从 [1, 2, 4, 8, 12, 32, 40, 56] 换成 recipe 完整验证过的 [1, 2, 4, 8, 12, 14, 16, 32, 40, 56, 64]。14、16、64 为新增点，其余各点逐点重调。
换镜像的关键不只是 ATOM 版本，还有 triton 的版本。0903 镜像带的是 triton
3.8.0，在该负载上 prefill TTFT 大约变差 2.1 倍——并发 1 的 p50 TTFT 从约 0.79 秒涨到约 1.5 秒，而 decode 完全不受影响；换 aiter、换 ATOM checkout、 换构建方式都救不回来。0907 回到 3.7.x。这个退化是静默的，因此脚本现在会把
triton 版本打进日志，并在不是 3.7.x 时告警。
服务分三档。并发 1-4 是时延下界：全部驻留 GPU，dcp 1，草稿 7 token 对应
golden AL 3.84。并发 8 到 16 用 dcp 8 把 decode 的 KV 读取分摊到 8 张卡， 由 LMCache DRAM 层承接分页 KV，草稿 3 token 对应 golden AL 3.00，并用 ReplaySSM 就地重建 KDA 循环状态。并发 32 及以上不加载草稿模型。两个接受长度
均取自仓库内已提交的 golden 曲线（7 -> 3.84，3 -> 3.00）；评测不传该参数， 使用真实接受率。
通过 --level 3、--cudagraph-mode FULL 和显式的 --cudagraph-capture-sizes， 在 [2 .. 窗口 x (1 + 草稿 token 数)] 这个稠密区间上抓取完整 CUDA graph。 窗口取 2 x 并发，(1 + 草稿) 这个系数覆盖 DSpark 一轮中的 verify 步——它在 被接受的 token 之上，每个草稿 token 再提交一行。并发 14 完整复用并发 16 的
server 配置（窗口固定为 32），只改客户端并发：在这里按 2 x 并发推出 28，
server 会在 graph warmup 阶段直接挂掉。
CPU state offload 层被移除。凡是保留草稿模型的并发点，KDA 状态都由 ReplaySSM 从 GPU 内的 checkpoint ring 就地重建，因此 OFFLOAD_STATE 及其配套变量、 --state-checkpoint-slots 覆盖项全部去掉，每 rank 的 CPU 预算整块留给分页 KV。 LMCACHE_MAX_LOCAL_CPU_SIZE 仍是每 rank 设置，所以按 agentic README 的要求 把聚合预算 TOTAL_CPU_DRAM_GB 除以 TP：dram-utilization 0.343 在并发 8 到 40 给到每 rank 128 GB，0.513 在并发 56 和 64 给到每 rank 192 GB。
recipe 里记录的两个失败模式现在由脚本兜住，而不是依赖镜像。草稿模型
Inferact/Kimi-K3-DSpark 在起 server 之前先落进共享 HF 缓存：仓库 id 未命中 缓存时，八个 rank 会同时去拉同一份 7 GB 权重，日志在加载 drafter 之后就停住，
某个集群上实测约 0.7 MB/s、约三小时，期间所有 GPU 都是 0%。另外 server 现在 用 --served-model-name $MODEL 注册，这正是 AgentX 客户端在协议上请求的名字。
Note
Medium Risk
Changes benchmark-representative serving recipes (KV offload sizing, CUDA graphs, spec decode, and concurrency matrix), so published AgentX numbers shift; risk is operational/tuning rather than auth or data handling.
Overview
Retunes the Kimi-K3 MXFP4 MI355X ATOM AgentX submission to image
kimi_k3_agentic_0907and expands the published concurrency set to [1, 2, 4, 8, 12, 14, 16, 32, 40, 56, 64] (adds 14, 16, 64; retunes the rest).amd-master.yamlnow uses two dram-utilization blocks (0.343 → 128 GB/rank LMCache, 0.513 → 192 GB/rank for 56/64) and documents three serving bands (interactive / mid / throughput).The launcher
kimik3_fp4_mi355x_atom_mtp.shis reworked to match the recipe: per-concurrency ReplaySSM, draft depth, batch limits, and full CUDA graphs (--level 3,--cudagraph-mode FULL, explicit capture sizes throughwindow × (1 + draft tokens)), including CONC=14 reusing the CONC=16 server with CUDA graph width pinned at 32. CPU KDA state offload and--state-checkpoint-slotsare removed; LMCache gets the full per-rank CPU budget, plus NUMA pinning and retained hybrid KV offload flags.Operational fixes in the script: unset leaked HTTP proxies from the 0907 image, prefetch the DSpark draft with retries before server start, and register
--served-model-name $MODELso AgentX warmup does not 404.perf-changelog.yamlrecords the refresh.Reviewed by Cursor Bugbot for commit 058a581. Bugbot is set up for automated code reviews on this repo. Configure here.