diff --git a/benchmarks/single_node/agentic/qwen3.8next_fp8_mi355x_sglang_mtp.sh b/benchmarks/single_node/agentic/qwen3.8next_fp8_mi355x_sglang_mtp.sh new file mode 100755 index 0000000000..c3130cf2e5 --- /dev/null +++ b/benchmarks/single_node/agentic/qwen3.8next_fp8_mi355x_sglang_mtp.sh @@ -0,0 +1,211 @@ +#!/usr/bin/env bash +set -euo pipefail +set -x + +# Agentic trace replay benchmark for Qwen3.8-Flash-Next FP8 on MI355X using +# SGLang with MTP speculative decoding. Day-zero recipe for the hybrid +# GDN + QSA architecture on CDNA4 (gfx950). +# +# MI355X is ROCm, so this arm is FP8 (Qwen/Qwen3.8-Flash-Next-FP8). +# NVFP4 is not available on AMD GPUs. Attention backend is aiter (ROCm). +# +# Image: lmsysorg/sglang-rocm:v0.5.19-rocm700-mi35x-20260909 includes +# SGLang PR #37500 (merged 2026-09-08), which adds full Qwen3.8-Flash-Next +# model support — no runtime patches needed. +# +# The SGLang cookbook (docs.sglang.io) verified flags for MI355X are: +# --tp-size 8 --attention-backend aiter --page-size 32 --kv-cache-dtype auto +# --chunked-prefill-size 16384 --watchdog-timeout 1200 --mem-fraction-static 0.9 +# +# Decode CUDA graphs are disabled due to a tilelang MFMA warp partition bug +# (N must be divisible by 16, but got 8) in the v0.5.19-rocm700-mi35x-20260909 +# image. This will be re-enabled once the upstream fix lands. +# +# Speculative decoding uses NEXTN (the Qwen3.8 Flash Next native MTP head), +# 3 steps, eagle-topk 1, 4 draft tokens — mirroring the H200 agentic recipe. +# +# Required env vars: +# MODEL, TP, CONC, EP_SIZE, KV_OFFLOADING, +# TOTAL_CPU_DRAM_GB, RESULT_DIR, DURATION + +source "$(dirname "$0")/../../benchmark_lib.sh" + +export EVAL_FRAMEWORK="lm-eval" + +check_env_vars \ + MODEL TP CONC EP_SIZE KV_OFFLOADING \ + TOTAL_CPU_DRAM_GB RESULT_DIR DURATION + +SCHEDULER_RECV_INTERVAL=${SCHEDULER_RECV_INTERVAL:-30} + +if [[ -n "${SLURM_JOB_ID:-}" ]]; then + echo "JOB $SLURM_JOB_ID running on ${SLURMD_NODENAME:-unknown}" +fi + +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 +else + hf download "$MODEL" + export MODEL_PATH="$MODEL" +fi + +rocm-smi || true +amd-smi || true + +export WEKA_LOADER_OVERRIDE=semianalysis_cc_traces_weka_with_subagents_256k +resolve_trace_source +install_agentic_deps + +export AIPERF_SERVER_METRICS_URLS="http://localhost:${PORT}/metrics" +export AIPERF_REQUIRED_SERVER_METRIC_PREFIX="sglang:" + +SERVER_LOG="$RESULT_DIR/server.log" +mkdir -p "$RESULT_DIR" + +SERVER_PID="" +cleanup_agentic_services() { + local exit_code=$? + trap - EXIT INT TERM + set +e + stop_background_process_tree "$SERVER_PID" "SGLang server" 60 + exit "$exit_code" +} +trap cleanup_agentic_services EXIT +trap 'exit 130' INT +trap 'exit 143' TERM + +# ---- HiCache (host-DRAM KV offload) ---------------------------------------- +CACHE_ARGS=() +if require_agentic_kv_offload_backend hicache; then + REQUESTED_HICACHE_TOTAL_GB="${HICACHE_TOTAL_CPU_DRAM_GB:-$TOTAL_CPU_DRAM_GB}" + if [ "$REQUESTED_HICACHE_TOTAL_GB" -gt "$TOTAL_CPU_DRAM_GB" ]; then + echo "Error: requested HiCache pool ${REQUESTED_HICACHE_TOTAL_GB} GB exceeds configured capacity ${TOTAL_CPU_DRAM_GB} GB" >&2 + exit 1 + fi + TOTAL_CPU_DRAM_GB="$REQUESTED_HICACHE_TOTAL_GB" + HICACHE_HOST_POOL_COUNT="${HICACHE_HOST_POOL_COUNT:-2}" + HICACHE_WRITE_POLICY="${HICACHE_WRITE_POLICY:-write_through_selective}" + MAX_HICACHE_SIZE_GB=$((TOTAL_CPU_DRAM_GB / TP / HICACHE_HOST_POOL_COUNT)) + HICACHE_SIZE_GB="${HICACHE_SIZE_GB:-$MAX_HICACHE_SIZE_GB}" + if [ "$HICACHE_SIZE_GB" -gt "$MAX_HICACHE_SIZE_GB" ]; then + echo "Error: HICACHE_SIZE_GB=$HICACHE_SIZE_GB exceeds configured per-pool limit $MAX_HICACHE_SIZE_GB" >&2 + exit 1 + fi + if [ "$HICACHE_SIZE_GB" -lt 1 ]; then + echo "Error: computed HICACHE_SIZE_GB=$HICACHE_SIZE_GB from TOTAL_CPU_DRAM_GB=$TOTAL_CPU_DRAM_GB, TP=$TP, HICACHE_HOST_POOL_COUNT=$HICACHE_HOST_POOL_COUNT" >&2 + exit 1 + fi + echo "HiCache CPU pool: ${HICACHE_SIZE_GB} GB per rank per host pool across TP=${TP}, host_pool_count=${HICACHE_HOST_POOL_COUNT}" + CACHE_ARGS=( + --enable-hierarchical-cache + --hicache-size "$HICACHE_SIZE_GB" + --hicache-io-backend kernel + --hicache-mem-layout page_first + --hicache-write-policy "$HICACHE_WRITE_POLICY" + ) +fi + +# ---- Parallelism ------------------------------------------------------------ +PARALLEL_ARGS=( + --tp-size "$TP" + --dp-size 1 + --ep-size "$EP_SIZE" +) + +TOKENIZER_ARGS=() +if [ "$TP" -ge 4 ]; then + TOKENIZER_ARGS=(--tokenizer-worker-num 6) +fi + +MAX_RUNNING_REQUESTS=$((2 * CONC)) +CUDA_GRAPH_MAX_BS="$CONC" +[ "$CUDA_GRAPH_MAX_BS" -gt 128 ] && CUDA_GRAPH_MAX_BS=128 + +# ---- ROCm / aiter environment ----------------------------------------------- +export PYTHONNOUSERSITE=1 +export SGLANG_USE_AITER=1 +export SGLANG_USE_AITER_UNIFIED_ATTN=1 +export AITER_FLYDSL_FORCE=1 +export ROCM_QUICK_REDUCE_QUANTIZATION=INT8 +export SGLANG_TIMEOUT_KEEP_ALIVE=1800 +export SGLANG_ENABLE_SPEC_V2=1 + +# ---- Speculative decoding (NEXTN / MTP) ------------------------------------- +SPEC_ARGS=( + --speculative-algorithm NEXTN + --speculative-num-steps 3 + --speculative-eagle-topk 1 + --speculative-num-draft-tokens 4 +) + +# golden_al_distribution/qwen3.8next_mtp.yaml: +# qwen3.8-flash-next-fp8.thinking_on[3] = 2.32. +# AgentX replays run with thinking on, so the thinking_on row is the right one. +if [ "${EVAL_ONLY:-false}" != "true" ]; then + export SGLANG_SIMULATE_ACC_LEN=2.32 + export SGLANG_SIMULATE_ACC_METHOD=match-expected + export SGLANG_SIMULATE_ACC_TOKEN_MODE=real-draft-token +fi + +# ---- multi_tokenizer cached_tokens_details patch ---------------------------- +SGLANG_MULTI_TOKENIZER=/sgl-workspace/sglang/python/sglang/srt/managers/multi_tokenizer_mixin.py +if [ -f "$SGLANG_MULTI_TOKENIZER" ]; then + if ! sed -n '/elif isinstance(output, BatchStrOutput):/,/input_token_logprobs_val=_extract_field_by_index/p' "$SGLANG_MULTI_TOKENIZER" \ + | grep -q 'cached_tokens_details=_extract_field_by_index'; then + sed -i '/elif isinstance(output, BatchStrOutput):/,/cached_tokens=_extract_field_by_index(output, "cached_tokens", i),/ { + /cached_tokens=_extract_field_by_index(output, "cached_tokens", i),/a\ + cached_tokens_details=_extract_field_by_index(\ + output, "cached_tokens_details", i\ + ), + }' "$SGLANG_MULTI_TOKENIZER" + fi +fi + +# ---- Launch server ----------------------------------------------------------- +SGLANG_CMD=( + python3 -m sglang.launch_server + --model-path "$MODEL_PATH" + --served-model-name "$MODEL" + --host 0.0.0.0 + --port "$PORT" + --trust-remote-code + "${PARALLEL_ARGS[@]}" + --attention-backend aiter + --page-size 32 + --kv-cache-dtype auto + --mem-fraction-static 0.90 + --model-loader-extra-config '{"enable_multithread_load": true}' + --watchdog-timeout 1200 + --chunked-prefill-size 16384 + --mamba-ssm-dtype bfloat16 + --cuda-graph-backend-decode disabled + --cuda-graph-max-bs "$CUDA_GRAPH_MAX_BS" + --max-running-requests "$MAX_RUNNING_REQUESTS" + --scheduler-recv-interval "$SCHEDULER_RECV_INTERVAL" + --stream-interval 50 + "${TOKENIZER_ARGS[@]}" + --tokenizer-path "$MODEL" + --reasoning-parser qwen3 + --tool-call-parser qwen3_coder + "${SPEC_ARGS[@]}" + --enable-metrics + --enable-cache-report + "${CACHE_ARGS[@]}" +) + +printf '%q ' "${SGLANG_CMD[@]}" | tee "$RESULT_DIR/sglang_command.txt" +printf '\n' | tee -a "$RESULT_DIR/sglang_command.txt" +"${SGLANG_CMD[@]}" > "$SERVER_LOG" 2>&1 & +SERVER_PID=$! + +wait_for_server_ready --port "$PORT" --server-log "$SERVER_LOG" --server-pid "$SERVER_PID" + +if [ "${EVAL_ONLY:-false}" = "true" ]; then + run_eval --port "$PORT" +else + build_replay_cmd "$RESULT_DIR" + REPLAY_CMD+=" --apply-chat-template" + run_agentic_replay_and_write_outputs "$RESULT_DIR" +fi diff --git a/configs/amd-master.yaml b/configs/amd-master.yaml index 2510dc3b8d..2d19d217bb 100644 --- a/configs/amd-master.yaml +++ b/configs/amd-master.yaml @@ -443,6 +443,20 @@ qwen3.5-fp8-mi300x-sglang-agentic-mtp: - { tp: 8, ep: 1, spec-decoding: mtp, kv-offloading: none, conc-list: [4, 16, 20, 24, 32] } - { tp: 8, ep: 8, spec-decoding: mtp, kv-offloading: none, conc-list: [4, 16, 20, 24, 32] } +qwen3.8-flash-next-fp8-mi355x-sglang-agentic-mtp: + image: lmsysorg/sglang-rocm:v0.5.19-rocm700-mi35x-20260909 + model: Qwen/Qwen3.8-Flash-Next-FP8 + model-prefix: qwen3.8-flash-next + runner: cluster:mi355x-amds + precision: fp8 + framework: sglang + multinode: false + scenarios: + agentic-coding: + - dram-utilization: 0.80 + search-space: + - { tp: 4, ep: 1, spec-decoding: mtp, kv-offloading: none, conc-list: [1, 4, 8, 12, 16, 20, 24, 28, 32] } + dsr1-fp8-mi355x-atom: image: rocm/atom:rocm7.2.3_ubuntu24.04_py3.12_pytorch_release_2.10.0_atom20260511 model: deepseek-ai/DeepSeek-R1-0528 diff --git a/perf-changelog.yaml b/perf-changelog.yaml index e0ef448195..3455390cf3 100644 --- a/perf-changelog.yaml +++ b/perf-changelog.yaml @@ -6947,3 +6947,12 @@ - "Update vLLM ROCm image from vllm/vllm-openai-rocm:v0.27.1 (v0.27.1 release) to vllm/vllm-openai-rocm:nightly-d9105ea8001e0a6d77a96327d17515bb5791fb36 (2026-09-07 upstream ROCm nightly, digest sha256:74d4a95f3ae672ecddf9acb7296917def82d9eca51687fa2862ae72b03ff1907, tag commit vllm-project/vllm@d9105ea8; Docker Hub last pushed 2026-09-07T05:26:48Z). benchmarks/single_node/agentic/minimaxm3_fp8_mi300x_mtp.sh is unchanged: TRITON_ATTN attention, fp8 KV, block-size 128, EAGLE3 speculative decoding with the Inferact MiniMax-M3 EAGLE3-GQA draft and the committed golden synthetic acceptance length, minimax_m3 tool-call and reasoning parsers; the search space is unchanged. The upstream commit-pinned ROCm nightly is the same vllm commit the B200 MiniMax-M3 AgentX recipe moved to in #2860. Note that vllm-openai-rocm commit-nightly tags have expired from Docker Hub within days in the past; node squash caches keep merged configs running, but a re-pin to a durable tag may be needed later." - "Add --compilation-config cudagraph_mode=FULL_DECODE_ONLY to the serve command. The upstream nightly does not torch-compile MiniMaxM3SparseForConditionalGeneration, so with VLLM_USE_BREAKABLE_CUDAGRAPH=0 the default FULL_AND_PIECEWISE mode aborts at engine init with piecewise CUDA graphs unavailable (first sweep, run 34174124043, eval cell); full decode-only graphs are what the MI355X MiniMax-M3 sibling runs on its nightly (#2825)." pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2873 + +- config-keys: + - qwen3.8-flash-next-fp8-mi355x-sglang-agentic-mtp + scenario-type: + - agentic-coding + description: + - "Add day-zero Qwen3.8-Flash-Next FP8 agentic benchmark on MI355X with NEXTN MTP speculative decoding. New config block qwen3.8-flash-next-fp8-mi355x-sglang-agentic-mtp using lmsysorg/sglang-rocm:v0.5.19-rocm700-mi35x-20260909 image (includes SGLang PR #37500, merged 2026-09-08, full Qwen3.8-Flash-Next model support) with model Qwen/Qwen3.8-Flash-Next-FP8 (176B total, 6B active, hybrid GDN+QSA architecture). TP=4, EP=1, NEXTN MTP (3 steps, eagle-topk 1, 4 draft tokens), aiter attention backend, page-size 32, kv-cache-dtype auto, mem-fraction-static 0.90, mamba-ssm-dtype bfloat16, golden acceptance length SGLANG_SIMULATE_ACC_LEN=2.32. Concurrency grid [1, 4, 8, 12, 16, 20, 24, 28, 32], no kv-offloading. Matches SGLang cookbook MI355X recipe (docs.sglang.io)." + - "Decode CUDA graphs are disabled (--cuda-graph-backend-decode disabled) due to a tilelang MFMA warp partition bug (N must be divisible by 16, but got 8) in the v0.5.19-rocm700-mi35x-20260909 image during decode graph capture. This affects all TP values. Will be re-enabled once the upstream fix lands." + pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2924