From fec0ceefe492df3be97c512a00973dca4cbf86d6 Mon Sep 17 00:00:00 2001
From: seungrokj <seungrok.jung@amd.com>
Date: Fri, 10 Jul 2026 15:51:16 +0900
Subject: [PATCH 01/11] [AMD][MI355X] Update DSv4 SGLang agentic launcher to
 v0.5.14 image with router and HiCache support

Align the agentic launcher with the fixed-seq-len sibling: dsv4 attention
backend, fp8_e4m3 KV cache, disable-radix-cache, DP-attention env exports
(SGLANG_SHARED_EXPERT_TP1, DP_USE_GATHERV, etc.), and two-batch-overlap.
Add SGLang router for DP-attention configs and HiCache KV offloading.
Sweep conc=48 across TP8 +/- DPA +/- HiCache.

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
---
 .../agentic/dsv4_fp4_mi355x_sglang.sh         | 240 +++++++++++-------
 configs/amd-master.yaml                       |  13 +-
 perf-changelog.yaml                           |  11 +
 3 files changed, 164 insertions(+), 100 deletions(-)

diff --git a/benchmarks/single_node/agentic/dsv4_fp4_mi355x_sglang.sh b/benchmarks/single_node/agentic/dsv4_fp4_mi355x_sglang.sh
index d8cba43f05..3cbebb5bbd 100755
--- a/benchmarks/single_node/agentic/dsv4_fp4_mi355x_sglang.sh
+++ b/benchmarks/single_node/agentic/dsv4_fp4_mi355x_sglang.sh
@@ -3,14 +3,14 @@ set -euo pipefail
 set -x
 
 # Agentic trace replay benchmark for DeepSeek-V4-Pro FP4 on MI355X using SGLang.
-# Adapted from benchmarks/single_node/dsv4_fp4_mi355x_sglang.sh (fixed-seq-len
-# sibling) with the agentic harness (build_replay_cmd / write_agentic_result_json
-# / analyze_benchmark_distributions) swapped in for run_benchmark_serving.
 #
-# This launcher only supports on-device KV cache.
+# KV_OFFLOADING=dram requires KV_OFFLOAD_BACKEND=hicache.
 #
 # Required env vars:
 #   MODEL, TP, CONC, KV_OFFLOADING, TOTAL_CPU_DRAM_GB, RESULT_DIR
+#
+# KV_OFFLOADING=dram requires one of these. 
+#   KV_OFFLOAD_BACKEND=hicache.
 
 source "$(dirname "$0")/../../benchmark_lib.sh"
 
@@ -43,115 +43,171 @@ amd-smi || true
 resolve_trace_source
 install_agentic_deps
 
-require_agentic_kv_offload_none
-
-# Transformers in the container doesn't recognize the `deepseek_v4` model_type.
-# PR #23608's fallback in hf_transformers_utils.get_config tries to handle this
-# by writing a patched config to /tmp, but in practice isn't catching the error
-# in this image. Patch the cached config.json directly instead: set model_type
-# to `deepseek_v3` so AutoConfig.from_pretrained succeeds, and keep
-# architectures=['DeepseekV4ForCausalLM'] so SGLang dispatches to its native
-# DSv4 model class (python/sglang/srt/models/deepseek_v4.py).
-python3 << PYEOF
-import json
-from huggingface_hub import hf_hub_download
-path = hf_hub_download(repo_id="$MODEL", filename="config.json")
-with open(path) as f:
-    config = json.load(f)
-if config.get("model_type") == "deepseek_v4":
-    config["model_type"] = "deepseek_v3"
-    with open(path, "w") as f:
-        json.dump(config, f, indent=2)
-    print(f"Patched {path}: model_type deepseek_v4 -> deepseek_v3")
-else:
-    print(f"No patch needed: model_type is {config.get('model_type')!r}")
-PYEOF
-
-# DSv4 FP4-experts path. Mirrors the env block in the fixed-seq-len sibling
-# (benchmarks/single_node/dsv4_fp4_mi355x_sglang.sh), which tracks the active
-# block in python/run_dsv4.sh on the amd/deepseek_v4 branch:
-#   SGLANG_DSV4_FP4_EXPERTS=True   -> route experts through FP4 kernels
-#   SGLANG_FORCE_TRITON_MOE_FP8=0  -> dispatch MoE through aiter and apply
-#                                    the swiglu_limit clamp in the triton
-#                                    MoE fallback path.
-export SGLANG_REASONING_EFFORT=max
-export SGLANG_OPT_USE_FUSED_COMPRESS=true
-export SGLANG_OPT_USE_OLD_COMPRESSOR=true
-export SGLANG_OPT_USE_TILELANG_SWA_PREPARE=false
-export SGLANG_OPT_USE_JIT_KERNEL_FUSED_TOPK=false
-export SGLANG_OPT_USE_FUSED_HASH_TOPK=false
-export SGLANG_OPT_DEEPGEMM_HC_PRENORM=false
-export SGLANG_OPT_USE_TILELANG_MHC_PRE=false
-export SGLANG_OPT_USE_TILELANG_MHC_POST=false
-export SGLANG_OPT_USE_AITER_MHC_PRE=true
-export SGLANG_OPT_USE_AITER_MHC_POST=true
-export SGLANG_ENABLE_THINKING=1
-export SGLANG_USE_AITER=1
-export SGLANG_USE_ROCM700A=1
-export SGLANG_TOPK_TRANSFORM_512_TORCH=0
-export SGLANG_FP8_PAGED_MQA_LOGITS_TORCH=1
-export SGLANG_DSV4_FP4_EXPERTS=True
-export SGLANG_OPT_DPSK_V4_RADIX=0
-export SGLANG_OPT_USE_OVERLAP_STORE_CACHE=false
-export SGLANG_OPT_USE_FUSED_STORE_CACHE=false
-export SGLANG_FORCE_TRITON_MOE_FP8=0
-export SGLANG_HACK_FLASHMLA_BACKEND=tilelang
-export SGLANG_OPT_USE_TILELANG_INDEXER=true
-export SGLANG_OPT_USE_TRITON_SWA_PREPARE=true
-
 # ---- Server config ----------------------------------------------------------
 SERVER_LOG="$RESULT_DIR/server.log"
 mkdir -p "$RESULT_DIR"
 
