From 1818c531505d0bd7b76b63f93564220176964fd6 Mon Sep 17 00:00:00 2001 From: seungrokj Date: Tue, 28 Jul 2026 14:19:27 +0900 Subject: [PATCH 1/8] [AMD] [AGENTX] dsv4-fp4-mi355x-vllm-agentic: DEP tuning + LMCache refresh Recipe: merge EP flag into PARALLEL_ARGS; add MODE_ARGS (ep-weight-filter, prefill-schedule-interval, DEP8 prefill budget --max-num-batched-tokens 8x / --long-prefill-token-threshold 16384); add --block-size 256, --max-model-len 1048576; switch cudagraph to FULL_DECODE_ONLY; gpu-mem-util 0.8 -> 0.95; enable VLLM_ROCM_QUICK_REDUCE_QUANTIZATION=INT4; drop the gist scheduler.py patch and the AITER_MOE export. LMCache: build v0.5.2 (was PR-3853 SHA), chunk-size 1024, undivided CPU pool. Fix TOTAL_CPU_DRAM_PARTITION_GB command substitution typo. amd-master.yaml: bump image, dram-utilization 0.60 -> 0.80, switch agentic search-space to DEP lmcache conc 76. Co-Authored-By: Claude Opus 4.6 --- .../agentic/dsv4_fp4_mi355x_vllm.sh | 73 +++++++++++++------ configs/amd-master.yaml | 11 +-- 2 files changed, 58 insertions(+), 26 deletions(-) diff --git a/benchmarks/single_node/agentic/dsv4_fp4_mi355x_vllm.sh b/benchmarks/single_node/agentic/dsv4_fp4_mi355x_vllm.sh index 241f042ff2..5792e1e4d2 100644 --- a/benchmarks/single_node/agentic/dsv4_fp4_mi355x_vllm.sh +++ b/benchmarks/single_node/agentic/dsv4_fp4_mi355x_vllm.sh @@ -31,6 +31,13 @@ source "$(dirname "$0")/../../benchmark_lib.sh" check_env_vars MODEL TP CONC KV_OFFLOADING TOTAL_CPU_DRAM_GB RESULT_DIR DURATION EP_SIZE DP_ATTENTION +GPU_COUNT=$TP +if [[ ! "$GPU_COUNT" =~ ^[1-9][0-9]*$ ]]; then + echo "Error: GPU_COUNT must be a positive integer, got '$GPU_COUNT'" >&2 + exit 1 +fi +export GPU_COUNT + if [[ -n "${SLURM_JOB_ID:-}" ]]; then echo "JOB $SLURM_JOB_ID running on ${SLURMD_NODENAME:-unknown}" fi @@ -109,7 +116,7 @@ case "${KV_OFFLOAD_BACKEND:-}" in # MI355X nodes have ~2.7 TiB of host DRAM available for offload; # reserve 2.5 TB for the offload pool (leaves ~200 GB headroom for # worker RSS / page cache / slurm cgroup). - TOTAL_CPU_DRAM_PARTITION_GB="$((TOTAL_CPU_DRAM_GB / (8 / TP)))" + TOTAL_CPU_DRAM_PARTITION_GB="$TOTAL_CPU_DRAM_GB" # Use vLLM's regular native KV-offload path (OffloadingConnector), # NOT the SimpleCPUOffloadConnector. The "vllm-native" backend resolves to # OffloadingConnector by default; setting VLLM_USE_SIMPLE_KV_OFFLOAD=1 @@ -132,7 +139,7 @@ case "${KV_OFFLOAD_BACKEND:-}" in # ---- Mooncake config ---------------------------------------------------------- # Embedded mode contributes one segment per GPU rank to a shared # distributed store, so pre-divide the aggregate host-memory budget. - PER_RANK_GB=$((TOTAL_CPU_DRAM_GB / TP)) + PER_RANK_GB=$((TOTAL_CPU_DRAM_GB / GPU_COUNT)) #MOONCAKE_VERSION=0.3.11.post1 #apt-get update && apt-get install -y libcurl4 libibverbs1 rdma-core librdmacm1 libnuma1 liburing2 @@ -276,16 +283,14 @@ EOF git clone https://github.com/LMCache/LMCache.git cd LMCache - # https://github.com/LMCache/LMCache/pull/3853 - git checkout 9229067cec0b3a63bb8a39368d101db7ac0bc3c1 + git checkout v0.5.2 pip install -r requirements/build.txt - pip install grpcio==1.78.0 CXX=hipcc BUILD_WITH_HIP=1 pip install -e . --no-build-isolation cd .. python3 -c "import lmcache.integration.vllm.lmcache_mp_connector" >/dev/null - TOTAL_CPU_DRAM_PARTITION_GB="$((TOTAL_CPU_DRAM_GB / (8 / TP)))" + TOTAL_CPU_DRAM_PARTITION_GB="$TOTAL_CPU_DRAM_GB" # Match the B200 Kimi LMCache setup: keep a 2.5 TB semantic CPU KV # pool, but let the external MP server own that pool so vLLM does not # split --kv-offloading-size across TP ranks through the integrated @@ -310,7 +315,7 @@ EOF # object present in L1 but no longer readable. Keep the 2.5 TB pool # size unchanged and only extend the lookup-to-retrieve lease. LMCACHE_L1_READ_TTL_SECONDS="${LMCACHE_L1_READ_TTL_SECONDS:-7200}" - LMCACHE_CHUNK_SIZE="${LMCACHE_CHUNK_SIZE:-256}" + LMCACHE_CHUNK_SIZE="${LMCACHE_CHUNK_SIZE:-1024}" LMCACHE_MAX_WORKERS="${LMCACHE_MAX_WORKERS:-$TP}" export PYTHONHASHSEED="${PYTHONHASHSEED:-0}" export LMCACHE_BLOCKING_TIMEOUT_SECS=1200 @@ -351,33 +356,57 @@ EOF esac fi +# ---- LLM server config ---------------------------------------------------------- PARALLEL_ARGS=(--tensor-parallel-size "$TP" --data-parallel-size 1) if [ "$DP_ATTENTION" = "true" ]; then PARALLEL_ARGS=(--tensor-parallel-size 1 --data-parallel-size "$TP") fi +if [ "$EP_SIZE" -gt 1 ]; then + PARALLEL_ARGS+=(--enable-expert-parallel) +fi + +# DEP8 (TP8 + DP-attention) is a high-concurrency arm tuned separately from the +# smaller DEP arm: larger prefill token budget and long-prefill chunking so +# decode latency stays bounded at high concurrency. +IS_DEP8=false +if [ "$DP_ATTENTION" = "true" ] && [ "$TP" -eq 8 ]; then + IS_DEP8=true +fi -EP_ARGS=() +MODE_ARGS=() +MNBT=8192 if [ "$EP_SIZE" -gt 1 ]; then - EP_ARGS=(--enable-expert-parallel) + MODE_ARGS+=(--enable-ep-weight-filter) +fi +if [ "$DP_ATTENTION" = "true" ]; then + MODE_ARGS+=(--prefill-schedule-interval 8) + if [ "$IS_DEP8" = "true" ]; then + MODE_ARGS+=( + --max-num-batched-tokens $((MNBT * 8)) + --long-prefill-token-threshold 16384 + ) + else + MODE_ARGS+=(--max-num-batched-tokens $MNBT) + fi fi -# AgentX concurrency counts live session trees, not individual requests. -# Subagent fan-out can push instantaneous request concurrency above CONC, so -# leave 2x headroom rather than clipping those bursts at the scheduler. +#if [ "$DP_ATTENTION" = "true" ]; then +# # The DEP source recipe enforces 2*CONC = DP_WORLD_SIZE*MAX_NUM_SEQS. +# MAX_NUM_SEQS=$((2 * CONC / TP)) +#else +# # Preserve the previous TP4 scheduler headroom for agentic fan-out. +# MAX_NUM_SEQS=$((2 * CONC)) +#fi MAX_NUM_SEQS=$((2 * CONC)) +CONTEXT_LEN=1048576 echo "Starting vllm server..." set -x export VLLM_ROCM_USE_AITER=1 -#export VLLM_ROCM_QUICK_REDUCE_QUANTIZATION=INT4 -export VLLM_ROCM_USE_AITER_MOE=1 +export VLLM_ROCM_QUICK_REDUCE_QUANTIZATION=INT4 sleep 180 -# https://github.com/vllm-project/vllm/pull/45497 -git clone https://gist.github.com/seungrokj/a37ff4d9a52db31752e2d5fa5b192e00 -cp a37ff4d9a52db31752e2d5fa5b192e00/gistfile1.txt /usr/local/lib/python3.12/dist-packages/vllm/v1/core/sched/scheduler.py - { set +x; } 2>/dev/null VLLM_CMD=( vllm serve "$MODEL_PATH" --served-model-name "$MODEL" @@ -387,11 +416,13 @@ VLLM_CMD=( --async-scheduling --distributed-executor-backend mp --kv-cache-dtype fp8 + --block-size 256 + --max-model-len "$CONTEXT_LEN" "${PARALLEL_ARGS[@]}" - "${EP_ARGS[@]}" - --gpu-memory-utilization 0.8 + "${MODE_ARGS[@]}" + --gpu-memory-utilization 0.95 --moe-backend aiter - --compilation-config '{"mode":3,"cudagraph_mode":"FULL_AND_PIECEWISE"}' + --compilation-config '{"mode":3,"cudagraph_mode":"FULL_DECODE_ONLY"}' --tokenizer-mode deepseek_v4 --tool-call-parser deepseek_v4 --reasoning-parser deepseek_v4 diff --git a/configs/amd-master.yaml b/configs/amd-master.yaml index 8e6bee505e..49098872f9 100644 --- a/configs/amd-master.yaml +++ b/configs/amd-master.yaml @@ -1516,7 +1516,7 @@ qwen3.5-fp8-mi355x-sglang-agentic-hicache: # DP-attention on/off and EP=8. dsv4-fp4-mi355x-vllm-agentic: - image: vllm/vllm-openai-rocm:nightly-09663abde0f50944a8d5ea30120666024b503faa + image: vllm/vllm-openai-rocm:nightly-49f31d7cee425a6d38f8c5bc76877986daf832ed model: deepseek-ai/DeepSeek-V4-Pro model-prefix: dsv4 runner: cluster:mi355x-amds @@ -1525,11 +1525,12 @@ dsv4-fp4-mi355x-vllm-agentic: multinode: false scenarios: agentic-coding: - - dram-utilization: 0.60 + - dram-utilization: 0.80 search-space: - - { tp: 8, ep: 1, dp-attn: false, kv-offloading: none, conc-list: [1, 4, 8, 16, 32, 40, 48] } - - { tp: 8, ep: 1, dp-attn: true, kv-offloading: none, conc-list: [64, 72], router: { name: vllm-router, version: "0.1.14" } } - - { tp: 8, ep: 1, dp-attn: false, kv-offloading: dram, kv-offload-backend: { name: lmcache, version: "9229067cec0b3a63bb8a39368d101db7ac0bc3c1" }, conc-list: [32, 40, 48] } + #- { tp: 8, ep: 1, dp-attn: false, kv-offloading: none, conc-list: [1, 4, 8, 16, 32, 40, 48] } + #- { tp: 8, ep: 1, dp-attn: true, kv-offloading: none, conc-list: [64, 72], router: { name: vllm-router, version: "0.1.14" } } + #- { tp: 8, ep: 1, dp-attn: false, kv-offloading: dram, kv-offload-backend: { name: lmcache, version: "9229067cec0b3a63bb8a39368d101db7ac0bc3c1" }, conc-list: [32, 40, 48] } + - { tp: 8, ep: 1, dp-attn: true, kv-offloading: dram, kv-offload-backend: { name: lmcache, version: "9229067cec0b3a63bb8a39368d101db7ac0bc3c1" }, conc-list: [76] } # DSv4-Pro FP4 on MI355X via SGLang. Uses a rocm720 mi35x image built off the # amd/deepseek_v4 branch in sgl-project/sglang; the SHA is encoded in the From e10d502d2d871a912f96a7a720bc10a5a225d8cc Mon Sep 17 00:00:00 2001 From: seungrokj Date: Tue, 28 Jul 2026 14:21:34 +0900 Subject: [PATCH 2/8] docs(perf-changelog): add dsv4-fp4-mi355x-vllm-agentic entry for #2378 Co-Authored-By: Claude Opus 4.6 --- perf-changelog.yaml | 12 +++++++++++- 1 file changed, 11 insertions(+), 1 deletion(-) diff --git a/perf-changelog.yaml b/perf-changelog.yaml index d969b6ce47..e23e6df31b 100644 --- a/perf-changelog.yaml +++ b/perf-changelog.yaml @@ -5128,4 +5128,14 @@ description: - "Re-sweep dsv4-fp4-mi355x-atom-mtp on latest atom-dev nightly (nightly_202607231538, was atom0.1.3 release). Align ATOM official MTP config (models.json): MTP3 dp-off conc 4-256 + DPA MTP3 (dp-attn) conc 32-1024 (added c32 dp-on; tp4 tried but removed — tp4+MTP+dp-on OOMs, KV budget negative). prefill-only TBO (--enable-tbo, argparse const=prefill so enable_tbo_decode=False) on dp-attn cells at conc>=256 (measured crossover run 30257759947 vs non-TBO 30238071409: TBO -10~14% output tput at c64/c128 but +8~14% at c256+; MTP removes low-conc latency so TBO overlap there is pure overhead); MTP-safe because only decode-TBO (--enable-tbo all) drops spec_decode_metadata. max_num_seqs=conc on dp-on cells and conc>=64 to avoid OOM. bench keeps --dsv4 (InferenceX bench uses encoding_dsv4.py; DSv4-Pro has no jinja chat_template). Fixed benchmark_lib source path (../ -> ../../)." pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2345 - + +- config-keys: + - dsv4-fp4-mi355x-vllm-agentic + description: + - "Tune the DEP arm: add MODE_ARGS with --enable-ep-weight-filter (EP>1), --prefill-schedule-interval 8 (dp-attn), and a DEP8 (TP8+dp-attn) prefill budget of --max-num-batched-tokens 65536 (8x the 8192 base) + --long-prefill-token-threshold 16384; non-DEP8 dp-attn keeps --max-num-batched-tokens 8192. Merge --enable-expert-parallel into PARALLEL_ARGS." + - "Serve flags: add --block-size 256 and --max-model-len 1048576; switch --compilation-config cudagraph_mode FULL_AND_PIECEWISE -> FULL_DECODE_ONLY; --gpu-memory-utilization 0.8 -> 0.95; enable VLLM_ROCM_QUICK_REDUCE_QUANTIZATION=INT4; drop the gist scheduler.py monkeypatch and the VLLM_ROCM_USE_AITER_MOE export." + - "LMCache: build tagged v0.5.2 (was the PR-3853 SHA 9229067) and drop the grpcio==1.78.0 pin; --chunk-size 256 -> 1024; give the external MP server the undivided CPU pool (TOTAL_CPU_DRAM_PARTITION_GB = full TOTAL_CPU_DRAM_GB instead of the /(8/TP) partition)." + - "Bug fix: TOTAL_CPU_DRAM_PARTITION_GB=\"$(TOTAL_CPU_DRAM_GB)\" (command substitution, yielded empty) -> \"$TOTAL_CPU_DRAM_GB\" in the vllm-native and lmcache branches." + - "amd-master.yaml: bump image to nightly-49f31d7cee425a6d38f8c5bc76877986daf832ed; dram-utilization 0.60 -> 0.80; switch agentic search-space to a single DEP lmcache arm at conc 76 (prior GPU-resident and dp-off arms commented out)." + pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2378 + From 52373498f10bf53a9a32b4773575a7304fb7b8a6 Mon Sep 17 00:00:00 2001 From: seungrokj Date: Tue, 28 Jul 2026 14:27:53 +0900 Subject: [PATCH 3/8] [AMD] [AGENTX] dsv4-fp4-mi355x-vllm-agentic: pin lmcache to v0.5.2 in master-config Align the agentic search-space kv-offload-backend version with the recipe (build tag v0.5.2, was the PR-3853 SHA 9229067). Co-Authored-By: Claude Opus 4.6 --- configs/amd-master.yaml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/configs/amd-master.yaml b/configs/amd-master.yaml index 49098872f9..0827a30060 100644 --- a/configs/amd-master.yaml +++ b/configs/amd-master.yaml @@ -1530,7 +1530,7 @@ dsv4-fp4-mi355x-vllm-agentic: #- { tp: 8, ep: 1, dp-attn: false, kv-offloading: none, conc-list: [1, 4, 8, 16, 32, 40, 48] } #- { tp: 8, ep: 1, dp-attn: true, kv-offloading: none, conc-list: [64, 72], router: { name: vllm-router, version: "0.1.14" } } #- { tp: 8, ep: 1, dp-attn: false, kv-offloading: dram, kv-offload-backend: { name: lmcache, version: "9229067cec0b3a63bb8a39368d101db7ac0bc3c1" }, conc-list: [32, 40, 48] } - - { tp: 8, ep: 1, dp-attn: true, kv-offloading: dram, kv-offload-backend: { name: lmcache, version: "9229067cec0b3a63bb8a39368d101db7ac0bc3c1" }, conc-list: [76] } + - { tp: 8, ep: 1, dp-attn: true, kv-offloading: dram, kv-offload-backend: { name: lmcache, version: "v0.5.2" }, conc-list: [76] } # DSv4-Pro FP4 on MI355X via SGLang. Uses a rocm720 mi35x image built off the # amd/deepseek_v4 branch in sgl-project/sglang; the SHA is encoded in the From c605e763eaf787ba549f4422e13e5fa8d54ff841 Mon Sep 17 00:00:00 2001 From: seungrokj Date: Tue, 28 Jul 2026 14:37:08 +0900 Subject: [PATCH 4/8] [AMD] [AGENTX] dsv4-fp4-mi355x-vllm-agentic: size max-num-seqs per DP rank --max-num-seqs is applied per scheduler. Under DP-attention each of the TP DP ranks runs its own scheduler (router spreads sessions across them), so use 2*CONC/TP (aggregate 2*CONC); pure-TP keeps 2*CONC. Mirrors the B300 recipe. Co-Authored-By: Claude Opus 4.6 --- .../agentic/dsv4_fp4_mi355x_vllm.sh | 18 ++++++++++-------- 1 file changed, 10 insertions(+), 8 deletions(-) diff --git a/benchmarks/single_node/agentic/dsv4_fp4_mi355x_vllm.sh b/benchmarks/single_node/agentic/dsv4_fp4_mi355x_vllm.sh index 5792e1e4d2..93b949c9fb 100644 --- a/benchmarks/single_node/agentic/dsv4_fp4_mi355x_vllm.sh +++ b/benchmarks/single_node/agentic/dsv4_fp4_mi355x_vllm.sh @@ -390,14 +390,16 @@ if [ "$DP_ATTENTION" = "true" ]; then fi fi -#if [ "$DP_ATTENTION" = "true" ]; then -# # The DEP source recipe enforces 2*CONC = DP_WORLD_SIZE*MAX_NUM_SEQS. -# MAX_NUM_SEQS=$((2 * CONC / TP)) -#else -# # Preserve the previous TP4 scheduler headroom for agentic fan-out. -# MAX_NUM_SEQS=$((2 * CONC)) -#fi -MAX_NUM_SEQS=$((2 * CONC)) +# --max-num-seqs is applied PER scheduler. Under DP-attention each of the $TP +# DP ranks runs its own scheduler and the router spreads sessions across them, +# so size the per-rank cap as 2*CONC/TP (aggregate = 2*CONC). In pure-TP there +# is a single scheduler across all GPUs that sees all CONC sessions, so use +# 2*CONC directly. +if [ "$DP_ATTENTION" = "true" ]; then + MAX_NUM_SEQS=$((2 * CONC / TP)) +else + MAX_NUM_SEQS=$((2 * CONC)) +fi CONTEXT_LEN=1048576 echo "Starting vllm server..." From 4bfe42a33e0b8a9b8e15c19f3d660d9190e11d20 Mon Sep 17 00:00:00 2001 From: seungrokj Date: Tue, 28 Jul 2026 14:59:22 +0900 Subject: [PATCH 5/8] [AMD] [AGENTX] dsv4-fp4-mi355x-vllm-agentic: derate gpu-mem-util to 0.85 for lmcache The LMCache MP server (kv_buffer_device=cuda) holds ~32 GiB GPU staging memory per device before vLLM starts (only ~256/288 GiB free at init), so --gpu-memory-utilization 0.95 requested 273.6 GiB and every DP worker failed init_device with "Free memory ... less than desired GPU memory utilization". Derate the lmcache arm to 0.85 (244.8 GiB, ~11 GiB headroom); other arms keep 0.95. Co-Authored-By: Claude Opus 4.6 --- .../single_node/agentic/dsv4_fp4_mi355x_vllm.sh | 14 +++++++++++++- 1 file changed, 13 insertions(+), 1 deletion(-) diff --git a/benchmarks/single_node/agentic/dsv4_fp4_mi355x_vllm.sh b/benchmarks/single_node/agentic/dsv4_fp4_mi355x_vllm.sh index 93b949c9fb..dba3bfa271 100644 --- a/benchmarks/single_node/agentic/dsv4_fp4_mi355x_vllm.sh +++ b/benchmarks/single_node/agentic/dsv4_fp4_mi355x_vllm.sh @@ -57,6 +57,7 @@ fi if [ -n "${ROCR_VISIBLE_DEVICES:-}" ]; then export HIP_VISIBLE_DEVICES="$ROCR_VISIBLE_DEVICES" fi +rocm-smi # ---- Resolve traces and install deps ---------------------------------------- resolve_trace_source @@ -402,6 +403,17 @@ else fi CONTEXT_LEN=1048576 +# The LMCache MP server (kv_buffer_device=cuda) holds GPU staging memory on +# every device before vLLM starts, so only ~256/288 GiB is free at init. At +# --gpu-memory-utilization 0.95 vLLM requests 273.6 GiB and hard-fails +# ("Free memory ... less than desired GPU memory utilization"). Derate the +# lmcache arm to leave headroom above the LMCache footprint; other arms keep +# the full 0.95. +GPU_MEM_UTIL=0.95 +if [ "${KV_OFFLOAD_BACKEND:-}" = "lmcache" ]; then + GPU_MEM_UTIL=0.85 +fi + echo "Starting vllm server..." set -x export VLLM_ROCM_USE_AITER=1 @@ -422,7 +434,7 @@ VLLM_CMD=( --max-model-len "$CONTEXT_LEN" "${PARALLEL_ARGS[@]}" "${MODE_ARGS[@]}" - --gpu-memory-utilization 0.95 + --gpu-memory-utilization "$GPU_MEM_UTIL" --moe-backend aiter --compilation-config '{"mode":3,"cudagraph_mode":"FULL_DECODE_ONLY"}' --tokenizer-mode deepseek_v4 From 2da5c487ac44547418586f0304698e142eaa48c9 Mon Sep 17 00:00:00 2001 From: seungrokj Date: Tue, 28 Jul 2026 18:41:20 +0900 Subject: [PATCH 6/8] [AMD] [AGENTX] dsv4-fp4-mi355x-vllm-agentic: force lm-eval; switch arm to GPU-resident c48 - dsv4_fp4_mi355x_vllm.sh: export EVAL_FRAMEWORK=lm-eval so run_eval never falls back to the agentic swebench default. - amd-master.yaml: switch the active search-space arm to TP8/DPA GPU-resident (kv-offloading: none) conc 48; comment out the lmcache dram c76 arm. Co-Authored-By: Claude Opus 4.6 --- benchmarks/single_node/agentic/dsv4_fp4_mi355x_vllm.sh | 7 +++++++ configs/amd-master.yaml | 3 ++- 2 files changed, 9 insertions(+), 1 deletion(-) diff --git a/benchmarks/single_node/agentic/dsv4_fp4_mi355x_vllm.sh b/benchmarks/single_node/agentic/dsv4_fp4_mi355x_vllm.sh index dba3bfa271..5708537fb5 100644 --- a/benchmarks/single_node/agentic/dsv4_fp4_mi355x_vllm.sh +++ b/benchmarks/single_node/agentic/dsv4_fp4_mi355x_vllm.sh @@ -29,6 +29,13 @@ set -x source "$(dirname "$0")/../../benchmark_lib.sh" +# Force the eval framework to lm-eval for this recipe. run_eval derives its +# default as swebench for agentic scenarios (scenario_default=swebench when +# IS_AGENTIC/SCENARIO_TYPE=agentic-coding), but EVAL_FRAMEWORK takes precedence +# over that default (benchmark_lib.sh: framework=${EVAL_FRAMEWORK:-...}), so +# setting it here makes the effective framework always lm-eval, never swebench. +export EVAL_FRAMEWORK="lm-eval" + check_env_vars MODEL TP CONC KV_OFFLOADING TOTAL_CPU_DRAM_GB RESULT_DIR DURATION EP_SIZE DP_ATTENTION GPU_COUNT=$TP diff --git a/configs/amd-master.yaml b/configs/amd-master.yaml index 0827a30060..e9f08cbc13 100644 --- a/configs/amd-master.yaml +++ b/configs/amd-master.yaml @@ -1530,7 +1530,8 @@ dsv4-fp4-mi355x-vllm-agentic: #- { tp: 8, ep: 1, dp-attn: false, kv-offloading: none, conc-list: [1, 4, 8, 16, 32, 40, 48] } #- { tp: 8, ep: 1, dp-attn: true, kv-offloading: none, conc-list: [64, 72], router: { name: vllm-router, version: "0.1.14" } } #- { tp: 8, ep: 1, dp-attn: false, kv-offloading: dram, kv-offload-backend: { name: lmcache, version: "9229067cec0b3a63bb8a39368d101db7ac0bc3c1" }, conc-list: [32, 40, 48] } - - { tp: 8, ep: 1, dp-attn: true, kv-offloading: dram, kv-offload-backend: { name: lmcache, version: "v0.5.2" }, conc-list: [76] } + #- { tp: 8, ep: 1, dp-attn: true, kv-offloading: dram, kv-offload-backend: { name: lmcache, version: "v0.5.2" }, conc-list: [76] } + - { tp: 8, ep: 1, dp-attn: true, kv-offloading: none, conc-list: [48] } # DSv4-Pro FP4 on MI355X via SGLang. Uses a rocm720 mi35x image built off the # amd/deepseek_v4 branch in sgl-project/sglang; the SHA is encoded in the From a649ac418b4dbba0c5b9de4364def72079872d6b Mon Sep 17 00:00:00 2001 From: seungrokj Date: Tue, 28 Jul 2026 20:08:21 +0900 Subject: [PATCH 7/8] [AMD] [AGENTX] dsv4-fp4-mi355x-vllm-agentic: cut DEP8 max-num-batched-tokens to MNBT*2 MNBT*8 (65536) drove a bf16 einsum/bmm of M=65536 at warmup that segfaulted inside rocBLASLt/TensileLite (hipblasLtMatmulAlgoGetHeuristic -> loadPlaceholderLibrary). Drop the DEP8 prefill budget to MNBT*2 (16384), matching the validated B300 recipe, to avoid the oversized GEMM shape. Co-Authored-By: Claude Opus 4.6 --- benchmarks/single_node/agentic/dsv4_fp4_mi355x_vllm.sh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/benchmarks/single_node/agentic/dsv4_fp4_mi355x_vllm.sh b/benchmarks/single_node/agentic/dsv4_fp4_mi355x_vllm.sh index 5708537fb5..f7da7f4352 100644 --- a/benchmarks/single_node/agentic/dsv4_fp4_mi355x_vllm.sh +++ b/benchmarks/single_node/agentic/dsv4_fp4_mi355x_vllm.sh @@ -390,7 +390,7 @@ if [ "$DP_ATTENTION" = "true" ]; then MODE_ARGS+=(--prefill-schedule-interval 8) if [ "$IS_DEP8" = "true" ]; then MODE_ARGS+=( - --max-num-batched-tokens $((MNBT * 8)) + --max-num-batched-tokens $((MNBT * 2)) --long-prefill-token-threshold 16384 ) else From 554c46a421051e10224e7234498a3f01991c549e Mon Sep 17 00:00:00 2001 From: seungrokj Date: Tue, 28 Jul 2026 21:01:33 +0900 Subject: [PATCH 8/8] [AMD] [AGENTX] dsv4-fp4-mi355x-vllm-agentic: cap gpu-mem-util at 0.85 for all arms The GPU-resident (kv-none) c48 arm also hit the startup OOM: these MI355X nodes have a stable ~32 GiB/GPU carveout (only ~256/288 GiB free at init, independent of the KV-offload backend), so 0.95 (273.6 GiB) fails every DP worker. Collapse the lmcache-only derate into a flat GPU_MEM_UTIL=0.85 (244.8 GiB, ~11 GiB headroom) covering both the lmcache and GPU-resident arms. Co-Authored-By: Claude Opus 4.6 --- .../single_node/agentic/dsv4_fp4_mi355x_vllm.sh | 17 +++++++---------- 1 file changed, 7 insertions(+), 10 deletions(-) diff --git a/benchmarks/single_node/agentic/dsv4_fp4_mi355x_vllm.sh b/benchmarks/single_node/agentic/dsv4_fp4_mi355x_vllm.sh index f7da7f4352..59e03089e5 100644 --- a/benchmarks/single_node/agentic/dsv4_fp4_mi355x_vllm.sh +++ b/benchmarks/single_node/agentic/dsv4_fp4_mi355x_vllm.sh @@ -410,16 +410,13 @@ else fi CONTEXT_LEN=1048576 -# The LMCache MP server (kv_buffer_device=cuda) holds GPU staging memory on -# every device before vLLM starts, so only ~256/288 GiB is free at init. At -# --gpu-memory-utilization 0.95 vLLM requests 273.6 GiB and hard-fails -# ("Free memory ... less than desired GPU memory utilization"). Derate the -# lmcache arm to leave headroom above the LMCache footprint; other arms keep -# the full 0.95. -GPU_MEM_UTIL=0.95 -if [ "${KV_OFFLOAD_BACKEND:-}" = "lmcache" ]; then - GPU_MEM_UTIL=0.85 -fi +# These MI355X nodes have a stable ~32 GiB/GPU carveout: only ~256/288 GiB is +# free at init, independent of the KV-offload backend (observed identically on +# both the lmcache and GPU-resident kv-none arms). At +# --gpu-memory-utilization 0.95 vLLM requests 273.6 GiB and every DP worker +# hard-fails ("Free memory ... less than desired GPU memory utilization"), so +# cap utilization at 0.85 (244.8 GiB, ~11 GiB headroom). +GPU_MEM_UTIL=0.85 echo "Starting vllm server..." set -x