Skip to content

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

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
seungrokj wants to merge 8 commits into
base: main
Choose a base branch
from
+169 −101
Open

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

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
8 commits
Select commit Hold shift + click to select a range
fec0cee
[AMD][MI355X] Update DSv4 SGLang agentic launcher to v0.5.14 image wi…
seungrokj Jul 10, 2026
3fe2c5e
[AMD][MI355X] Add dram-utilization 0.80 to DSv4 SGLang agentic sweep …
seungrokj Jul 10, 2026
e282bf7
Update perf-changelog with PR link for DSv4 SGLang agentic entry
seungrokj Jul 10, 2026
6acad04
[AMD][MI355X] Remove --context-length and comment out extra sweep con…
seungrokj Jul 10, 2026
7e1c332
[AMD][MI355X] DSv4 SGLang agentic: use compressed attention, 2x max-r…
seungrokj Jul 10, 2026
2ecf4e9
[AMD][MI355X] Fix cudaHostRegister OOM: lower hicache-ratio to 4, adj…
seungrokj Jul 10, 2026
f504c26
Merge branch 'main' into amd/agentx_dsv4_sgl
seungrokj Jul 10, 2026
fcffec7
[AMD][MI355X] DSv4 SGLang agentic: focus sweep on DPA configs, scale …
seungrokj Jul 10, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
244 changes: 152 additions & 92 deletions benchmarks/single_node/agentic/dsv4_fp4_mi355x_sglang.sh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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"

Expand Down Expand Up @@ -43,115 +43,175 @@ 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. 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

HICACHE_RATIO="${HICACHE_RATIO:-$DEFAULT_HICACHE_RATIO}"
HICACHE_WRITE_POLICY="${HICACHE_WRITE_POLICY:-write_through}"
HICACHE_IO_BACKEND="${HICACHE_IO_BACKEND:-direct}"
HICACHE_MEM_LAYOUT="${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: native, mooncake, lmcache)" >&2
exit 1
Comment on lines +83 to +85

@claude claude Bot Jul 10, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 The default arm of the KV_OFFLOAD_BACKEND case at line 85 prints (expected: native, mooncake, lmcache), but the case above only accepts hicache — the listed values look copy-pasted from the sibling vLLM launchers (which do accept native/mooncake/lmcache) and would themselves fall into the same default arm here. Sweep configs only use hicache so CI never trips this branch, but anyone debugging by hand will be misdirected. Fix: change the list to (expected: hicache).

Extended reasoning...

What the bug is

The case "${KV_OFFLOAD_BACKEND:-}" block at lines 56-88 of benchmarks/single_node/agentic/dsv4_fp4_mi355x_sglang.sh has exactly one accepting arm: hicache). Anything else falls into the default *) arm at lines 84-87, which emits:

Error: unsupported KV_OFFLOAD_BACKEND '<value>' (expected: native, mooncake, lmcache)

None of native, mooncake, or lmcache are recognized by this case statement — they are all backends supported by other launchers in the repo (native for kimik2.5-fp4-mi355x-vllm-agentic, mooncake for minimaxm3-fp8-mi300x-vllm-agentic, lmcache for the sibling dsv4-fp4-mi355x-vllm-agentic). The message is a copy-paste from those vLLM-family launchers.

Why the message is self-contradictory

A user who sees unsupported KV_OFFLOAD_BACKEND 'foo' and reads the parenthetical hint will pick one of native, mooncake, or lmcache — and hit the exact same default arm and the exact same message. In the case where they picked native, the output literally reads unsupported KV_OFFLOAD_BACKEND 'native' (expected: native, mooncake, lmcache), which is self-contradictory.

Step-by-step proof

  1. Operator invokes the launcher manually with KV_OFFLOADING=dram KV_OFFLOAD_BACKEND=native (a natural attempt after seeing another dsv4 launcher use lmcache).
  2. agentic_kv_offload_enabled returns true because KV_OFFLOADING=dram.
  3. Bash enters the case at line 56 and evaluates the arms in order. The only arm is hicache).
  4. native does not match hicache), so control falls to *) at line 84.
  5. Line 85 prints: Error: unsupported KV_OFFLOAD_BACKEND 'native' (expected: native, mooncake, lmcache).
  6. Line 86 exits 1.
  7. Operator, taking the message at face value, retries with mooncake or lmcache. Same path — steps 3-6 repeat.

Why existing code doesn't prevent it

The header comment at line 7 states "KV_OFFLOADING=dram requires KV_OFFLOAD_BACKEND=hicache", so the canonical documented value is hicache. The error text just wasn't updated to match when this launcher was adapted from the vLLM sibling.

Impact

Cosmetic misdirection in a diagnostic-only path. The four sweep entries in configs/amd-master.yaml for this recipe all pass kv-offload-backend: hicache, so CI never hits the default arm. The cost is a minute or two of wasted debug time for someone experimenting with a non-hicache backend by hand.

Fix

Change line 85's parenthetical to (expected: hicache):

echo "Error: unsupported KV_OFFLOAD_BACKEND '${KV_OFFLOAD_BACKEND:-}' (expected: hicache)" >&2

On the refutations

Both refutations are duplicate-management: bug_002 and bug_004 target the same defect at the same lines with the same fix, and the refuters explicitly note the finding itself is real — they only refuted one of the two to avoid posting the same comment twice. The synthesis agent merged them into a single report, which resolves that concern.

;;
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 ----------------------------------------------------------

# 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
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")
METRICS_ARGS=(--enable-metrics)
MEM_FRACTION_STATIC=0.90
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[@]}"
--attention-backend compressed
--cuda-graph-max-bs "$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[@]}"
"${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

if [ "${#METRICS_ARGS[@]}" -gt 0 ]; then
capture_cache_metrics
trap capture_cache_metrics EXIT
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"
14 changes: 5 additions & 9 deletions configs/amd-master.yaml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2321,13 +2321,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
Expand All @@ -2336,9 +2331,10 @@ 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, 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
Expand Down
12 changes: 12 additions & 0 deletions perf-changelog.yaml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4717,6 +4717,18 @@
- "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"
pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2146

- config-keys:
- dsv4-fp4-mi355x-sglang-disagg-agentic-hicache
description:
Expand Down