-# Parallelism: pure TP, TP+EP, or DEP (DP-attn + EP). Matches the dsv4 b200
-# vllm agentic launcher so the agentic sweep can probe both interactivity and
-# throughput regimes.
+export SGLANG_ENABLE_UNIFIED_RADIX_TREE=1
+export SGLANG_OPT_UNIFIED_CACHE_FREE_OUT_OF_WINDOW_SLOTS=1
+
+CACHE_ARGS=()
+
+if agentic_kv_offload_enabled; then
+case "${KV_OFFLOAD_BACKEND:-}" in
+  hicache)
+        # ---- Hicache config ----------------------------------------------------------
+        # DSv4 allocates several physical host sub-pools for each logical host
+        # token. MI355X nodes have ~3 TB of host DRAM (similar to B200's 3.8
+        # TiB), so ratio=8 at TP≥8 provides a large useful CPU tier within the
+        # node budget. Lower TP configs use higher ratios to maintain adequate
+        # host token capacity without exceeding DRAM limits.
+
+        DEFAULT_HICACHE_RATIO=8
+        HICACHE_RATIO="${HICACHE_RATIO:-$DEFAULT_HICACHE_RATIO}"
+        if [ "$HICACHE_RATIO" -gt "$DEFAULT_HICACHE_RATIO" ]; then
+            echo "Error: HICACHE_RATIO=$HICACHE_RATIO exceeds configured limit $DEFAULT_HICACHE_RATIO" >&2
+            exit 1
+        fi
+        HICACHE_WRITE_POLICY="${HICACHE_WRITE_POLICY:-write_through}"
+        HICACHE_IO_BACKEND="${HICACHE_IO_BACKEND:-direct}"
+        HICACHE_MEM_LAYOUT="${HICACHE_MEM_LAYOUT:-page_first_direct}"
+        export SGLANG_ENABLE_UNIFIED_RADIX_TREE=1
+        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"
+        )
+        echo "HiCache DSv4 CPU tier: ratio=$HICACHE_RATIO, write_policy=$HICACHE_WRITE_POLICY, io_backend=$HICACHE_IO_BACKEND, mem_layout=$HICACHE_MEM_LAYOUT"
+        ;;
+  *)
+    echo "Error: unsupported KV_OFFLOAD_BACKEND '${KV_OFFLOAD_BACKEND:-}' (expected: native, mooncake, lmcache)" >&2
+    exit 1
+    ;;
+esac
+fi
+
+USE_SGLANG_ROUTER=false
+SGLANG_BACKEND_PORT="$PORT"
+ROUTER_LOG="$RESULT_DIR/router.log"
+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))
+fi
+
+# ---- LLM server config ----------------------------------------------------------
+
+CUDA_GRAPH_MAX_BS="$CONC"
+[ "$CUDA_GRAPH_MAX_BS" -gt 64 ] && CUDA_GRAPH_MAX_BS=64
+
+export SGLANG_DEFAULT_THINKING=1
+export SGLANG_DSV4_REASONING_EFFORT=max
+export SGLANG_USE_ROCM700A=0
+export SGLANG_HACK_FLASHMLA_BACKEND=unified_kv_triton
+export AITER_BF16_FP8_MOE_BOUND=0
+
 PARALLEL_ARGS=(--tensor-parallel-size "$TP")
+CHUNKED_PREFILL_SIZE=8192
 if [ "$DP_ATTENTION" = "true" ]; then
+    export SGLANG_SHARED_EXPERT_TP1=1
+    export SGLANG_DP_SHARED_EXPERT_LOCAL=1
+    export SGLANG_DP_USE_GATHERV=1
+    export SGLANG_DP_USE_REDUCE_SCATTER=1
+    export GPU_MAX_HW_QUEUES=5
+
+    CHUNKED_PREFILL_SIZE=$((8192 * TP))
     PARALLEL_ARGS+=(
         --dp "$TP"
         --enable-dp-attention
         --enable-prefill-delayer
+        --enable-two-batch-overlap
     )
 fi
 if [ "${EP_SIZE:-1}" -gt 1 ]; then
     PARALLEL_ARGS+=(--ep-size "$EP_SIZE")
 fi
 
