diff --git a/benchmarks/single_node/agentic/dsv4_fp4_mi355x_sglang.sh b/benchmarks/single_node/agentic/dsv4_fp4_mi355x_sglang.sh index d8cba43f0..09378a8dc 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" @@ -25,9 +25,6 @@ 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 [[ ! -d "$MODEL_PATH" || -z "$(ls -A "$MODEL_PATH" 2>/dev/null)" ]]; then hf download "$MODEL" --local-dir "$MODEL_PATH" @@ -43,115 +40,165 @@ 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. -PARALLEL_ARGS=(--tensor-parallel-size "$TP") +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 + # 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 + +# ---- LLM server config ---------------------------------------------------------- +USE_SGLANG_ROUTER=false +SGLANG_BACKEND_PORT="$PORT" +ROUTER_LOG="$RESULT_DIR/router.log" +MEM_FRACTION_STATIC=0.95 +CHUNKED_PREFILL_SIZE=16384 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 -fi +# 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 -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 & +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) +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 "$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" +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..." + "${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 + +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" diff --git a/configs/amd-master.yaml b/configs/amd-master.yaml index 4b1ea4aed..07d64ec9e 100644 --- a/configs/amd-master.yaml +++ b/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.15-rocm700-mi35x-20260713 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: [80, 96] } # MiniMax-M3 MXFP8 MI355X recipe: # https://github.com/vllm-project/recipes/commit/2a3728ed9892debfd767a72a58ebc90b33f186e5 diff --git a/perf-changelog.yaml b/perf-changelog.yaml index 57d510dd1..25ad1ee95 100644 --- a/perf-changelog.yaml +++ b/perf-changelog.yaml @@ -4750,3 +4750,15 @@ - "Image: lmsysorg/sglang:nightly-dev-cu13-20260709-074bb928" - "6 topologies across 1k/1k and 8k/1k: 1P1D TP4 STP + wide-EP (DEP4 prefill / DEP16 decode) from 1P1D up to 8P1D, recipes under benchmarks/multi_node/srt-slurm-recipes/sglang/qwen3.5/gb300-fp8/" pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2137 + +- config-keys: + - dsv4-fp4-mi355x-sglang-agentic-hicache + description: + - "Bump image to lmsysorg/sglang-rocm:v0.5.15-rocm700-mi35x-20260713" + - "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 diff --git a/runners/launch_mi355x-amds.sh b/runners/launch_mi355x-amds.sh index 8cb92b7a1..ec99d1698 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)"