[AMD][MI355X] DSv4 SGLang agentic: v0.5.14 image, router, HiCache

benchmarks/single_node/agentic/dsv4_fp4_mi355x_sglang.sh

@@ -1,34 +1,31 @@
 #!/usr/bin/env bash
-set -euo pipefail
+set -eo 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"
 
 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:-unknown}"
+if [[ -n "$SLURM_JOB_ID" ]]; then
+    echo "JOB $SLURM_JOB_ID running on $SLURMD_NODENAME"
 fi
 
 # ROCR/HIP visibility under slurm cgroups.
-if [ -n "${ROCR_VISIBLE_DEVICES:-}" ]; then
+if [ -n "$ROCR_VISIBLE_DEVICES" ]; then
     export HIP_VISIBLE_DEVICES="$ROCR_VISIBLE_DEVICES"
 fi
 
-# `hf download` creates the target dir if missing and is itself idempotent.
-# When MODEL_PATH is unset (stand-alone runs), fall back to the HF_HUB_CACHE
-# Either way, MODEL_PATH is what the server is launched with.
-if [[ -n "${MODEL_PATH:-}" ]]; then
+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
@@ -43,115 +40,153 @@ 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.
+CACHE_ARGS=()
+if agentic_kv_offload_enabled; then
+    # HiCache config — https://lmsysorg.mintlify.app/cookbook/autoregressive/DeepSeek/DeepSeek-V4
+    case "$KV_OFFLOAD_BACKEND" in
+        hicache)
+            HICACHE_RATIO=4
+            HICACHE_WRITE_POLICY="write_through"
+            HICACHE_IO_BACKEND="direct"
+            HICACHE_MEM_LAYOUT="page_first_direct"
+            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: hicache)" >&2
+            exit 1
+            ;;
+    esac
+fi
+# ---- Client config ----------------------------------------------------------
+export AIPERF_HTTP_TCP_USER_TIMEOUT=1000000 
+
+# ---- LLM server config ----------------------------------------------------------
+USE_SGLANG_ROUTER=false
+SGLANG_BACKEND_PORT="$PORT"
+ROUTER_LOG="$RESULT_DIR/router.log"
+MEM_FRACTION_STATIC=0.90
+CHUNKED_PREFILL_SIZE=8192
 PARALLEL_ARGS=(--tensor-parallel-size "$TP")
 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)
+
+    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
+if [ "$EP_SIZE" -gt 1 ]; then
+    PARALLEL_ARGS+=(--ep-size "$EP_SIZE")
 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 &
+# 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 128 ] && CUDA_GRAPH_MAX_BS=128
+
+export SGLANG_DEFAULT_THINKING=1
+export SGLANG_DSV4_REASONING_EFFORT=high
+export SGLANG_USE_ROCM700A=0
+export SGLANG_HACK_FLASHMLA_BACKEND=unified_kv_triton
+export AITER_BF16_FP8_MOE_BOUND=0
+
+# sglang kv cache
+export SGLANG_ENABLE_UNIFIED_RADIX_TREE=1
+export SGLANG_OPT_UNIFIED_CACHE_FREE_OUT_OF_WINDOW_SLOTS=1
+
+METRICS_ARGS=(--enable-metrics)
+SPEC_ARGS=()
+
+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[@]}"
+    --attention-backend compressed
+    --cuda-graph-max-bs-decode "$CUDA_GRAPH_MAX_BS"
+    --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"
+    --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
+    "${METRICS_ARGS[@]}"
+    "${SPEC_ARGS[@]}"
+    "${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"
+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..."
+    "${SGLANG_ROUTER_CMD[@]}" \
+        --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
 
 # ---- 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"

configs/amd-master.yaml

@@ -2341,13 +2341,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
+dsv4-fp4-mi355x-sglang-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
@@ -2356,9 +2351,11 @@ dsv4-fp4-mi355x-sglang-agentic:
   multinode: false
   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] }
+    - 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] }
+      - { tp: 8, dp-attn: true, kv-offloading: dram, kv-offload-backend: { name: hicache }, conc-list: [64] }
 
 # MiniMax-M3 MXFP8 MI355X recipe:
 # https://github.com/vllm-project/recipes/commit/2a3728ed9892debfd767a72a58ebc90b33f186e5

perf-changelog.yaml

@@ -4782,6 +4782,18 @@
     - "Clean the export envs"
   pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2198
 
+- config-keys:
+    - dsv4-fp4-mi355x-sglang-agentic-hicache
+  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"
+  pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2146
+
 - config-keys:
     - kimik2.5-fp4-b300-vllm-agentic-mtp
   description:

runners/launch_mi355x-amds.sh

@@ -208,11 +208,7 @@ else
     export GPU_COUNT="${GPU_COUNT:-${TP:?TP must be set}}"
 
     set -x
-    # Exclude known-bad mi355x compute nodes (KLAUD_DEBUG §5.1 / §5.2):
-    #   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 --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)"