-# --max-running-requests is per-engine. With DP-attn each DP engine handles
-# only CONC/$TP sequences in steady state (the agentic harness load-balances
-# users across DP ranks), so size the per-engine cap to that.
-# Pure TP is a single engine and sees all CONC sequences itself.
-if [ "$DP_ATTENTION" = "true" ]; then
-    PER_ENGINE_MAX_RUNNING=$(( CONC / TP ))
-    [ "$PER_ENGINE_MAX_RUNNING" -lt 1 ] && PER_ENGINE_MAX_RUNNING=1
-else
-    PER_ENGINE_MAX_RUNNING=$CONC
-fi
-
-echo "Starting sglang server..."
-python3 -m sglang.launch_server \
-    --model-path "$MODEL_PATH" --served-model-name "$MODEL" \
-    --host=0.0.0.0 \
-    --port "$PORT" \
-    "${PARALLEL_ARGS[@]}" \
-    --trust-remote-code \
-    --attention-backend compressed \
-    --max-running-requests "$PER_ENGINE_MAX_RUNNING" \
-    --cuda-graph-max-bs "$PER_ENGINE_MAX_RUNNING" \
-    --page-size 256 \
-    --chunked-prefill-size 8192 \
-    --disable-shared-experts-fusion \
-    --tool-call-parser deepseekv4 \
-    --reasoning-parser deepseek-v4 \
-    --chat-template "$(dirname "$0")/../chat_templates/deepseek_v4_thinking.jinja" \
-    --watchdog-timeout 1800 > "$SERVER_LOG" 2>&1 &
+SGLANG_CMD=(
+    python3 -m sglang.launch_server
+    --model-path "$MODEL_PATH"
+    --served-model-name "$MODEL"
+    --host 0.0.0.0
+    --port "$SGLANG_BACKEND_PORT"
+    --trust-remote-code
+    "${PARALLEL_ARGS[@]}"
+    --disable-radix-cache
+    --attention-backend dsv4
+    --cuda-graph-max-bs "$CUDA_GRAPH_MAX_BS"
+    --max-running-requests "$CONC"
+    --mem-fraction-static 0.90
+    --swa-full-tokens-ratio 0.15
+    --page-size 256
+    --kv-cache-dtype fp8_e4m3
+    --context-length "$MAX_MODEL_LEN"
+    --chunked-prefill-size "$CHUNKED_PREFILL_SIZE"
+    --disable-shared-experts-fusion
+    --tool-call-parser deepseekv4
+    --reasoning-parser deepseek-v4
+    --chat-template "$(dirname "$0")/../chat_templates/deepseek_v4_thinking.jinja"
+    --watchdog-timeout 1800
+    --enable-metrics
+    "${CACHE_ARGS[@]}"
+)
+
+printf '%q ' "${SGLANG_CMD[@]}" | tee "$RESULT_DIR/sglang_command.txt"
+printf '\n' | tee -a "$RESULT_DIR/sglang_command.txt"
+
+{
+    echo "=== SGLANG_* env vars at launch ==="
+    env | grep -E '^SGLANG_' | sort
+    echo "==================================="
+} | tee "$SERVER_LOG"
+
+echo "Starting SGLang server for MI355X..."
+"${SGLANG_CMD[@]}" >> "$SERVER_LOG" 2>&1 &
 SERVER_PID=$!
 echo "Server PID: $SERVER_PID"
 
-wait_for_server_ready --port "$PORT" --server-log "$SERVER_LOG" --server-pid "$SERVER_PID"
+capture_cache_metrics() {
+    {
+        echo "=== SGLang cache metrics snapshot $(date --iso-8601=seconds) ==="
+        curl -fsS "http://localhost:$SGLANG_BACKEND_PORT/metrics" 2>/dev/null \
+            | grep -E '^(sglang:(cache_hit_rate|cached_tokens_total|prompt_tokens_total|hicache_host_used_tokens|hicache_host_total_tokens|token_usage|num_requests_running|num_requests_waiting))' \
+            || true
+        echo "============================================================"
+    } >> "$SERVER_LOG"
+}
+
+wait_for_server_ready --port "$SGLANG_BACKEND_PORT" --server-log "$SERVER_LOG" --server-pid "$SERVER_PID"
+
+if [ "$USE_SGLANG_ROUTER" = "true" ]; then
+    echo "Starting SGLang router on port $PORT for $TP DP ranks..."
+    python3 -m sglang_router.launch_router \
+        --worker-urls "http://localhost:$SGLANG_BACKEND_PORT" \
+        --policy consistent_hashing \
+        --request-id-headers x-correlation-id \
+        --dp-aware \
+        --host 0.0.0.0 \
+        --port "$PORT" \
+        --prometheus-host 127.0.0.1 \
+        --prometheus-port "$SGLANG_ROUTER_METRICS_PORT" \
+        --connect-timeout-secs 900 \
+        --request-timeout-secs 14400 \
+        --disable-health-check \
+        --disable-retries > "$ROUTER_LOG" 2>&1 &
+    ROUTER_PID=$!
+    echo "Router PID: $ROUTER_PID"
+    wait_for_server_ready --port "$PORT" --server-log "$ROUTER_LOG" --server-pid "$ROUTER_PID"
+fi
+
+capture_cache_metrics
+trap capture_cache_metrics EXIT
 
 # ---- Run benchmark ----------------------------------------------------------
 build_replay_cmd "$RESULT_DIR"
+REPLAY_CMD+=" --server-metrics http://localhost:$SGLANG_BACKEND_PORT/metrics"
 
 run_agentic_replay_and_write_outputs "$RESULT_DIR"
diff --git a/configs/amd-master.yaml b/configs/amd-master.yaml
index 687439a98f..f458606980 100644
--- a/configs/amd-master.yaml
+++ b/configs/amd-master.yaml
@@ -2266,13 +2266,8 @@ dsr1-fp4-mi355x-sglang-disagg-mtp:
           - "DECODE_NODES=1"
           - "DECODE_MTP_SIZE=1"
 
-# 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
-# image tag, so bumping sglang is just an image tag bump here. Sweeps
-# DP-attention on/off and EP=8.
-
 dsv4-fp4-mi355x-sglang-agentic:
-  image: rocm/sgl-dev:rocm720-mi35x-0363e6c-20260509-DSv4
+  image: lmsysorg/sglang-rocm:v0.5.14-rocm720-mi35x-20260710
   model: deepseek-ai/DeepSeek-V4-Pro
   model-prefix: dsv4
   runner: cluster:mi355x-amds
