perf(agentx): switch K3 MI355X DCP8 arm to MTP spec-decoding - #2809
perf(agentx): switch K3 MI355X DCP8 arm to MTP spec-decoding#2809seungrokj wants to merge 8 commits into
Conversation
Use the merged Kimi-K3 recipe with vLLM nightly 7c5dc571 and LMCache 0.5.5rc3 for C1, C8, C14, and TP8/DCP8 C40. Keep the validated C40 profile with GPU memory utilization reduced to 0.88. Assisted-by: OpenAI Codex <codex@openai.com>
Assisted-by: OpenAI Codex <codex@openai.com>
Assisted-by: OpenAI Codex <codex@openai.com>
为 Kimi-K3 MI355X LMCache 曲线添加 C44 和 C48，并沿用 C40 的 TP8/DCP8、无推测解码和 GMU 0.88 配置。 Assisted-by: OpenAI Codex <codex@openai.com>
将 LMCache MP 心跳超时提高到 90 秒，并将 worker 回收超时提高到 300 秒，避免 C14 大批量传输期间的短暂阻塞触发错误恢复路径。 Assisted-by: OpenAI Codex <codex@openai.com>
将 C1、C8 和 C14 的 gpu-memory-utilization 从 0.90 降至 0.88，使全部六个测试点使用相同的 GMU。移除上一轮未证实的 heartbeat workaround，以便单独验证显存利用率变化。 Assisted-by: OpenAI Codex <codex@openai.com>
Move the second kimik3-fp4-mi355x-vllm-agentic-mtp search-space arm from a no-spec DCP8 curve (conc 40/44/48, LMCache 0.5.5rc3+rocm7.2) to an MTP spec-decoding DCP8 curve (conc 44/48/52, LMCache 0.5.5rc3), and refresh the agentic trace-replay benchmark script. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
|
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 关于重新运行失败任务的文档 |
2 similar comments
|
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 关于重新运行失败任务的文档 |
|
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 关于重新运行失败任务的文档 |
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
![claude[bot]](https://avatars.githubusercontent.com/in/1236702?s=60&v=4){kind=small}
| case "$CONC" in | ||
| # No KV offload; the working set fits in HBM. | ||
| 1) | ||
| SYNTHETIC_ACCEPT_LEN=3.75 | ||
| SPEC_NUM_TOKENS=6 | ||
| GPU_MEM_UTIL=0.9 | ||
| MAX_NUM_BATCHED_TOKENS=16384 | ||
| ;; | ||
| 2|4|8|10|12|14) | ||
| SYNTHETIC_ACCEPT_LEN=3.00 | ||
| SPEC_NUM_TOKENS=3 | ||
| GPU_MEM_UTIL=0.9 | ||
| MAX_NUM_BATCHED_TOKENS=8192 | ||
| ;; | ||
| *) | ||
| SPEC_NUM_TOKENS=0 | ||
| GPU_MEM_UTIL=0.9 | ||
| MAX_NUM_BATCHED_TOKENS=8192 | ||
| ;; | ||
| esac |
There was a problem hiding this comment.
🔴 amd-master.yaml now sets spec-decoding: mtp with conc-list [44, 48, 52] for this arm, but the script's case "$CONC" only enables speculative decoding for CONC in {1,2,4,8,10,12,14}; 44/48/52 all fall into the *) branch which sets SPEC_NUM_TOKENS=0, so SPEC_ARGS stays empty and the server runs with no --speculative-config at all. Fix: extend the case statement (or key off spec-decoding/SPEC_DECODE rather than a hardcoded CONC allow-list) so every conc value in the new search-space arm actually gets MTP speculative decoding, not silently falls back to none.
Extended reasoning...
configs/amd-master.yaml line 640 changes the second kimik3-fp4-mi355x-vllm-agentic-mtp search-space entry from spec-decoding: none / conc-list [40,44,48] to spec-decoding: mtp / conc-list [44,48,52]. The sweep launches the script with CONC=44 (or 48/52); at lines 251-270 the case statement has explicit arms only for 1 and 2|4|8|10|12|14, so these values hit the default *) arm (SPEC_NUM_TOKENS=0). At line 273 if [ "$SPEC_NUM_TOKENS" -gt 0 ] is false, so SPEC_ARGS is never populated and no --speculative-config is passed to vllm serve. The benchmark silently runs the no-spec baseline while being labeled/reported as the MTP arm, producing misleading perf numbers for the exact change this PR claims to make.
Verification: normal. The new arm in configs/amd-master.yaml:640 sets conc-list: [44, 48, 52] with spec-decoding: mtp (previously [40,44,48] / none). The sweep launches the rewritten script once per conc value (CONC=44/48/52). In the script, speculative decoding is gated solely on CONC via the case statement at lines 251-270: only 1 and 2|4|8|10|12|14 set a nonzero SPEC_NUM_TOKENS; 44/48/52 all…
| - { tp: 8, kv-offloading: none, conc-list: [1] , spec-decoding: mtp} | ||
| - { tp: 8, ep: 1, kv-offloading: dram, kv-offload-backend: { name: lmcache, version: "0.5.5.dev60+rocm7.2" }, conc-list: [4, 8, 10, 12, 14], spec-decoding: mtp } | ||
| - { tp: 8, ep: 1, kv-offloading: dram, kv-offload-backend: { name: lmcache, version: "0.5.5rc3+rocm7.2" }, conc-list: [1, 8, 14], spec-decoding: mtp } | ||
| - { tp: 8, ep: 1, dcp-size: 8, kv-offloading: dram, kv-offload-backend: { name: lmcache, version: "0.5.5rc3" }, conc-list: [44, 48, 52], spec-decoding: mtp } |
There was a problem hiding this comment.
🔴 This diff changes a benchmark search-space arm's spec-decoding, conc-list and LMCache version, and rewrites the paired script's runtime knobs, but adds no perf-changelog.yaml entry, violating AGENTS.md's non-negotiable rule that every change affecting benchmark performance or recipe modification requires a new append-only changelog entry. Fix: append a new perf-changelog.yaml entry (matching the existing kimik3-fp4-mi355x-vllm-agentic-mtp entries' format) describing this DCP8-arm MTP switch, without altering any existing bytes.
Extended reasoning...
perf-changelog.yaml already has three entries keyed on kimik3-fp4-mi355x-vllm-agentic-mtp (lines ~5429, 6622, 6826), establishing that this config-key's changes are tracked there. This PR modifies the same config key's search-space (spec-decoding none->mtp, conc-list, lmcache version) and heavily rewrites the associated benchmark script, but perf-changelog.yaml is untouched in the diff. Merging as-is leaves downstream consumers of the changelog (perf history/audit tooling) unaware this arm's behavior changed, breaking the append-only audit trail AGENTS.md requires for every performance-affecting change.
Verification: normal. AGENTS.md states a non-negotiable invariant: "Every change that can affect benchmark performance and every recipe addition or modification requires a new perf-changelog.yaml entry. The file is append-only... append only at the tail." The diff (git diff a720a04..HEAD) touches only 2 files — configs/amd-master.yaml and benchmarks/single_node/agentic/glm5.2_fp4_mi355x_sglang_mtp.sh —…
| echo "Starting vllm server..." | ||
| export PYTHONNOUSERSITE=1 | ||
| export VLLM_EXECUTE_MODEL_TIMEOUT_SECONDS="${VLLM_EXECUTE_MODEL_TIMEOUT_SECONDS:-1200}" | ||
|
|
||
|
|
||
| # ---- DCP ------------------------------------------------------------ | ||
| # DCP shards decode KV across the TP ranks, so it must divide TP. | ||
| DCP_SIZE="${DCP_SIZE:-8}" | ||
| if [ $((TP % DCP_SIZE)) -ne 0 ]; then | ||
| echo "Error: TP='$TP' must be divisible by DCP_SIZE='$DCP_SIZE'" >&2 | ||
| exit 1 | ||
| fi | ||
| # 2×CONC in-flight slots: MTP draft+verify transiently batches more tokens | ||
| # than CONC sessions; headroom prevents scheduler stalls under burst. | ||
| MAX_RUNNING_REQUESTS=$((2 * CONC)) | ||
| [ "$MAX_RUNNING_REQUESTS" -gt 256 ] && MAX_RUNNING_REQUESTS=256 | ||
| # SGLang interpolates a bs list [1..max_bs] automatically; cap at 64 to | ||
| # keep graph-capture memory bounded without giving up coverage. | ||
| CUDA_GRAPH_MAX_BS=$(( MAX_RUNNING_REQUESTS < 64 ? MAX_RUNNING_REQUESTS : 64 )) | ||
|
|
||
| if [ "${EVAL_ONLY:-false}" != "true" ]; then | ||
| export SGLANG_SIMULATE_ACC_LEN=3.61 | ||
| export SGLANG_SIMULATE_ACC_METHOD=match-expected | ||
| export SGLANG_SIMULATE_ACC_TOKEN_MODE=real-draft-token | ||
| CP_ARGS=() | ||
| ATTN_BE_ARGS=() | ||
| if [ "$DCP_SIZE" -gt 1 ]; then | ||
| CP_ARGS+=(--decode-context-parallel-size "$DCP_SIZE" --dcp-comm-backend a2a) | ||
| ATTN_BE_ARGS+=(--attention-backend ROCM_AITER_MLA) | ||
| fi | ||
|
|
||
| SGLANG_CMD=( | ||
| python3 -m sglang.launch_server | ||
| --model-path "$MODEL_PATH" | ||
| --served-model-name "$MODEL" | ||
| export VLLM_USE_DIRECT_DCP_A2A=0 | ||
| export VLLM_USE_DIRECT_DCP_Q_GATHER=0 | ||
| export VLLM_USE_DIRECT_DCP_KV_GATHER=0 | ||
|
|
||
| { set +x; } 2>/dev/null | ||
| VLLM_CMD=( | ||
| vllm serve "$MODEL_PATH" --served-model-name "$MODEL" | ||
| --host 0.0.0.0 | ||
| --port "$SGLANG_BACKEND_PORT" | ||
| --port "$PORT" | ||
| --trust-remote-code | ||
| "${PARALLEL_ARGS[@]}" | ||
| --kv-cache-dtype fp8_e4m3 | ||
| --dsa-prefill-backend tilelang | ||
| --dsa-decode-backend tilelang | ||
| # GLM-5.2 emits the GLM-4.7-style tool-call format; glm47 is required for | ||
| # structured message.tool_calls (SWE-bench agentic evals die without it). | ||
| # The glm45 reasoning parser keeps hybrid thinking in reasoning_content. | ||
| --tool-call-parser glm47 | ||
| --reasoning-parser glm45 | ||
| --chunked-prefill-size "$CHUNKED_PREFILL_SIZE" | ||
| --mem-fraction-static "$MEM_FRACTION_STATIC" | ||
| --max-running-requests "$MAX_RUNNING_REQUESTS" | ||
| --cuda-graph-max-bs "$CUDA_GRAPH_MAX_BS" | ||
| --speculative-algorithm EAGLE | ||
| --speculative-num-steps 5 | ||
| --speculative-eagle-topk 1 | ||
| --speculative-num-draft-tokens 6 | ||
| "${CACHE_ARGS[@]}" | ||
| --watchdog-timeout 1800 | ||
| --enable-metrics | ||
| --moe-backend auto | ||
| --tensor-parallel-size "$TP" | ||
| "${EP_ARGS[@]}" | ||
| --load-format fastsafetensors | ||
| --gpu-memory-utilization "$GPU_MEM_UTIL" | ||
| --language-model-only | ||
| --max-num-seqs "$MAX_NUM_SEQS" | ||
| --enable-auto-tool-choice | ||
| --tool-call-parser kimi_k3 |
There was a problem hiding this comment.
🔴 This file's path is resolved (via runners/launch_mi355x-amds.sh's ${model-prefix}_${precision}_mi355x_${framework}.sh scheme) by config glm5.2-fp4-mi355x-sglang-agentic-mtp (sglang framework, model amd/GLM-5.2-MXFP4, TP4/EP4 and TP8/EP1 arms) -- but the rewrite replaces the sglang GLM-5.2 launch with a hardcoded vllm serve command using Kimi-K3-only flags (--tool-call-parser kimi_k3, --reasoning-parser kimi_k3, speculative model Inferact/Kimi-K3-DSpark) and a hard TP -ne 8 exit, so that real benchmark now launches the wrong binary with the wrong model config and always exits at TP=4. Fix: revert this file to the GLM-5.2/sglang recipe; put the new Kimi-K3/vLLM content in whatever file kimik3-fp4-mi355x-vllm-agentic-mtp actually resolves to.
Extended reasoning...
runners/launch_mi355x-amds.sh:307-309 computes SCRIPT_FW="kimik3_fp4_mi355x_vllm_mtp.sh" for the PR's target config kimik3-fp4-mi355x-vllm-agentic-mtp (model-prefix kimik3, framework vllm); that file does not exist, so it falls back to the existing, untouched kimik3_fp4_mi355x_mtp.sh -- meaning this diff's script rewrite is never even used by the config change it was meant to support. Meanwhile config glm5.2-fp4-mi355x-sglang-agentic-mtp (amd-master.yaml:1710-1723, not modified by this diff) resolves to exactly agentic/glm5.2_fp4_mi355x_sglang_mtp.sh per the same naming scheme (model-prefix=glm5.2, precision=fp4, framework=sglang). Its TP4/EP4 arm now trips if [ "$TP" -ne 8 ]; then exit 1 (line ~78 in the new file) and any TP=8 run launches vllm serve with GLM-5.2 weights but Kimi-K3 tool/reasoning parsers and a Kimi-K3 DSpark draft model path -- both arms of that production benchmark now fail outright where they previously worked.
Verification: normal. The rewrite landed in the wrong file. glm5.2_fp4_mi355x_sglang_mtp.sh is the launch script for master-config arm glm5.2-fp4-mi355x-sglang-agentic-mtp (configs/amd-master.yaml:1710-1723: model-prefix glm5.2, framework sglang, precision fp4, runner mi355x-amds, agentic-coding). runners/launch_mi355x-amds.sh:307-313 builds… | normal (regression). The runner resolves benchmark scripts by…
| # Perf-search knobs. Each defaults to the reference command's value, so an | ||
| # otherwise-unset run reproduces the reference exactly: | ||
| # GPU_MEM_UTIL 0.95 (reference) | ||
| # MAX_NUM_BATCHED_TOKENS 8192 (default) | ||
| # AITER_A8W4 1 (reference; 0 = aiter a16w4 MoE path) | ||
| # LANGUAGE_MODEL_ONLY true | ||
| # KV_CACHE_DTYPE fp8 (default for every arm; =auto for a bf16 A/B) | ||
| # KV_BLOCK_SIZE unset (unset -> vLLM sizes the page; 128 under fp8) | ||
| # MAX_MODEL_LEN 1M | ||
| # SPEC_DECODE true (this is the _mtp DSpark recipe; =false for a no-spec A/B) | ||
| # SPEC_NUM_TOKENS 2 (DSpark draft length; validated by the _mtp config) | ||
|
|
||
| source "$(dirname "$0")/../../benchmark_lib.sh" | ||
|
|
||
| export EVAL_FRAMEWORK="lm-eval" | ||
| check_env_vars MODEL TP CONC KV_OFFLOADING TOTAL_CPU_DRAM_GB RESULT_DIR DURATION EP_SIZE DP_ATTENTION | ||
| if [[ -n "$SLURM_JOB_ID" ]]; then | ||
| echo "JOB $SLURM_JOB_ID running on $SLURMD_NODENAME" | ||
| wait_for_amd_gpu_clean | ||
|
|
||
| check_env_vars MODEL TP CONC KV_OFFLOADING TOTAL_CPU_DRAM_GB RESULT_DIR DURATION EP_SIZE | ||
|
|
||
| if [[ -n "${SLURM_JOB_ID:-}" ]]; then | ||
| echo "JOB $SLURM_JOB_ID running on ${SLURMD_NODENAME:-unknown}" | ||
| fi | ||
|
|
||
| # ROCR/HIP visibility under slurm cgroups. | ||
| if [ -n "$ROCR_VISIBLE_DEVICES" ]; then | ||
|
|
||
| if [ "$TP" -ne 8 ]; then | ||
| echo "Error: Kimi-K3 MXFP4 is a 1.56 TB checkpoint and only fits at TP=8 on" >&2 | ||
| echo " 288 GB gfx950 parts (~195 GB/GPU). Got TP=$TP." >&2 | ||
| exit 1 | ||
| fi | ||
|
|
||
| # ROCR/HIP visibility for vLLM 0.14+ | ||
| if [ -n "${ROCR_VISIBLE_DEVICES:-}" ]; then | ||
| export HIP_VISIBLE_DEVICES="$ROCR_VISIBLE_DEVICES" | ||
| fi | ||
|
|
||
|
|
||
| if [[ -n "$MODEL_PATH" ]]; then | ||
|
|
||
| # `hf download` creates the target dir if missing and is itself idempotent. The | ||
| # 1.56 TB checkpoint is normally pre-staged, so these calls are a no-op there. | ||
| if [[ -n "${MODEL_PATH:-}" ]]; then | ||
| if [[ ! -d "$MODEL_PATH" || -z "$(ls -A "$MODEL_PATH" 2>/dev/null)" ]]; then | ||
| hf download "$MODEL" --local-dir "$MODEL_PATH" | ||
| fi | ||
| else | ||
| hf download "$MODEL" | ||
| export MODEL_PATH="$MODEL" | ||
| fi | ||
|
|
||
| rocm-smi || true | ||
| amd-smi || true | ||
|
|
||
| # A server killed on this node minutes earlier (previous job, crashed run) | ||
| # can still be draining its ~1.4 TB of HBM: KFD reclaim takes minutes, and | ||
| # booting into a half-drained node fails RCCL init with HIP 'unhandled cuda | ||
| # error' / 'invalid argument' (observed as the mooncake-c64 CI failure). | ||
| # Wait for the GPUs to come back before launching. | ||
| # Per-GPU threshold: idle nodes hold a small driver/firmware VRAM baseline | ||
| # (observed up to ~4%/GPU, node-dependent), while a draining or occupied | ||
| # GPU sits at 50-90%. Require every GPU <= 10%. | ||
| GPU_CLEAN=false | ||
| for i in $(seq 1 90); do | ||
| VRAM_MAX=$(rocm-smi --showmemuse 2>/dev/null | grep -oE "GPU Memory Allocated \(VRAM%\): [0-9]+" | awk '{if ($NF > m) m = $NF} END {print m+0}') | ||
| if [ "${VRAM_MAX:-0}" -le 10 ]; then echo "GPUs clean (vram%max=$VRAM_MAX after $((i*10))s)"; GPU_CLEAN=true; break; fi | ||
| echo "waiting for prior-job GPU memory reclaim: vram%max=$VRAM_MAX"; sleep 10 | ||
| done | ||
| [ "$GPU_CLEAN" = "true" ] || { echo "Error: GPUs still draining prior job's memory after 15min" >&2; exit 1; } | ||
|
|
||
|
|
||
| # ---- Resolve traces and install deps ---------------------------------------- | ||
| resolve_trace_source | ||
| install_agentic_deps | ||
|
|
||
|
|
||
| # ---- Reference env block ---------------------------------------------------- | ||
| export VLLM_ROCM_AITER_MLA_ASM_PADDING=asm | ||
| export VLLM_ROCM_USE_AITER=1 | ||
| export SAFETENSORS_FAST_GPU=1 | ||
| export VLLM_ROCM_USE_AITER_MOE_SITUV2_A8W4=1 | ||
| export AITER_SITUV2_A8W4=1 | ||
| export AITER_BF16_FP8_MOE_BOUND=0 | ||
| export VLLM_USE_BREAKABLE_CUDAGRAPH=0 | ||
| export AITER_QUICK_REDUCE_QUANTIZATION=INT4 | ||
|
|
||
| # Workaround for MEC FW <177 RCCL memory reclaim issue (shared with the other | ||
| # gfx950 recipes in this tree). | ||
| mec_version=$(rocm-smi --showfw 2>/dev/null | grep MEC | head -n 1 | awk '{print $NF}') | ||
| if [[ "$mec_version" == "" || ${mec_version:-0} -lt 177 ]]; then | ||
| export HSA_NO_SCRATCH_RECLAIM=1 | ||
| fi | ||
|
|
||
| # 2.8T of weights off a shared/NFS mount takes far longer than the default. | ||
| export VLLM_ENGINE_READY_TIMEOUT_S="${VLLM_ENGINE_READY_TIMEOUT_S:-7200}" | ||
|
|
||
| # Long agentic turns against a 1M context: keep the client from timing out | ||
| # mid-request while the server is prefill-bound. | ||
| export AIPERF_HTTP_TCP_USER_TIMEOUT=900000 | ||
|
|
||
| # ---- Server config ---------------------------------------------------------- | ||
| SERVER_LOG="$RESULT_DIR/server.log" | ||
| ROUTER_LOG="$RESULT_DIR/router.log" | ||
| mkdir -p "$RESULT_DIR" | ||
|
|
||
| export PYTHONNOUSERSITE=1 | ||
| # Agentic warmup dispatches hundreds of large prompts at once; allow up to | ||
| # 15 minutes of TCP progress before AIPerf declares a connection dead. | ||
| export AIPERF_HTTP_TCP_USER_TIMEOUT=900000 | ||
| # AIPerf pins one pooled keep-alive connection per session (client-side | ||
| # keep-alive 300s) while uvicorn's default SGLANG_TIMEOUT_KEEP_ALIVE is 5s; | ||
| # inter-turn idle gaps can reuse a socket exactly as the server closes it. | ||
| # Outlast the client pool so the race cannot occur. | ||
| export SGLANG_TIMEOUT_KEEP_ALIVE=900 | ||
| # The DSA indexer's top-k v2 kernel (default since v0.5.14) is JIT-compiled | ||
| # from CUDA-only source (cooperative_groups.h) and cannot build for gfx950; | ||
| # v1 dispatches to the precompiled HIP op in sgl-kernel (upstream MI355X CI | ||
| # runs DSA models the same way). | ||
| export SGLANG_OPT_USE_TOPK_V2=false | ||
|
|
||
| # HiCache L2 (host DRAM), optionally extended with Mooncake L3. | ||
| # KV_OFFLOADING=dram requires KV_OFFLOAD_BACKEND=hicache or mooncake. | ||
| # | ||
| # Per-arm L2 ratio (sizing rationale below) applies to both backends unless | ||
| # overridden via HICACHE_RATIO. TP arm (182.7 GB/rank device pool): the | ||
| # agentic-coding corpus saturates any fixed DRAM pool at conc ≥ 10; ratio 1.5 | ||
| # (~2.9 TB pinned) is the safe default for cluster:mi355x-amds nodes (~3.0 TB | ||
| # available DRAM per runners.yaml). ratio=2.5 (~4.8 TB) yields higher | ||
| # throughput at conc 10-12 but exceeds physical DRAM on these nodes and must | ||
| # be set via HICACHE_RATIO env-var override on nodes that can accommodate it. | ||
| # The DP-attention arm (159.4 GB/rank) only runs at conc >= 32, where the host | ||
| # tier just absorbs overflow - ratio 0.5 (~1.2 TB pinned, ~1.8 TB of load | ||
| # headroom) at negligible hit-rate cost (ratio 1.5 OOMs the host mid-storm at | ||
| # conc 48). | ||
| CACHE_ARGS=() | ||
| if agentic_kv_offload_enabled; then | ||
| if [ "$DP_ATTENTION" = "true" ]; then | ||
| HICACHE_RATIO="${HICACHE_RATIO:-0.5}" | ||
| else | ||
| # ratio=1.5 (~2.9 TB pinned): safe default within the ~3.0 TB DRAM | ||
| # available on cluster:mi355x-amds nodes. Set HICACHE_RATIO=2.5 via | ||
| # env-var override for maximum throughput on nodes with >4 TB DRAM. | ||
| HICACHE_RATIO="${HICACHE_RATIO:-1.5}" | ||
| fi | ||
| # write_through_selective skips DRAM writes for non-reusable KV blocks, | ||
| # reducing host-bus traffic without affecting the cache hit rate. | ||
| HICACHE_WRITE_POLICY="${HICACHE_WRITE_POLICY:-write_through_selective}" | ||
| HICACHE_IO_BACKEND="${HICACHE_IO_BACKEND:-direct}" | ||
| HICACHE_MEM_LAYOUT="${HICACHE_MEM_LAYOUT:-page_first_direct}" | ||
| case "$KV_OFFLOAD_BACKEND" in | ||
| hicache) | ||
| echo "HiCache (GPU+host DRAM only): ratio=$HICACHE_RATIO, write_policy=$HICACHE_WRITE_POLICY, io_backend=$HICACHE_IO_BACKEND, mem_layout=$HICACHE_MEM_LAYOUT" | ||
| CACHE_ARGS=( | ||
| --enable-hierarchical-cache | ||
| --hicache-ratio "$HICACHE_RATIO" | ||
| --hicache-write-policy "$HICACHE_WRITE_POLICY" | ||
| --hicache-io-backend "$HICACHE_IO_BACKEND" | ||
| --hicache-mem-layout "$HICACHE_MEM_LAYOUT" | ||
| ) | ||
| ;; | ||
| mooncake) | ||
| L3_PER_RANK_GB="${L3_PER_RANK_GB:-40}" | ||
| python3 -c "from mooncake.store import MooncakeDistributedStore" >/dev/null | ||
| MOONCAKE_MASTER_PORT=$((PORT + 12000)) | ||
| MOONCAKE_MASTER_LOG="$RESULT_DIR/mooncake_master.log" | ||
| MOONCAKE_CONFIG_PATH="$RESULT_DIR/mooncake_config.json" | ||
| cat > "$MOONCAKE_CONFIG_PATH" <<EOF | ||
| { | ||
| "local_hostname": "127.0.0.1", | ||
| "metadata_server": "P2PHANDSHAKE", | ||
| "master_server_address": "127.0.0.1:$MOONCAKE_MASTER_PORT", | ||
| "global_segment_size": "${L3_PER_RANK_GB}gb", | ||
| "local_buffer_size": "4gb", | ||
| "protocol": "tcp", | ||
| "device_name": "" | ||
|
|
||
| SERVER_PID="" | ||
| LMCACHE_PID="" | ||
|
|
||
| cleanup_agentic_services() { | ||
| local exit_code=$? | ||
| trap - EXIT INT TERM | ||
| set +e | ||
| stop_background_process_tree "$SERVER_PID" "vLLM server" 60 | ||
| stop_background_process_tree "$LMCACHE_PID" "LMCache server" | ||
| exit "$exit_code" | ||
| } | ||
| EOF | ||
| export SGLANG_HICACHE_MOONCAKE_CONFIG_PATH="$MOONCAKE_CONFIG_PATH" | ||
| mooncake_master --port "$MOONCAKE_MASTER_PORT" \ | ||
| --default_kv_lease_ttl=120s \ | ||
| --eviction_high_watermark_ratio=0.80 \ | ||
| --eviction_ratio=0.10 > "$MOONCAKE_MASTER_LOG" 2>&1 & | ||
| MOONCAKE_MASTER_PID=$! | ||
| sleep 2 | ||
| kill -0 "$MOONCAKE_MASTER_PID" | ||
| echo "HiCache+Mooncake: ratio=$HICACHE_RATIO, l3_per_rank=${L3_PER_RANK_GB} GB, dram_budget=${TOTAL_CPU_DRAM_GB} GB" | ||
| CACHE_ARGS=( | ||
| --enable-hierarchical-cache | ||
| --hicache-ratio "$HICACHE_RATIO" | ||
| --hicache-size 0 | ||
| --hicache-write-policy "$HICACHE_WRITE_POLICY" | ||
| --hicache-io-backend "$HICACHE_IO_BACKEND" | ||
| --hicache-mem-layout "$HICACHE_MEM_LAYOUT" | ||
| --hicache-storage-backend mooncake | ||
| --hicache-storage-prefetch-policy wait_complete | ||
| ) | ||
| ;; | ||
| *) | ||
| echo "Error: unsupported KV_OFFLOAD_BACKEND '$KV_OFFLOAD_BACKEND' (expected: hicache or mooncake)" >&2 | ||
| exit 1 | ||
| ;; | ||
| esac | ||
| trap cleanup_agentic_services EXIT | ||
| trap 'exit 130' INT | ||
| trap 'exit 143' TERM | ||
|
|
||
| # ---- KV offload ------------------------------------------------------------- | ||
| # TOTAL_CPU_DRAM_GB is the aggregate host-DRAM budget the matrix generator | ||
| # derives from dram-utilization and the runner's available-cpu-dram-mib, capped | ||
| # at the 3,095,781 MiB (3 TB decimal) agentic limit. Per | ||
| # benchmarks/single_node/agentic/README.md it must be consumed as given and | ||
| # never replaced with a model-specific constant. | ||
| OFFLOAD_ARGS=() | ||
|
|
||
| if agentic_kv_offload_enabled; then | ||
| case "${KV_OFFLOAD_BACKEND:-}" in | ||
| vllm-simple) | ||
| require_agentic_kv_offload_backend "$KV_OFFLOAD_BACKEND" | ||
| CPU_BYTES_PER_RANK=$(( TOTAL_CPU_DRAM_GB * 1000 * 1000 * 1000 / TP )) | ||
| # Identical prefixes must hash to identical block keys across ranks. | ||
| export PYTHONHASHSEED=42 | ||
| SIMPLE_LAZY_OFFLOAD="${SIMPLE_LAZY_OFFLOAD:-false}" | ||
| OFFLOAD_ARGS=( | ||
| --kv-transfer-config | ||
| "{\"kv_connector\":\"SimpleCPUOffloadConnector\",\"kv_role\":\"kv_both\",\"kv_connector_extra_config\":{\"cpu_bytes_to_use_per_rank\":$CPU_BYTES_PER_RANK,\"lazy_offload\":$SIMPLE_LAZY_OFFLOAD}}" | ||
| ) | ||
| echo "SimpleCPUOffloadConnector: ${CPU_BYTES_PER_RANK} B/rank x ${TP} ranks, lazy_offload=$SIMPLE_LAZY_OFFLOAD" | ||
| ;; | ||
| lmcache) | ||
| require_agentic_kv_offload_backend "$KV_OFFLOAD_BACKEND" | ||
|
|
||
| # Keep the image's tested torch/ROCm stack and install only LMCache's | ||
| # missing runtime dependencies, same as the MiniMax-M3 lmcache arm. | ||
| LMCACHE_VERSION="0.5.5rc3" | ||
| LMCACHE_ROCM_INDEX="https://github.com/LMCache/LMCache/releases/expanded_assets/nightly-rocm" | ||
| agentic_pip_install --quiet --no-cache-dir --no-deps \ | ||
| "sortedcontainers==2.4.0" \ | ||
| "opentelemetry-exporter-prometheus==0.61b0" \ | ||
| "cupy-rocm-7-0==14.1.1" \ | ||
| "lmcache==${LMCACHE_VERSION}" --find-links "$LMCACHE_ROCM_INDEX" | ||
|
|
||
| # LMCache 0.5.5's transfer-channel layer eagerly imports the Mooncake | ||
| # backend (mooncake_te_impl.py -> `from mooncake.engine import | ||
| # TransferEngine`), whose native .so resolves all of its DT_NEEDED libs at | ||
| # import. The vLLM ROCm image ships none of them, so the import sanity | ||
| # check below (and the LMCache server) would otherwise fail with | ||
| # "ImportError: lib*.so: cannot open shared object file" (first libglog, | ||
| # then libjsoncpp, ...). Provision Mooncake's full runtime lib set from the | ||
| # distro before importing. apt-get install is idempotent, so run it | ||
| # whenever any of the libs is still missing rather than gating on one. | ||
| LMCACHE_NATIVE_LIBS=(libglog.so.0 libjsoncpp.so.25 libibverbs.so.1 librdmacm.so.1 libnuma.so.1) | ||
| for lib in "${LMCACHE_NATIVE_LIBS[@]}"; do | ||
| if ! ldconfig -p | grep -q "$lib"; then | ||
| apt-get update | ||
| apt-get install -y \ | ||
| libgoogle-glog0v5 libjsoncpp25 libibverbs1 librdmacm1 libnuma1 | ||
| break | ||
| fi | ||
| done | ||
| python3 -c \ | ||
| "import cupy; import lmcache.integration.vllm.lmcache_mp_connector; import opentelemetry.exporter.prometheus" \ | ||
| >/dev/null | ||
|
|
||
| # One MP server for the node, per the Kimi-K3 recipe | ||
| # (docs.lmcache.ai/recipes/kimi_k3.html), with --chunk-size sized for | ||
| # THIS stack rather than the recipe's CUDA-path 768: the connector | ||
| # requires the chunk to be a multiple of every engine KV group's | ||
| # tokens_per_block, and the hybrid KDA/MLA layout here registers | ||
| # attention groups at 1536 ("Setting attention block size to 1536", | ||
| # run 31644990546) plus a KDA state group at 3072 (run 31645828378), | ||
| # so 3072 is the minimum valid chunk. The multi-group layout also | ||
| # requires one object group per sliding-window size: | ||
| # --separate-object-groups. | ||
| LMCACHE_PORT=6555 | ||
| LMCACHE_HTTP_PORT=8090 | ||
| LMCACHE_LOG="$RESULT_DIR/lmcache_server.log" | ||
|
|
||
| LMCACHE_L1_SIZE_GB="$TOTAL_CPU_DRAM_GB" | ||
|
|
||
| LMCACHE_CMD=( | ||
| lmcache server | ||
| --host 127.0.0.1 | ||
| --port "$LMCACHE_PORT" | ||
| --http-host 127.0.0.1 | ||
| --http-port "$LMCACHE_HTTP_PORT" | ||
| --l1-size-gb "$LMCACHE_L1_SIZE_GB" | ||
| --l1-init-size-gb 10 | ||
| --chunk-size 12288 | ||
| --separate-object-groups | ||
| --enable-extra-logging | ||
| --extra-logging-interval 30 | ||
| --max-cpu-workers 8 | ||
| --max-gpu-workers 8 | ||
| --eviction-policy LRU | ||
| --supported-transfer-mode lmcache_driven | ||
| --shm-name "" | ||
| ) | ||
| append_command "$RESULT_DIR/lmcache_command.txt" "${LMCACHE_CMD[@]}" | ||
| "${LMCACHE_CMD[@]}" > "$LMCACHE_LOG" 2>&1 & | ||
| LMCACHE_PID=$! | ||
| wait_for_ready \ | ||
| --endpoint "http://127.0.0.1:${LMCACHE_HTTP_PORT}/healthcheck" \ | ||
| --log "$LMCACHE_LOG" \ | ||
| --pid "$LMCACHE_PID" \ | ||
| --sleep-interval 1 \ | ||
| --timeout 600 | ||
|
|
||
| # 100k-330k-token agentic prefixes make single retrieves large; use the | ||
| # same MQ timeout headroom as the MiniMax-M3 arm. | ||
| OFFLOAD_ARGS=( | ||
| --kv-transfer-config | ||
| "{\"kv_connector\":\"LMCacheMPConnector\",\"kv_connector_module_path\":\"lmcache.integration.vllm.lmcache_mp_connector\",\"kv_role\":\"kv_both\",\"kv_connector_extra_config\":{\"lmcache.mp.port\":$LMCACHE_PORT,\"lmcache.mp.mq_timeout\":6000.0}}" | ||
| ) | ||
| ;; | ||
| *) | ||
| echo "Error: unsupported KV_OFFLOAD_BACKEND='$KV_OFFLOAD_BACKEND' (expected vllm-simple or lmcache)" >&2 | ||
| exit 1 | ||
| ;; | ||
| esac | ||
| fi | ||
|
|
||
| # Arm selection. TP arm keeps the FP8 sibling's cookbook batch-shaping | ||
| # bands. | ||
| # | ||
| # NOTE: the DP-attention path below is currently DORMANT (no dp-attn arms | ||
| # in amd-master.yaml): DSA + dp-attention hangs a collective under | ||
| # long-context prefill on ROCm v0.5.14 (watchdog kills the scheduler with | ||
| # zero completions; reproduced with and without HiCache, with and without | ||
| # the DSv4 DP collective envs; short prompts are fine). Re-enable the | ||
| # config arm once upstream fixes the DSA DP prefill path. | ||
| # | ||
| # When active, the DP-attention (DEP) arm fronts the DP ranks with sglang-router | ||
| # using consistent hashing on the AIPerf correlation id so multi-turn | ||
| # sessions stay on the DP rank holding their radix/hicache prefix, and | ||
| # widens chunked-prefill (whole-engine, /dp ranks) like the B300 sibling. | ||
| USE_SGLANG_ROUTER=false | ||
| SGLANG_BACKEND_PORT="$PORT" | ||
| PARALLEL_ARGS=(--tp "$TP" --ep-size "$EP_SIZE") | ||
| MEM_FRACTION_STATIC=0.85 | ||
| if [ "$DP_ATTENTION" = "true" ]; then | ||
| USE_SGLANG_ROUTER=true | ||
| export AIPERF_HTTP_X_SMG_ROUTING_KEY_FROM_CORRELATION_ID=true | ||
| SGLANG_BACKEND_PORT=$((PORT + 1)) | ||
| SGLANG_ROUTER_METRICS_PORT=$((PORT + 10000)) | ||
| SGLANG_ROUTER_CMD=(python3 -m sglang_router.launch_router) | ||
| PARALLEL_ARGS+=(--dp "$TP" --enable-dp-attention) | ||
| CHUNKED_PREFILL_SIZE=32768 | ||
| export AGENTIC_WARMUP_GRACE_PERIOD=3600 | ||
| # Swap the DP gather collectives to gatherv/reduce-scatter on ROCm | ||
| # (dsv4_fp4_mi355x_sglang.sh precedent - the only green DP-attention | ||
| # config on this cluster/image): with the defaults the DSA DP path | ||
| # hangs a collective under long-context prefill load until the | ||
| # watchdog kills the scheduler (0/96 storm completions, twice). | ||
| export SGLANG_DP_USE_GATHERV=1 | ||
| export SGLANG_DP_USE_REDUCE_SCATTER=1 | ||
| export GPU_MAX_HW_QUEUES=5 | ||
| elif [ "$CONC" -le 16 ]; then | ||
| # Chunked prefill 32k: smaller chunks let the scheduler interleave decode | ||
| # steps between prefill chunks, reducing TPOT for concurrent sessions | ||
| # (improved interactivity vs the original 131072-token chunk). The reduced | ||
| # chunk size drops per-chunk activation headroom from ~7 GiB/rank to | ||
| # ~1.7 GiB/rank, so mem-fraction 0.85 is safe (0.85 OOMed at 131k: | ||
| # "Tried to allocate 6.86 GiB ... 5.15 GiB is free", run 29751563205). | ||
| CHUNKED_PREFILL_SIZE=32768 | ||
| MEM_FRACTION_STATIC=0.85 | ||
|
|
||
| # ---- LLM server ------------------------------------------------------------ | ||
|
|
||
| # ---- Parallelism ------------------------------------------------------------ | ||
| EP_ARGS=() | ||
| if [ "$EP_SIZE" -gt 1 ]; then | ||
| EP_ARGS=(--enable-expert-parallel) | ||
| fi | ||
|
|
||
| # ---- Speculative / Util------------------------------------------------------ | ||
| case "$CONC" in | ||
| # No KV offload; the working set fits in HBM. | ||
| 1) | ||
| SYNTHETIC_ACCEPT_LEN=3.75 | ||
| SPEC_NUM_TOKENS=6 | ||
| GPU_MEM_UTIL=0.9 | ||
| MAX_NUM_BATCHED_TOKENS=16384 | ||
| ;; | ||
| 2|4|8|10|12|14) | ||
| SYNTHETIC_ACCEPT_LEN=3.00 | ||
| SPEC_NUM_TOKENS=3 | ||
| GPU_MEM_UTIL=0.9 | ||
| MAX_NUM_BATCHED_TOKENS=8192 | ||
| ;; | ||
| *) | ||
| SPEC_NUM_TOKENS=0 | ||
| GPU_MEM_UTIL=0.9 | ||
| MAX_NUM_BATCHED_TOKENS=8192 |
There was a problem hiding this comment.
🔴 Header (lines 33-43) documents GPU_MEM_UTIL, MAX_NUM_BATCHED_TOKENS, AITER_A8W4, LANGUAGE_MODEL_ONLY, KV_CACHE_DTYPE, KV_BLOCK_SIZE, MAX_MODEL_LEN, SPEC_DECODE as overridable perf-search knobs defaulting to reference values, but GPU_MEM_UTIL/MAX_NUM_BATCHED_TOKENS are hardcoded per-CONC (0.9, not the claimed 0.95) and the other five are never read anywhere in the script. Fix: read each var via ${VAR:-default} at its actual use site, or drop the unimplemented knobs from the header. Pre-existing, copied verbatim from kimik3_fp4_mi355x_mtp.sh.
Extended reasoning...
Line 35 claims GPU_MEM_UTIL defaults to the reference 0.95 and 'an otherwise-unset run reproduces the reference exactly', but lines 256/262/267 unconditionally set GPU_MEM_UTIL=0.9 inside the case "$CONC" block, overwriting any pre-set env var; the vllm command (line 325) always launches with --gpu-memory-utilization 0.9, never 0.95, and there is no way to override it. Same for MAX_NUM_BATCHED_TOKENS (lines 257/263/268). AITER_A8W4 (37), LANGUAGE_MODEL_ONLY (38), KV_CACHE_DTYPE (39), KV_BLOCK_SIZE (40), MAX_MODEL_LEN (41), SPEC_DECODE (42) never appear again in the file; --language-model-only (326), --kv-cache-dtype "fp8" (333) and --max-model-len 1048576 (331) are hardcoded literals disconnected from any env var, so setting these documented vars silently does nothing. The identical unwired documentation and hardcoded values already exist in benchmarks/single_node/agentic/kimik3_fp4_mi355x_mtp.sh (lines 33-42, 247-264), so this diff copies a pre-existing defect into a new file rather than introducing a new mistake from scratch.
Verification: nit. The header (lines 33-43) documents 8 "perf-search knobs" that "each default to the reference command's value, so an otherwise-unset run reproduces the reference exactly," but the code contradicts this: - GPU_MEM_UTIL: header line 35 says "0.95 (reference)", yet lines 256/262/267 unconditionally assign GPU_MEM_UTIL=0.9 in every case branch (direct assignment, not ${GPU_MEM_UTIL:-...}).…
|
|
||
| wait_for_server_ready --port "$PORT" --server-log "$SERVER_LOG" --server-pid "$SERVER_PID" | ||
|
|
||
| if [ "${EVAL_ONLY}" = "true" ]; then |
There was a problem hiding this comment.
🔴 The diff turns on set -u (line 2) but line 350 still references bare ${EVAL_ONLY} (no :- default), unlike the guarded use at line 274 and the other vars this same diff fixed (SLURM_JOB_ID, ROCR_VISIBLE_DEVICES, MODEL_PATH). If EVAL_ONLY is unset, bash aborts with "unbound variable" right after the server starts, before the replay/eval ever runs, whereas the old set -eo pipefail (no -u) silently treated it as empty/false. Fix: use ${EVAL_ONLY:-false} at line 350 (matching line 274) so an unset EVAL_ONLY behaves as false instead of crashing.
Extended reasoning...
The script's own header lists 'Required env vars' as MODEL, TP, CONC, KV_OFFLOADING, TOTAL_CPU_DRAM_GB, RESULT_DIR, DURATION, EP_SIZE — EVAL_ONLY is not in that list and check_env_vars does not require it. A caller (e.g. a developer manually reproducing a run per the documented required-vars list, or any invocation path that doesn't happen to export EVAL_ONLY) runs the script with only the documented vars set. With set -u now active, expanding ${EVAL_ONLY} at line 350 when the variable was never set triggers bash's nounset error and the script exits nonzero immediately after wait_for_server_ready succeeds, killing the server via the cleanup trap without ever running run_agentic_replay_and_write_outputs or run_eval. In the base script (set -eo pipefail, no -u) the same unset variable simply expanded to empty string and the comparison evaluated false, so the run proceeded normally.
Verification: normal (regression). Line 2 newly adds set -u (set -euo pipefail; base was set -eo pipefail), while line 350 keeps a bare if [ "${EVAL_ONLY}" = "true" ]; then with no :- default. On the base branch the equivalent bare reference (base line 275) was harmless because there was no -u; adding -u makes an unset EVAL_ONLY abort with "unbound variable". EVAL_ONLY is genuinely not…
2 participants
Summary
kimik3-fp4-mi355x-vllm-agentic-mtpsearch-space arm from a no-spec DCP8 curve (conc 40/44/48, LMCache0.5.5rc3+rocm7.2) to an MTP spec-decoding DCP8 curve (conc 44/48/52, LMCache0.5.5rc3).Stacked on #2804.
Test plan
kimik3-fp4-mi355x-vllm-agentic-mtp🤖 Generated with Claude Code
Note
Medium Risk
Changes the perf matrix and a large, multi-service benchmark launcher (vLLM + LMCache + DCP + speculative decoding) for a memory-constrained TP8 K3 recipe; mis-tuned conc/spec or offload settings would skew or break cluster runs rather than affecting app auth/data paths.
Overview
Updates
kimik3-fp4-mi355x-vllm-agentic-mtpso the high-concurrency TP8 / DCP8 + LMCache DRAM arm runs MTP (DSpark) instead of no speculative decoding: concurrency moves to 44/48/52,spec-decoding: mtp, and LMCache is pinned to0.5.5rc3(dropping the+rocm7.2suffix on that arm).The agentic launch script in the diff is rewritten from an SGLang GLM-5.2 path to the Kimi-K3 vLLM recipe:
vllm servewith AITER/ROCm env, TP=8 enforcement,wait_for_amd_gpu_clean, exit traps for vLLM and LMCache, SimpleCPUOffload or LMCache MP KV offload (replacing HiCache/Mooncake), DSpark--speculative-configtuned by concurrency, DCP8 decode-context parallel, and CUDA graph compilation sizing tied to concurrency and draft depth.perf-changelog.yamldocuments the search-space and script refresh (PR #2809).Reviewed by Cursor Bugbot for commit 82b3e60. Bugbot is set up for automated code reviews on this repo. Configure here.