perf: MI355X DSv4.1 Flash graph capture sized to concurrency, sweep extended to 128 / MI355X DSv4.1 Flash 按并发设定 graph capture 规模并将并发扫描扩展至 128 - #3068
Conversation
The deepseekv41-flash-0909 tag predates vllm-project/vllm#56503, which moves the mHC delayed pre block off the eager Torch reference and onto AITER. That block is 85% of the decoder's kernel launches for this model, so the release tag leaves most of the decode cost unaddressed. Tag is nightly-TBD until the first nightly containing that commit is published; it is then pinned to nightly-<sha> like the other ROCm entries.
AITER's Triton MoE GEMM logs "GLUON backend not available. Using TRITON backend!!!" on every call. Gluon supports only gfx1250, so on gfx950 the message reports a fixed property of the hardware rather than a condition, and it dominates the server log: 410,984 of 417,504 lines in a gsm8k run, 29.8 MiB of 32 MiB. Set AITER_TRITON_LOG_LEVEL=ERROR, the knob aiter.ops.triton.utils.logger reads. Every warning that module emits is about Gluon availability -- the two Gluon GEMM depth fallbacks are inside gfx1250-only paths, and fused_mxfp4_quant warns only when Gluon is requested explicitly -- so raising the threshold hides nothing actionable on this architecture. Log hygiene rather than performance: the emits run at 61/s per worker and 4.33 us each, or 0.026% of wall time. Signed-off-by: Fangzhou Ai <fangzhou@semianalysis.ai> Co-authored-by: Cursor <cursoragent@cursor.com>
The 2*CONC cap sat below AgentX's subagent fan-out. At CONC=1 the engine admitted 2 requests and queued the rest on scheduling capacity, with vLLM reporting queue times up to 16.2 s. Removing the cap takes queue time to zero at identical output.
The nightly containing vllm-project/vllm#56503 is out: nightly-eed1f3d0c6043bd494424a22443ee198dd56f657 (sha256:960228cf..., published 2026-09-12). Replace the nightly-TBD placeholder in the config, recipe header, and both docs, and record the upstream recipe PR that moves vllm-project/recipes to the same image. Co-authored-by: Cursor <cursoragent@cursor.com>
The recipe pinned --moe-backend aiter_triton_mxfp4_bf16, which names a single W4A16 entry and dispatched _moe_gemm_a16w4. Passing plain aiter opens vLLM's priority list, whose head is the Composable Kernel backend, and CK quantizes activations to FP8 internally: the decode trace switches to mfma_moe1_silu_mul_afp8_wfp4_bf16 and mfma_moe2_afp8_wfp4_bf16, the same a8w4 expert family the DSV4-Pro MI355X recipe already runs. The BF16 in the backend name and this checkpoint's activation_scheme=dynamic both suggest otherwise, hence the comment at the flag. Measured on one MI355X half, TP4, 131k context, CUDA graphs on, both arms captured on the same four GPUs under identical profiling overhead: conc MoE GEMM ms/step whole step ms 1 2.72 -> 2.30 19.11 -> 18.31 (4.2%) 4 4.91 -> 4.27 19.42 -> 18.40 (5.3%) 16 9.06 -> 7.18 28.38 -> 25.92 (8.7%) AgentX replay at conc 1, each arm solo on a quiet node, improves mean ITL from 4.01 to 3.91 ms and mean TTFT from 681 to 639 ms. gsm8k is unchanged: 0.9719 +/- 0.0045 strict-match, 0.9712 +/- 0.0046 flexible, identical to the Triton arm. Signed-off-by: Fangzhou Ai <31551580+Fangzhou-Ai@users.noreply.github.com> Co-authored-by: Cursor <cursoragent@cursor.com>
![claude[bot]](https://avatars.githubusercontent.com/in/1236702?s=60&v=4){kind=small}
There was a problem hiding this comment.
Nothing blocking. The comments below are optional suggestions. There is no need to push a fix for them before merging.
Beyond the inline finding, I checked whether pinning MAX_NUM_SEQS to 128 (dropping the 2*CONC fan-out headroom used by dsv4_fp4_mi355x_vllm_mtp.sh) starves the new conc=64/128 points — it doesn't look like a regression: dsv41flash_fp4_b300_vllm_mtp.sh never passes --max-num-seqs at all and relies on the same vLLM engine default of 128, so the scheduler cap is now identical across the two arms as the PR claims, not something unique to MI355X.
Extended reasoning...
This run's inline finding (missing perf-changelog coverage for the image bump, MAX_NUM_SEQS change, and CAPTURE_SIZE derivation switch, per AGENTS.md's changelog requirement) already flags something a human should confirm before merge, so no separate top-level verdict is warranted beyond a brief note on an adjacent concern I checked and ruled out: whether removing 2x scheduler headroom for AgentX subagent fan-out at the two newly-added top-of-sweep concurrencies (64, 128) leaves MI355X worse off than its CUDA siblings. Reading dsv41flash_fp4_b300_vllm_mtp.sh shows it never sets --max-num-seqs at all, so it already runs at the vLLM engine default of 128 with no fan-out headroom either — confirming the PR description's claim that pinning MI355X to 128 makes the scheduler cap identical across arms rather than introducing a MI355X-specific regression.
| - config-keys: | ||
| - glm5.2-fp4-mi355x-sglang-agentic-mtp | ||
| scenario-type: | ||
| - agentic-coding | ||
| description: | ||
| - "Update the SGLang ROCm image from lmsysorg/sglang-rocm:v0.5.16-rocm720-mi35x-20260728 to lmsysorg/sglang-rocm:v0.5.19-rocm720-mi35x-20260910." | ||
| - "Pick up our recent SGLang main-branch optimizations for GLM-5.2-MXFP4 serving." | ||
| - "Lower the TP HiCache host-pool ratio from 1.5 to 1.0 (DP-attention arm stays at 0.5). Eval and sweep still share the same ratio-based sizing; the smaller pool cuts pinned host DRAM on the TP4 arm (~453 GB vs ~679 GB). Keep the write policy at write_through: evals-only run 34569310922 passed with ratio 1.0 and write_through, so write_through_selective was unrelated to the earlier failure." | ||
| - "Remove the SGLANG_HEALTH_CHECK_TIMEOUT=120 override and its stale first-request JIT explanation; successful evals-only run 34569310922 used the default health-check timeout." | ||
| - "Replace the removed decode-only alias --cuda-graph-max-bs with its canonical name --cuda-graph-max-bs-decode, retaining the same min(2*CONC, 64) cap. The 20260910 image includes sgl-project/sglang#38375, so argparse interprets the removed spelling as an ambiguous prefix of the decode and prefill options and exits before serving. Leave --cuda-graph-max-bs-prefill unset: it controls a separate prefill token-capture ceiling, not request concurrency." | ||
| - "Switch DSA prefill and decode from tilelang to triton (sgl-project/sglang#30575 Fast Triton Sparse MLA). The 20260910 image already contains that backend; kv-cache-dtype stays fp8_e4m3, which the Triton DSA path requires." | ||
| pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2900 | ||
|
|
||
| - config-keys: | ||
| - glm5.1-fp8-b200-tilert | ||
| scenario-type: | ||
| - fixed-seq-len | ||
| description: | ||
| - Require native power only for TileRT 8k1k; keep native launches inside the owned | ||
| allocation and retain both role exit statuses. | ||
| - 仅对 TileRT 8k1k 要求原生功耗；原生启动始终使用本任务分配，并保留两个角色退出状态。 | ||
| pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/3057 | ||
|
|
||
| - config-keys: | ||
| - glm5.1-fp8-b200-tilert | ||
| description: | ||
| - "Preserve an explicit prepared GLM-5.1 model path and mount shared HF cache blobs in native TileRT containers." | ||
| - "保留显式指定的 GLM-5.1 预备模型路径，并在原生 TileRT 容器中挂载共享 HF 缓存 blobs。" | ||
| pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/3057 | ||
|
|
||
| - config-keys: | ||
| - glm5.2-fp4-gb200-dynamo-sglang-agentic-agg | ||
| scenario-type: | ||
| - agentic-coding | ||
| - dsv41flash-fp4-mi355x-vllm-agentic-dspark | ||
| description: | ||
| - "Collect shared measured GPU power for the existing two-node GB200 GLM-5.2 aggregate AgentX recipe with producer 80d7203e424f903c9017de4608ee2044afce9574, exact selected-concurrency windows, and post-job adapter validation. Preserve native status and invalid diagnostics before failure. Ordinary Slurm limits, serving settings, other AgentX recipes, and fixed-sequence producer pins are unchanged." | ||
| - "为现有双节点 GB200 GLM-5.2 聚合 AgentX 配方接入共享 GPU 实测功耗，按所选并发绑定测量窗口，并在任务结束后通过适配器校验。返回失败前保留原生状态及无效诊断；普通 Slurm 时限、服务配置、其他 AgentX 配方和固定序列 producer 版本保持不变。" | ||
| pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2928 | ||
| - "Switch --moe-backend from aiter_triton_mxfp4_bf16 to aiter, which selects the CK a8w4 experts (mfma_moe1_silu_mul_afp8_wfp4_bf16 / mfma_moe2_afp8_wfp4_bf16) instead of the W4A16 _moe_gemm_a16w4 Triton kernel, matching the DSV4-Pro MI355X recipe. Decode traces at 131k context show MoE GEMM falling from 2.72 to 2.30 ms/step at conc 1, 4.91 to 4.27 at conc 4 and 9.06 to 7.18 at conc 16, for whole-step gains of 4.2%, 5.3% and 8.7%. AgentX conc 1 improves mean ITL 4.01 to 3.91 ms and mean TTFT 681 to 639 ms; gsm8k is unchanged at 0.9719 strict-match." | ||
| - "将 --moe-backend 从 aiter_triton_mxfp4_bf16 改为 aiter，从而选择 CK a8w4 专家内核（mfma_moe1_silu_mul_afp8_wfp4_bf16 / mfma_moe2_afp8_wfp4_bf16），取代 W4A16 的 _moe_gemm_a16w4 Triton 内核，与 DSV4-Pro MI355X 配方保持一致。131k 上下文的解码 trace 显示 MoE GEMM 在并发 1 时从 2.72 降至 2.30 ms/step，并发 4 从 4.91 降至 4.27，并发 16 从 9.06 降至 7.18，整步分别加速 4.2%、5.3% 和 8.7%。AgentX 并发 1 的平均 ITL 从 4.01 降至 3.91 ms，平均 TTFT 从 681 降至 639 ms；gsm8k strict-match 保持 0.9719 不变。" | ||
| pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/3058 |
There was a problem hiding this comment.
🟡 (optional) AGENTS.md requires a perf-changelog entry for every performance-affecting change and every recipe modification, but this diff only logs the --moe-backend switch. The image bump to the nightly tag, the MAX_NUM_SEQS change from 2*CONC to 128, the CAPTURE_SIZE derivation switch from MAX_NUM_SEQS to CONC (which changes KV-cache memory and max concurrency per the PR's own numbers), and the conc-list extension to 64/128 are all undocumented. Fix: add one perf-changelog.yaml entry per remaining change, matching the one-entry-per-change convention used elsewhere in the file.
Extended reasoning...
AGENTS.md line 32 states: 'Every change that can affect benchmark performance and every recipe addition or modification requires a new perf-changelog.yaml entry.' The diff to benchmarks/single_node/agentic/dsv41flash_fp4_mi355x_vllm_mtp.sh changes MAX_NUM_SEQS (2CONC -> 128), the CAPTURE_SIZE loop bound (MAX_NUM_SEQS(1+NUM_SPEC_TOKENS) -> CONC*(1+NUM_SPEC_TOKENS)), and configs/amd-master.yaml changes the image tag and extends conc-list to [64, 128] -- all of which the PR description itself claims materially affect KV-cache memory, queue times, and comparability with B300. Only the moe-backend change got a perf-changelog.yaml entry (added at line 7432-7437); readers relying on perf-changelog.yaml as the source of truth for what changed and why will miss the graph-capture-sizing fix, the image bump rationale, and the sweep extension, breaking the changelog's function as an audit trail for benchmark-affecting changes.
Verification: normal (repo non-negotiable invariant, introduced by this PR). AGENTS.md: "Every change that can affect benchmark performance and every recipe addition or modification requires a new perf-changelog.yaml entry." The net diff fa902f1..HEAD adds exactly ONE changelog entry (perf-changelog.yaml lines 7431-7437, pr-link #3058) documenting only the --moe-backend switch. The same diff makes multiple…
4fd3bec to
04b355a
Compare
…to 128 Deriving CAPTURE_SIZE from MAX_NUM_SEQS pinned --max-cudagraph-capture-size at 1024 for every point in the sweep. At CONC=1 the engine never sees a decode batch larger than 6, yet it still captured graphs out to 1024: 18.71 GiB of PIECEWISE plus 19.40 GiB of FULL capture memory per GPU, roughly 30% of what is left after weights, spent on shapes the workload cannot reach. This is hygiene, not a throughput fix. Reclaiming that memory grows the KV pool from 33.6M tokens, but KV was never the constraint at these points: peak occupancy was 1.6% at conc 1 and 13.0% at conc 32, with the prefix cache hit rate at 99.7%, so nothing was evicted and nothing refilled. Extrapolating the observed doubling, conc 128 would have sat near half of even the old pool, so the ladder could have been extended without this change. The reasons to make it are that capture should describe the point being measured, that 38.1 GiB per GPU and the startup time to capture 1024-wide shapes buy nothing, and that the headroom at the new points is then not something to reason about. Derive it from the workload: each of the CONC lanes submits one real token plus NUM_SPEC_TOKENS drafts per step, so the largest decode batch is CONC * (1 + NUM_SPEC_TOKENS), rounded up to a power of two. That gives 8 through 1024 across the ladder rather than a flat 1024, and matches the derivation the other DeepSeek-V4.1-Flash agentic recipes already use. MAX_NUM_SEQS stays at the upstream default of 128, so the scheduler cap is unchanged. Extend conc-list to [1, 2, 4, 8, 16, 32, 64, 128]. The published sweep is still climbing steeply at the old top of the ladder: output throughput per GPU rises 62% from conc 16 to 32 with no sign of a knee, energy per output token is still improving at 3.81 -> 2.60 J, TTFT is flat near 0.9 s and KV occupancy peaked at 13%. The sweep ended before it found the throughput plateau, which is the region that determines serving cost. Also correct the Engram note: passing --engram-config on ROCm fails a strict current_platform.is_cuda() check in EngramConfig.verify_model_config, but omitting it is not a downgrade, because validation is skipped entirely when the flag is absent and the Engram module then defaults to cpu_offload=True. Every MI355X run already logs the tables in pinned host memory at 23.60 GiB per rank per layer.
…ncy ladder Co-authored-by: Cursor <cursoragent@cursor.com>
04b355a to
6b7f2cd
Compare
|
see unofficial run visualizer at https://inferencex.semianalysis.com/inference?unofficialRun=34736117923 |
2 participants
Supersedes #3058 and targets
maindirectly, so the MI355X DSv4.1 Flash recipe lands in one self-consistent state. All of #3058's commits are carried here unchanged; this PR adds two on top. #3058 (and #3035 behind it) can be closed once this merges.Carried from #3058:
configs/amd-master.yaml: imagedeepseekv41-flash-0909→nightly-eed1f3d0c6043bd494424a22443ee198dd56f657, matching [DeepSeek] V4.1-Flash: point the AMD image at a ROCm nightly and select the CK a8w4 MoE experts vllm-project/recipes#962AITER_TRITON_LOG_LEVEL=ERROR, which removed 98% of the server log with no throughput effectMAX_NUM_SEQSfrom2 * CONCto the upstream default of 128, which took queue times at CONC=1 from 16.2 s to zero--moe-backend aiterrather thanaiter_triton_mxfp4_bf16, selecting the CK a8w4 experts1. Size CUDA graph capture to the point being measured
CAPTURE_SIZEwas derived fromMAX_NUM_SEQS, which is a constant, so--max-cudagraph-capture-sizewas 1024 at every point in the sweep regardless of what was being measured. At CONC=1 the engine never sees a decode batch larger than 6, yet it still captured graphs out to 1024: 18.71 GiB of PIECEWISE plus 19.40 GiB of FULL capture memory per GPU, roughly 30% of what is left after weights, spent on shapes the workload cannot reach.This is hygiene, not a throughput fix, and it is not expected to move any published metric. Reclaiming that memory grows the KV pool from 33.6M tokens, but KV was never the constraint at these points (run 34710937012 server logs):
Nothing was evicted and nothing refilled. Extrapolating the observed doubling, conc 128 would sit near half of even the old pool, so the ladder could have been extended without this change. The reasons to make it are that capture should describe the point being measured, that 38.1 GiB per GPU plus the startup time to capture 1024-wide shapes buys nothing, and that the headroom at the new points is then not something a reader has to reason about.
The derivation comes from the workload: each of the
CONClanes submits one real token plusNUM_SPEC_TOKENSdrafts per step, so the largest decode batch isCONC * (1 + NUM_SPEC_TOKENS), rounded up to a power of two.This matches the derivation the other DSv4.1 Flash agentic recipes already use, so nothing here is novel to MI355X — the AMD arm was simply deriving it from the wrong variable.
2. Extend the ladder to 64 and 128
conc-listwas[1, 2, 4, 8, 16, 32], which stops before the curve does anything interesting. From the published sweep:At the top of the ladder throughput is still climbing steeply — output throughput per GPU rises 62% from conc 16 to 32 with no sign of a knee — and energy per output token is still improving, 3.81 → 2.60 J. TTFT is flat near 0.9 s and KV occupancy peaked at 13%, so nothing indicates the engine is near saturation.
There is a sharper reason than the shape of the curve. AgentX is a closed-loop harness: each of the
CONClanes issues a request, waits for it, then spends time on the client side of the agent turn, so aggregate throughput is the product of per-request output rate and how many requests are actually in flight. Computing average in-flight requests assum(e2el) / windowshows how much of the cap the sweep is using:Up to conc 8 the engine is served by a quarter of the lanes, so the aggregate figure is set by trace demand rather than by the engine, and per-request latency is absorbed by lanes sitting idle between turns. Occupancy only starts climbing at 16 and 32. The points where lane occupancy approaches the cap — and where aggregate throughput therefore becomes a statement about the engine rather than about the harness — are above the current top of the ladder, which is exactly the region that determines serving cost.
conc-listbecomes[1, 2, 4, 8, 16, 32, 64, 128], the same ladder the other DSv4.1 Flash agentic recipes run.3. Engram note corrected
The previous wording ("upstream currently rejects it on ROCm") was accurate but incomplete, and invited the reading that MI355X therefore keeps the Engram tables in HBM. It does not.
EngramConfig.verify_model_configgates oncurrent_platform.is_cuda(), which is strictlyPlatformEnum.CUDA, so passing--engram-configon ROCm raisesValueErrorand the engine refuses to start. ButVllmConfig._resolve_and_verify_engram_configreturns before validation when the flag is absent, andmodels/deepseek_v4_1/common/engram.pythen defaults tocpu_offload=True. Every MI355X run already logs the result:Four ranks × two Engram layers, so ~188 GiB in pinned host memory. Omitting the flag is therefore not a downgrade, and the capture-memory figures in section 1 have nothing to do with Engram placement.
For the record, the underlying machinery is not CUDA-only:
is_uva_available()requires only pinned memory and returns True here, andget_accelerator_view_from_cpu_tensor()already dispatches onis_cuda_alike(), which includes ROCm, intotorch.ops._C.get_cuda_view_from_cpu_tensor— an op that is registered in this ROCm build. The upstream gate looks like a conservative whitelist rather than a capability limit.Not changed
MAX_NUM_SEQSstays pinned at 128, which isSchedulerConfig.DEFAULT_MAX_NUM_SEQS, so the scheduler cap is unchanged.--gpu-memory-utilizationstays at 0.9 per upstream recipe #946.dram-utilizationstays at 0.60 since it only takes effect forkv-offloading: dram.Validation
python -m pytest utils/matrix_logic/ -v: 319 passed.python -m infx.workflows.validate_perf_changelogpasses locally againstorigin/main.bash -non the recipe passes; the capture derivation was executed for all eight concurrencies and produces the table in section 1.configs/amd-master.yamlparses and the resolved search space istp: 4, kv-offloading: none, spec-decoding: mtpwith the newconc-list.Runtime evidence for the two new points is pending. Note that the MI355X runners have been failing independently of this recipe: legs die with a worker killed by signal during PIECEWISE capture (
Worker proc VllmWorker-N died unexpectedly (exit code: None)) on both MoE backends and on four different nodes, while the identical serve command completes on a dedicated MI355X box. Capturing far fewer graph shapes shortens that window but is not a fix for it.