@@ -2282,8 +2277,10 @@ dsv4-fp4-mi355x-sglang-agentic:
   scenarios:
     agentic-coding:
     - search-space:
-      - { tp: 8, kv-offloading: none, conc-list: [16, 32, 64] }
-      - { tp: 8, dp-attn: true, kv-offloading: none, conc-list: [64, 128, 256] }
+      - { tp: 8, kv-offloading: none, conc-list: [48] }
+      - { tp: 8, kv-offloading: dram, kv-offload-backend: hicache, conc-list: [48] }
+      - { tp: 8, dp-attn: true, kv-offloading: none, conc-list: [48] }
+      - { tp: 8, dp-attn: true, kv-offloading: dram, kv-offload-backend: hicache, conc-list: [48] }
 
 # MiniMax-M3 MXFP8 MI355X recipe:
 # https://github.com/vllm-project/recipes/commit/2a3728ed9892debfd767a72a58ebc90b33f186e5
diff --git a/perf-changelog.yaml b/perf-changelog.yaml
index c1a3c3aee8..e822cee68a 100644
--- a/perf-changelog.yaml
+++ b/perf-changelog.yaml
@@ -4699,3 +4699,14 @@
     - "Clean the export envs"
     - "Enable two batch overlap"
   pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2093
+
+- config-keys:
+    - dsv4-fp4-mi355x-sglang-agentic
+  description:
+    - "Bump image to lmsysorg/sglang-rocm:v0.5.14-rocm720-mi35x-20260710"
+    - "Align launcher env vars and server args with fixed-seq-len sibling (dsv4 attention backend, fp8_e4m3 kv-cache, disable-radix-cache, cuda-graph-max-bs, DP-attention exports, two-batch-overlap)"
+    - "Add SGLang router for DP-attention configs (consistent-hashing, dp-aware, correlation-id routing)"
+    - "Add HiCache KV offloading support"
+    - "Add SGLANG_CMD array pattern with command logging and env-var dump"
+    - "Add capture_cache_metrics for pre/post-benchmark cache stats"
+    - "Sweep conc=48 across TP8 +/- DPA +/- HiCache"

From 3fe2c5e37042109b74000b9691e1287fb26c27bc Mon Sep 17 00:00:00 2001
From: seungrokj <seungrok.jung@amd.com>
Date: Fri, 10 Jul 2026 16:24:32 +0900
Subject: [PATCH 02/11] [AMD][MI355X] Add dram-utilization 0.80 to DSv4 SGLang
 agentic sweep config

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
---
 configs/amd-master.yaml | 3 ++-
 1 file changed, 2 insertions(+), 1 deletion(-)

diff --git a/configs/amd-master.yaml b/configs/amd-master.yaml
index f458606980..0bea15f1fe 100644
--- a/configs/amd-master.yaml
+++ b/configs/amd-master.yaml
@@ -2276,7 +2276,8 @@ dsv4-fp4-mi355x-sglang-agentic:
   multinode: false
   scenarios:
     agentic-coding:
-    - search-space:
+    - dram-utilization: 0.80
+      search-space:
       - { tp: 8, kv-offloading: none, conc-list: [48] }
       - { tp: 8, kv-offloading: dram, kv-offload-backend: hicache, conc-list: [48] }
       - { tp: 8, dp-attn: true, kv-offloading: none, conc-list: [48] }

From e282bf705839b34e9df8d67dd27f810e82676ecf Mon Sep 17 00:00:00 2001
From: seungrokj <seungrok.jung@amd.com>
Date: Fri, 10 Jul 2026 16:27:12 +0900
Subject: [PATCH 03/11] Update perf-changelog with PR link for DSv4 SGLang
 agentic entry

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
---
 perf-changelog.yaml | 1 +
 1 file changed, 1 insertion(+)

diff --git a/perf-changelog.yaml b/perf-changelog.yaml
index e822cee68a..f78480edac 100644
--- a/perf-changelog.yaml
+++ b/perf-changelog.yaml
@@ -4710,3 +4710,4 @@
     - "Add SGLANG_CMD array pattern with command logging and env-var dump"
     - "Add capture_cache_metrics for pre/post-benchmark cache stats"
     - "Sweep conc=48 across TP8 +/- DPA +/- HiCache"
+  pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2146

From 6acad049038984f92a69ca353c34198544030786 Mon Sep 17 00:00:00 2001
From: seungrokj <seungrok.jung@amd.com>
Date: Fri, 10 Jul 2026 16:36:33 +0900
Subject: [PATCH 04/11] [AMD][MI355X] Remove --context-length and comment out
 extra sweep configs for initial validation

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
---
 benchmarks/single_node/agentic/dsv4_fp4_mi355x_sglang.sh | 1 -
 configs/amd-master.yaml                                  | 7 ++++---
 2 files changed, 4 insertions(+), 4 deletions(-)

diff --git a/benchmarks/single_node/agentic/dsv4_fp4_mi355x_sglang.sh b/benchmarks/single_node/agentic/dsv4_fp4_mi355x_sglang.sh
index 3cbebb5bbd..6dc348c66f 100755
--- a/benchmarks/single_node/agentic/dsv4_fp4_mi355x_sglang.sh
+++ b/benchmarks/single_node/agentic/dsv4_fp4_mi355x_sglang.sh
@@ -146,7 +146,6 @@ SGLANG_CMD=(
     --swa-full-tokens-ratio 0.15
     --page-size 256
     --kv-cache-dtype fp8_e4m3
-    --context-length "$MAX_MODEL_LEN"
     --chunked-prefill-size "$CHUNKED_PREFILL_SIZE"
     --disable-shared-experts-fusion
     --tool-call-parser deepseekv4
diff --git a/configs/amd-master.yaml b/configs/amd-master.yaml
index 0bea15f1fe..7cd5ede798 100644
--- a/configs/amd-master.yaml
+++ b/configs/amd-master.yaml
@@ -2279,9 +2279,10 @@ dsv4-fp4-mi355x-sglang-agentic:
     - dram-utilization: 0.80
       search-space:
       - { tp: 8, kv-offloading: none, conc-list: [48] }
-      - { tp: 8, kv-offloading: dram, kv-offload-backend: hicache, conc-list: [48] }
-      - { tp: 8, dp-attn: true, kv-offloading: none, conc-list: [48] }
-      - { tp: 8, dp-attn: true, kv-offloading: dram, kv-offload-backend: hicache, conc-list: [48] }
+      #- { tp: 8, kv-offloading: none, conc-list: [48] }
+      #- { tp: 8, kv-offloading: dram, kv-offload-backend: hicache, conc-list: [48] }
+      #- { tp: 8, dp-attn: true, kv-offloading: none, conc-list: [48] }
+      #- { tp: 8, dp-attn: true, kv-offloading: dram, kv-offload-backend: hicache, conc-list: [48] }
 
 # MiniMax-M3 MXFP8 MI355X recipe:
 # https://github.com/vllm-project/recipes/commit/2a3728ed9892debfd767a72a58ebc90b33f186e5

From 7e1c3321053b28cf3ebb920fe83152721ae2fb94 Mon Sep 17 00:00:00 2001
From: seungrokj <seungrok.jung@amd.com>
Date: Fri, 10 Jul 2026 17:55:53 +0900
Subject: [PATCH 05/11] [AMD][MI355X] DSv4 SGLang agentic: use compressed
 attention, 2x max-running-requests, re-enable all sweep configs

- Fix attention-backend typo (compresseed -> compressed)
- Remove duplicate CUDA_GRAPH_MAX_BS block
- Set MAX_RUNNING_REQUESTS to 2*CONC for subagent fan-out headroom
- Use MEM_FRACTION_STATIC=0.90 and swa-full-tokens-ratio=0.10
- Add METRICS_ARGS array, guard capture_cache_metrics on metrics enabled
- Remove duplicate SGLANG_ENABLE_UNIFIED_RADIX_TREE export from hicache block
- Uncomment all 4 sweep configs (TP8 +/- DPA +/- HiCache)

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
---
 .../agentic/dsv4_fp4_mi355x_sglang.sh         | 25 +++++++++++--------
 configs/amd-master.yaml                       |  7 +++---
 2 files changed, 18 insertions(+), 14 deletions(-)

diff --git a/benchmarks/single_node/agentic/dsv4_fp4_mi355x_sglang.sh b/benchmarks/single_node/agentic/dsv4_fp4_mi355x_sglang.sh
index 6dc348c66f..8aab580582 100755
--- a/benchmarks/single_node/agentic/dsv4_fp4_mi355x_sglang.sh
+++ b/benchmarks/single_node/agentic/dsv4_fp4_mi355x_sglang.sh
@@ -71,7 +71,6 @@ case "${KV_OFFLOAD_BACKEND:-}" in
         HICACHE_WRITE_POLICY="${HICACHE_WRITE_POLICY:-write_through}"
         HICACHE_IO_BACKEND="${HICACHE_IO_BACKEND:-direct}"
         HICACHE_MEM_LAYOUT="${HICACHE_MEM_LAYOUT:-page_first_direct}"
-        export SGLANG_ENABLE_UNIFIED_RADIX_TREE=1
         CACHE_ARGS=(
             --enable-hierarchical-cache
             --hicache-ratio "$HICACHE_RATIO"
@@ -100,7 +99,10 @@ fi
 
 # ---- LLM server config ----------------------------------------------------------
 
-CUDA_GRAPH_MAX_BS="$CONC"
+# AgentX concurrency counts live session trees, not individual requests.
+# Allow subagent fan-out to exceed CONC without clipping request bursts.
+MAX_RUNNING_REQUESTS=$((2 * CONC))
+CUDA_GRAPH_MAX_BS=$CONC
 [ "$CUDA_GRAPH_MAX_BS" -gt 64 ] && CUDA_GRAPH_MAX_BS=64
 
 export SGLANG_DEFAULT_THINKING=1
@@ -110,6 +112,8 @@ export SGLANG_HACK_FLASHMLA_BACKEND=unified_kv_triton
 export AITER_BF16_FP8_MOE_BOUND=0
 
 PARALLEL_ARGS=(--tensor-parallel-size "$TP")
+METRICS_ARGS=(--enable-metrics)
+MEM_FRACTION_STATIC=0.90
 CHUNKED_PREFILL_SIZE=8192
 if [ "$DP_ATTENTION" = "true" ]; then
     export SGLANG_SHARED_EXPERT_TP1=1
@@ -138,12 +142,11 @@ SGLANG_CMD=(
     --port "$SGLANG_BACKEND_PORT"
     --trust-remote-code
     "${PARALLEL_ARGS[@]}"
-    --disable-radix-cache
-    --attention-backend dsv4
+    --attention-backend compressed
     --cuda-graph-max-bs "$CUDA_GRAPH_MAX_BS"
-    --max-running-requests "$CONC"
-    --mem-fraction-static 0.90
-    --swa-full-tokens-ratio 0.15
+    --max-running-requests "$MAX_RUNNING_REQUESTS"
+    --mem-fraction-static "$MEM_FRACTION_STATIC"
+    --swa-full-tokens-ratio 0.10
     --page-size 256
     --kv-cache-dtype fp8_e4m3
     --chunked-prefill-size "$CHUNKED_PREFILL_SIZE"
@@ -152,7 +155,7 @@ SGLANG_CMD=(
     --reasoning-parser deepseek-v4
     --chat-template "$(dirname "$0")/../chat_templates/deepseek_v4_thinking.jinja"
     --watchdog-timeout 1800
-    --enable-metrics
+    "${METRICS_ARGS[@]}"
     "${CACHE_ARGS[@]}"
 )
 
@@ -202,8 +205,10 @@ if [ "$USE_SGLANG_ROUTER" = "true" ]; then
     wait_for_server_ready --port "$PORT" --server-log "$ROUTER_LOG" --server-pid "$ROUTER_PID"
 fi
 
-capture_cache_metrics
-trap capture_cache_metrics EXIT
+if [ "${#METRICS_ARGS[@]}" -gt 0 ]; then
+    capture_cache_metrics
+    trap capture_cache_metrics EXIT
+fi
 
 # ---- Run benchmark ----------------------------------------------------------
 build_replay_cmd "$RESULT_DIR"
diff --git a/configs/amd-master.yaml b/configs/amd-master.yaml
index 7cd5ede798..0bea15f1fe 100644
--- a/configs/amd-master.yaml
+++ b/configs/amd-master.yaml
@@ -2279,10 +2279,9 @@ dsv4-fp4-mi355x-sglang-agentic:
     - dram-utilization: 0.80
       search-space:
       - { tp: 8, kv-offloading: none, conc-list: [48] }
-      #- { tp: 8, kv-offloading: none, conc-list: [48] }
-      #- { tp: 8, kv-offloading: dram, kv-offload-backend: hicache, conc-list: [48] }
-      #- { tp: 8, dp-attn: true, kv-offloading: none, conc-list: [48] }
-      #- { tp: 8, dp-attn: true, kv-offloading: dram, kv-offload-backend: hicache, conc-list: [48] }
+      - { tp: 8, kv-offloading: dram, kv-offload-backend: hicache, conc-list: [48] }
+      - { tp: 8, dp-attn: true, kv-offloading: none, conc-list: [48] }
+      - { tp: 8, dp-attn: true, kv-offloading: dram, kv-offload-backend: hicache, conc-list: [48] }
 
 # MiniMax-M3 MXFP8 MI355X recipe:
 # https://github.com/vllm-project/recipes/commit/2a3728ed9892debfd767a72a58ebc90b33f186e5

From 2ecf4e9f3b7311371fc0859722f397a8d1e3911e Mon Sep 17 00:00:00 2001
From: seungrokj <seungrok.jung@amd.com>
Date: Fri, 10 Jul 2026 21:55:18 +0900
Subject: [PATCH 06/11] [AMD][MI355X] Fix cudaHostRegister OOM: lower
 hicache-ratio to 4, adjust sweep concurrency

hicache-ratio=8 tried to pin ~495 GB via cudaHostRegister, exceeding
ROCm's per-process pinnable memory limit. Lower to ratio=4 and sweep
graduated concurrency levels across configs.

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
---
 .../agentic/dsv4_fp4_mi355x_sglang.sh          | 18 +++++++++---------
 configs/amd-master.yaml                        |  8 ++++----
 2 files changed, 13 insertions(+), 13 deletions(-)

diff --git a/benchmarks/single_node/agentic/dsv4_fp4_mi355x_sglang.sh b/benchmarks/single_node/agentic/dsv4_fp4_mi355x_sglang.sh
index 8aab580582..4dda76ab45 100755
--- a/benchmarks/single_node/agentic/dsv4_fp4_mi355x_sglang.sh
+++ b/benchmarks/single_node/agentic/dsv4_fp4_mi355x_sglang.sh
@@ -57,17 +57,17 @@ case "${KV_OFFLOAD_BACKEND:-}" in
   hicache)
         # ---- Hicache config ----------------------------------------------------------
         # DSv4 allocates several physical host sub-pools for each logical host
-        # token. MI355X nodes have ~3 TB of host DRAM (similar to B200's 3.8
-        # TiB), so ratio=8 at TP≥8 provides a large useful CPU tier within the
-        # node budget. Lower TP configs use higher ratios to maintain adequate
-        # host token capacity without exceeding DRAM limits.
+        # token. ROCm's cudaHostRegister has a per-process pinnable-memory
+        # ceiling well below the node's ~2.4 TiB total DRAM. At TP8 with
+        # mem-fraction-static=0.90 each GPU's KV pool is ~62 GB, so ratio=8
+        # would try to pin ~495 GB (TP8 pure) or ~495 GB per engine (DPA),
+        # both of which blow the limit. Ratio=2 keeps pinned memory safely
+        # under the ceiling while still providing a meaningful CPU KV tier.
+
+        # RuntimeError: cudaHostRegister failed (rc=1, invalid argument) for ptr=0x6f6388400000 size=531470745600; host buffer is not pinned and device transfers may silently return stale data 
+        DEFAULT_HICACHE_RATIO=4
 
-        DEFAULT_HICACHE_RATIO=8
         HICACHE_RATIO="${HICACHE_RATIO:-$DEFAULT_HICACHE_RATIO}"
-        if [ "$HICACHE_RATIO" -gt "$DEFAULT_HICACHE_RATIO" ]; then
-            echo "Error: HICACHE_RATIO=$HICACHE_RATIO exceeds configured limit $DEFAULT_HICACHE_RATIO" >&2
-            exit 1
-        fi
         HICACHE_WRITE_POLICY="${HICACHE_WRITE_POLICY:-write_through}"
         HICACHE_IO_BACKEND="${HICACHE_IO_BACKEND:-direct}"
         HICACHE_MEM_LAYOUT="${HICACHE_MEM_LAYOUT:-page_first_direct}"
diff --git a/configs/amd-master.yaml b/configs/amd-master.yaml
index 0bea15f1fe..7b1e6a63fe 100644
--- a/configs/amd-master.yaml
+++ b/configs/amd-master.yaml
@@ -2278,10 +2278,10 @@ dsv4-fp4-mi355x-sglang-agentic:
     agentic-coding:
     - dram-utilization: 0.80
       search-space:
-      - { tp: 8, kv-offloading: none, conc-list: [48] }
-      - { tp: 8, kv-offloading: dram, kv-offload-backend: hicache, conc-list: [48] }
-      - { tp: 8, dp-attn: true, kv-offloading: none, conc-list: [48] }
-      - { tp: 8, dp-attn: true, kv-offloading: dram, kv-offload-backend: hicache, conc-list: [48] }
+      - { tp: 8, kv-offloading: none, conc-list: [16] }
+      - { tp: 8, kv-offloading: dram, kv-offload-backend: hicache, conc-list: [16, 32] }
+      - { tp: 8, dp-attn: true, kv-offloading: none, conc-list: [16, 32] }
+      - { tp: 8, dp-attn: true, kv-offloading: dram, kv-offload-backend: hicache, conc-list: [16, 32, 48] }
 
 # MiniMax-M3 MXFP8 MI355X recipe:
 # https://github.com/vllm-project/recipes/commit/2a3728ed9892debfd767a72a58ebc90b33f186e5

From fcffec7d8aba655d9a6671aa0bbf56c31b2756d4 Mon Sep 17 00:00:00 2001
From: seungrokj <seungrok.jung@amd.com>
Date: Sat, 11 Jul 2026 01:48:50 +0900
Subject: [PATCH 07/11] [AMD][MI355X] DSv4 SGLang agentic: focus sweep on DPA
 configs, scale concurrency to 96

Device KV utilization peaked at only 27% even at conc=48, and non-DPA
configs underperformed DPA in both throughput and completed requests.
Drop non-DPA sweep entries and extend DPA concurrency range to [16..96]
(none) and [48..96] (hicache) to find the saturation point.

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
---
 configs/amd-master.yaml | 6 ++----
 1 file changed, 2 insertions(+), 4 deletions(-)

diff --git a/configs/amd-master.yaml b/configs/amd-master.yaml
index df1eda40fb..c125630708 100644
--- a/configs/amd-master.yaml
+++ b/configs/amd-master.yaml
@@ -2333,10 +2333,8 @@ dsv4-fp4-mi355x-sglang-agentic:
     agentic-coding:
     - dram-utilization: 0.80
       search-space:
-      - { tp: 8, kv-offloading: none, conc-list: [16] }
-      - { tp: 8, kv-offloading: dram, kv-offload-backend: hicache, conc-list: [16, 32] }
-      - { tp: 8, dp-attn: true, kv-offloading: none, conc-list: [16, 32] }
-      - { tp: 8, dp-attn: true, kv-offloading: dram, kv-offload-backend: hicache, conc-list: [16, 32, 48] }
+      - { tp: 8, dp-attn: true, kv-offloading: none, conc-list: [16, 32, 48, 64, 80, 96] }
+      - { tp: 8, dp-attn: true, kv-offloading: dram, kv-offload-backend: hicache, conc-list: [48, 64, 80, 96] }
 
 # MiniMax-M3 MXFP8 MI355X recipe:
 # https://github.com/vllm-project/recipes/commit/2a3728ed9892debfd767a72a58ebc90b33f186e5

From 2294501c4b7ab4a4cc82f776ad19c2676feb25a3 Mon Sep 17 00:00:00 2001
From: seungrokj <seungrok.jung@amd.com>
Date: Sat, 11 Jul 2026 11:04:52 +0900
Subject: [PATCH 08/11] [AMD][MI355X] DSv4 SGLang agentic: add non-DPA baseline
 config (tp=8, conc 1-8)

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
---
 configs/amd-master.yaml | 1 +
 1 file changed, 1 insertion(+)

diff --git a/configs/amd-master.yaml b/configs/amd-master.yaml
index c125630708..d1f7e4eb49 100644
--- a/configs/amd-master.yaml
+++ b/configs/amd-master.yaml
@@ -2333,6 +2333,7 @@ dsv4-fp4-mi355x-sglang-agentic:
     agentic-coding:
     - dram-utilization: 0.80
       search-space:
+      - { tp: 8, kv-offloading: none, conc-list: [1, 2, 4, 8] }
       - { tp: 8, dp-attn: true, kv-offloading: none, conc-list: [16, 32, 48, 64, 80, 96] }
       - { tp: 8, dp-attn: true, kv-offloading: dram, kv-offload-backend: hicache, conc-list: [48, 64, 80, 96] }
 

From 0c6bcc33e7c4797644e221e6687698d3c35fd077 Mon Sep 17 00:00:00 2001
From: seungrokj <144636725+seungrokj@users.noreply.github.com>
Date: Sat, 11 Jul 2026 11:16:07 +0900
Subject: [PATCH 09/11] Update perf-changelog.yaml

---
 perf-changelog.yaml | 7 -------
 1 file changed, 7 deletions(-)

diff --git a/perf-changelog.yaml b/perf-changelog.yaml
index 9ab1125b3e..8822b0541a 100644
--- a/perf-changelog.yaml
+++ b/perf-changelog.yaml
@@ -4728,10 +4728,3 @@
     - "Add capture_cache_metrics for pre/post-benchmark cache stats"
     - "Sweep conc=48 across TP8 +/- DPA +/- HiCache"
   pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2146
-
-- config-keys:
-    - dsv4-fp4-mi355x-sglang-disagg-agentic-hicache
-  description:
-    - "Bump the SGLang ROCm image from 20260706 to 20260709 and retire the DeepSeek-V4 compress-state, DSA paged-MQA, AITER instruction-shape, and HiCache host-pool shims that are already included upstream."
-    - "Remove the SWA re-prefill and unified-KV HiCache overlays after their upstream PRs merged; retain only patch code without a merged upstream equivalent."
-  pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2127

From a91aa69a3fb907261e8be13caadf26bd673f4a7d Mon Sep 17 00:00:00 2001
From: seungrokj <seungrok.jung@amd.com>
Date: Sat, 11 Jul 2026 11:22:36 +0900
Subject: [PATCH 10/11] [AMD][MI355X] exclude mia1-p01-g12 from agentic salloc

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
---
 runners/launch_mi355x-amds.sh | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/runners/launch_mi355x-amds.sh b/runners/launch_mi355x-amds.sh
index 8cb92b7a16..ec99d1698d 100644
--- a/runners/launch_mi355x-amds.sh
+++ b/runners/launch_mi355x-amds.sh
@@ -212,7 +212,7 @@ else
     #   mia1-p01-g09: pyxis broken (persistently fails to create container filesystem)
     #   mia1-p01-g11: docker.sock permissions denied (cluster-cleanup step fails)
     # Both have been root-caused via #1431/#1432/#1440/#1441/#1443 sweep failures.
-    salloc --partition=$PARTITION --exclude=mia1-p01-g09,mia1-p01-g11 --gres=gpu:$GPU_COUNT --exclusive --cpus-per-task=128 --time=500 --no-shell --job-name="$RUNNER_NAME"
+    salloc --partition=$PARTITION --exclude=mia1-p01-g09,mia1-p01-g11,mia1-p01-g12 --gres=gpu:$GPU_COUNT --exclusive --cpus-per-task=128 --time=500 --no-shell --job-name="$RUNNER_NAME"
     JOB_ID=$(squeue --name="$RUNNER_NAME" -h -o %A | head -n1)
 
     srun --jobid=$JOB_ID bash -c "docker stop \$(docker ps -a -q)"

From a1da9ec6c5f1bdc7b025878b6bb43f051e6f3e08 Mon Sep 17 00:00:00 2001
From: seungrokj <seungrok.jung@amd.com>
Date: Sat, 11 Jul 2026 11:38:31 +0900
Subject: [PATCH 11/11] [AMD][MI355X] DSv4 SGLang: narrow DPA sweep ranges, add
 disagg agentic hicache recipe

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
---
 configs/amd-master.yaml | 37 +++++++++++++++++++++++++++++++++++--
 1 file changed, 35 insertions(+), 2 deletions(-)

diff --git a/configs/amd-master.yaml b/configs/amd-master.yaml
index 912c151caa..71d8cf303c 100644
--- a/configs/amd-master.yaml
+++ b/configs/amd-master.yaml
@@ -2334,8 +2334,8 @@ dsv4-fp4-mi355x-sglang-agentic:
     - dram-utilization: 0.80
       search-space:
       - { tp: 8, kv-offloading: none, conc-list: [1, 2, 4, 8] }
-      - { tp: 8, dp-attn: true, kv-offloading: none, conc-list: [16, 32, 48, 64, 80, 96] }
-      - { tp: 8, dp-attn: true, kv-offloading: dram, kv-offload-backend: hicache, conc-list: [48, 64, 80, 96] }
+      - { tp: 8, dp-attn: true, kv-offloading: none, conc-list: [16, 32, 48, 64] }
+      - { tp: 8, dp-attn: true, kv-offloading: dram, kv-offload-backend: hicache, conc-list: [80, 96] }
 
 # MiniMax-M3 MXFP8 MI355X recipe:
 # https://github.com/vllm-project/recipes/commit/2a3728ed9892debfd767a72a58ebc90b33f186e5
@@ -3119,3 +3119,36 @@ minimaxm3-fp8-mi325x-vllm-agentic:
       - { tp: 4, ep: 4, kv-offloading: dram, kv-offload-backend: native, conc-list: [3, 4, 5, 6, 7, 8, 9, 10, 12, 14, 16] }
       - { tp: 8, ep: 8, kv-offloading: dram, kv-offload-backend: native, conc-list: [10, 12, 14, 16, 18, 20, 22, 24, 26, 28, 32] }
       - { tp: 8, ep: 8, dp-attn: true, kv-offloading: dram, kv-offload-backend: native, conc-list: [24, 32, 36, 40, 44, 48, 52, 56, 60, 64, 72, 80, 96] }
+
+dsv4-fp4-mi355x-sglang-disagg-agentic-hicache:
+  image: lmsysorg/sglang-rocm:v0.5.14-rocm720-mi35x-20260710
+  model: deepseek-ai/DeepSeek-V4-Pro
+  model-prefix: dsv4
+  runner: cluster:mi355x-amds
+  precision: fp4
+  framework: sglang-disagg
+  multinode: true
+  disagg: true
+  scenarios:
+    agentic-coding:
+    - dram-utilization: 0.80
+      search-space:
+      - spec-decoding: "none"
+        conc-list: [ 2,4,8,16,32 ]
+        kv-offloading: dram
+        kv-offload-backend: hicache
+        prefill:
+          num-worker: 1
+          tp: 8
+          ep: 1
+          dp-attn: false
+          additional-settings:
+          - "PREFILL_NODES=1"
+        decode:
+          num-worker: 1
+          tp: 8
+          ep: 1
+          dp-attn: false
+          additional-settings:
+          - "DECODE_NODES=1"
+          - "DECODE_MTP_SIZE=0"