diff --git a/benchmarks/multi_node/agentic/glm5.2_fp4_mi355x_sglang-disagg.sh b/benchmarks/multi_node/agentic/glm5.2_fp4_mi355x_sglang-disagg.sh new file mode 100755 index 0000000000..684a8dce7b --- /dev/null +++ b/benchmarks/multi_node/agentic/glm5.2_fp4_mi355x_sglang-disagg.sh @@ -0,0 +1,176 @@ +#!/usr/bin/env bash + +# Agentic trace-replay recipe for a disaggregated SGLang server on MI355X +# (GLM-5.2 MXFP4, 1P1D TP8, HiCache DRAM KV offload, built-in MTP/EAGLE). +# +# CI-style sibling of dsv4_fp4_mi355x_sglang-disagg.sh: driven entirely by +# environment variables and submits a SLURM job via submit.sh. HiCache defaults: +# ratio 1.5 + page_first_direct + write_through_selective (L2; avoids Mori RDMA race). + +SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" +source "$SCRIPT_DIR/../../benchmark_lib.sh" + +# MI355X amd-aim: default model root matches launch_mi355x-amds.sh and job.slurm +# (MODEL_DIR=/it-share/data on mia1*). Override MODEL_PATH/MODEL_DIR when weights +# live elsewhere (e.g. /it-share/hf_cache). +export MODEL_NAME="${MODEL_NAME:-GLM-5.2-MXFP4}" +export MODEL_PATH="${MODEL_PATH:-/it-share/data}" +export MODEL_DIR="${MODEL_DIR:-$MODEL_PATH}" + +check_env_vars \ + CONC_LIST \ + ISL \ + OSL \ + IMAGE \ + SPEC_DECODING \ + PREFILL_NUM_WORKERS \ + PREFILL_TP \ + PREFILL_EP \ + PREFILL_DP_ATTN \ + DECODE_NUM_WORKERS \ + DECODE_TP \ + DECODE_EP \ + DECODE_DP_ATTN \ + PREFILL_NODES \ + DECODE_NODES \ + RANDOM_RANGE_RATIO \ + DURATION \ + KV_OFFLOADING \ + IS_AGENTIC \ + FRAMEWORK + +if [[ -n "$SLURM_JOB_ID" ]]; then + echo "JOB $SLURM_JOB_ID running on $SLURMD_NODENAME" +fi + +set -x + +cd "$GITHUB_WORKSPACE/benchmarks/multi_node/amd_utils" || exit 1 + +export TIME_LIMIT="${TIME_LIMIT:-08:00:00}" +export MODEL_PATH="${MODEL_PATH}" +export MODEL_DIR="${MODEL_DIR:-$MODEL_PATH}" +export MODEL_NAME="${MODEL_NAME}" +export CONTAINER_IMAGE=$IMAGE +export CLIENT_IMAGE="${CLIENT_IMAGE:-$IMAGE}" + +export MODEL_PREFIX="${MODEL_PREFIX:-glm5.2}" +export PRECISION="${PRECISION:-fp4}" +export RESULT_FILENAME="${RESULT_FILENAME:-${RUNNER_NAME:-glm5.2-fp4-agentic}}" + +export DURATION="${DURATION:-1800}" +export MAX_MODEL_LEN="${MAX_MODEL_LEN:-1000000}" + +# GLM-5.2 is glm_moe_dsa (pure DSA/MLA); skip the hybrid-state mori_conn patch. +export MORI_CONN_PATCH="${MORI_CONN_PATCH:-skip}" + +export DISABLE_CUSTOM_ALL_REDUCE="${DISABLE_CUSTOM_ALL_REDUCE:-0}" + +export KV_OFFLOADING="${KV_OFFLOADING:-none}" +if [[ "$KV_OFFLOADING" != "none" ]]; then + export KV_OFFLOAD_BACKEND="${KV_OFFLOAD_BACKEND:-hicache}" +fi +if [[ "$KV_OFFLOADING" != "none" && "${KV_OFFLOAD_BACKEND:-}" == "hicache" ]]; then + export HICACHE_TIER="${HICACHE_TIER:-L2}" + export HICACHE_HOST_POOL_COUNT="${HICACHE_HOST_POOL_COUNT:-1}" + export HICACHE_PAGE_SIZE="${HICACHE_PAGE_SIZE:-1}" + export HICACHE_RATIO="${HICACHE_RATIO:-1.5}" + export HICACHE_PREFETCH_POLICY="${HICACHE_PREFETCH_POLICY:-best_effort}" + + if [[ "${HICACHE_TIER^^}" == "L3" ]]; then + export HICACHE_MEM_LAYOUT="${HICACHE_MEM_LAYOUT:-page_first}" + export HICACHE_IO_BACKEND="${HICACHE_IO_BACKEND:-direct}" + export HICACHE_WRITE_POLICY="${HICACHE_WRITE_POLICY:-write_through}" + export HICACHE_STORAGE_BACKEND="${HICACHE_STORAGE_BACKEND:-mooncake}" + else + # write_through_selective evicts only under GPU memory pressure, avoiding the + # mori RDMA race that write_back / write_through can trigger under disagg load. + export HICACHE_MEM_LAYOUT="${HICACHE_MEM_LAYOUT:-page_first_direct}" + export HICACHE_IO_BACKEND="${HICACHE_IO_BACKEND:-direct}" + export HICACHE_WRITE_POLICY="${HICACHE_WRITE_POLICY:-write_through_selective}" + export HICACHE_STORAGE_BACKEND="${HICACHE_STORAGE_BACKEND:-}" + fi + export MC_MASTER_PORT="${MC_MASTER_PORT:-58137}" + export MC_METADATA_PORT="${MC_METADATA_PORT:-8080}" + export MC_METRICS_PORT="${MC_METRICS_PORT:-19003}" + export MC_MASTER_THREADS="${MC_MASTER_THREADS:-64}" + export MC_EVICTION_HIGH_WATERMARK="${MC_EVICTION_HIGH_WATERMARK:-0.95}" + export MC_PATCH_HOSTPOOL="${MC_PATCH_HOSTPOOL:-1}" + export MC_PROTOCOL="${MC_PROTOCOL:-tcp}" + export MC_GLOBAL_SEG="${MC_GLOBAL_SEG:-64gb}" + export MC_DEVICE="${MC_DEVICE:-}" + export MC_MASTER_ADDR="${MC_MASTER_ADDR:-}" + export MC_METADATA_SERVER="${MC_METADATA_SERVER:-}" +fi + +# MoRI IO QP tuning for this recipe: amd_utils/env.sh when MODEL_NAME=GLM-5.2-MXFP4 +# and DISAGG=true (ionic MSN-safe overrides; other models/topologies unchanged). + +export PREFILL_ROUTER_POLICY="${PREFILL_ROUTER_POLICY:-cache_aware}" +export ENABLE_METRICS="${ENABLE_METRICS:-1}" + +if [[ "${SPEC_DECODING:-none}" == "mtp" || "${DECODE_MTP_SIZE:-0}" -gt 0 ]]; then + # Do not use ${DECODE_MTP_SIZE:-2}: wrappers may pre-set 0 and block :- defaulting. + if [[ "${DECODE_MTP_SIZE:-0}" -le 0 ]]; then + export DECODE_MTP_SIZE=2 + fi + # Throughput: synthetic acceptance from golden_al_distribution/glm5.2_mtp.yaml + # (thinking_on, num_speculative_tokens=2 -> AL 2.50). EVAL_ONLY runs use real MTP. + if [[ "${EVAL_ONLY:-false}" != "true" ]]; then + export SGLANG_SIMULATE_ACC_LEN="${SGLANG_SIMULATE_ACC_LEN:-2.50}" + export SGLANG_SIMULATE_ACC_METHOD="${SGLANG_SIMULATE_ACC_METHOD:-match-expected}" + export SGLANG_SIMULATE_ACC_TOKEN_MODE="${SGLANG_SIMULATE_ACC_TOKEN_MODE:-real-draft-token}" + fi +else + export DECODE_MTP_SIZE="${DECODE_MTP_SIZE:-0}" +fi + +# AIPerf reuses keep-alive sockets across agentic turns; outlast the default +# 5s SGLang keep-alive (same mitigation as glm5.2_fp4_b300_sglang.sh). +export SGLANG_TIMEOUT_KEEP_ALIVE="${SGLANG_TIMEOUT_KEEP_ALIVE:-900}" +export AIPERF_HTTP_TCP_USER_TIMEOUT="${AIPERF_HTTP_TCP_USER_TIMEOUT:-900000}" + +if [[ "${PREFILL_EP:-1}" -eq 1 ]]; then +export PREFILL_ENABLE_EP=false +else +export PREFILL_ENABLE_EP=true +fi + +if [[ "$PREFILL_DP_ATTN" == "true" ]]; then +export PREFILL_ENABLE_DP=true +else +export PREFILL_ENABLE_DP=false +fi + +if [[ "${DECODE_EP:-1}" -eq 1 ]]; then +export DECODE_ENABLE_EP=false +else +export DECODE_ENABLE_EP=true +fi + +if [[ "$DECODE_DP_ATTN" == "true" ]]; then +export DECODE_ENABLE_DP=true +else +export DECODE_ENABLE_DP=false +fi + +SUBMIT_ARGS=( + "$PREFILL_NODES" "$PREFILL_NUM_WORKERS" "$DECODE_NODES" "$DECODE_NUM_WORKERS" + "$ISL" "$OSL" "${CONC_LIST// /x}" inf + "${PREFILL_ENABLE_EP}" "${PREFILL_ENABLE_DP}" + "${DECODE_ENABLE_EP}" "${DECODE_ENABLE_DP}" + "${PREFILL_TP}" "${DECODE_TP}" + "${RANDOM_RANGE_RATIO}" +) +if [[ -n "${NODE_LIST:-}" ]]; then + SUBMIT_ARGS+=("$NODE_LIST") +fi + +JOB_ID=$(bash ./submit.sh "${SUBMIT_ARGS[@]}") + +if [[ $? -ne 0 ]]; then + echo "Failed to submit job" >&2 + exit 1 +fi + +echo "$JOB_ID" diff --git a/benchmarks/multi_node/amd_utils/env.sh b/benchmarks/multi_node/amd_utils/env.sh index 0cf3978bbf..fb70b904c1 100755 --- a/benchmarks/multi_node/amd_utils/env.sh +++ b/benchmarks/multi_node/amd_utils/env.sh @@ -295,6 +295,19 @@ else export NCCL_BLOCKING_WAIT="${NCCL_BLOCKING_WAIT:-1}" # export NCCL_DEBUG="${NCCL_DEBUG:-INFO}" + # ========================================================================= + # GLM-5.2-MXFP4 disagg (benchmarks/multi_node/agentic/glm5.2_fp4_mi355x_sglang-disagg.sh) + # Ionic MSN-safe MoRI IO QP tuning; overrides global MoRI defaults only when + # this model runs in PD/disagg mode (DISAGG=true). Other models/topologies unchanged. + # ========================================================================= + if [[ "$MODEL_NAME" == "GLM-5.2-MXFP4" && "${DISAGG:-false}" == "true" ]]; then + export MORI_IO_SQ_BACKOFF_TIMEOUT_US="${MORI_IO_SQ_BACKOFF_TIMEOUT_US:-500000}" + export MORI_IO_QP_MAX_SEND_WR="${MORI_IO_QP_MAX_SEND_WR:-8192}" + export MORI_IO_QP_MAX_CQE="${MORI_IO_QP_MAX_CQE:-16384}" + export MORI_IO_QP_MAX_SGE="${MORI_IO_QP_MAX_SGE:-2}" + export MORI_IO_TC_DISABLE="${MORI_IO_TC_DISABLE:-0}" + fi + # ========================================================================= # DeepSeek-V4-Pro PD recipe overrides # Placed at the end of the SGLang env block so it wins over the global @@ -374,4 +387,4 @@ else export GPU_MAX_HW_QUEUES=5 fi -fi \ No newline at end of file +fi diff --git a/benchmarks/multi_node/amd_utils/job.slurm b/benchmarks/multi_node/amd_utils/job.slurm index 3c8a6f16a8..35a0622337 100755 --- a/benchmarks/multi_node/amd_utils/job.slurm +++ b/benchmarks/multi_node/amd_utils/job.slurm @@ -557,6 +557,8 @@ MC_GLOBAL_SEG=${MC_GLOBAL_SEG:-} MC_DEVICE=${MC_DEVICE:-} MC_MASTER_ADDR=${MC_MASTER_ADDR:-} MC_METADATA_SERVER=${MC_METADATA_SERVER:-} +IBDEVICES=${IBDEVICES:-} +MORI_RDMA_TC=${MORI_RDMA_TC:-} EOF echo "[config] wrote HiCache/Mooncake settings -> $HICACHE_MC_CONFIG" diff --git a/benchmarks/multi_node/amd_utils/models.yaml b/benchmarks/multi_node/amd_utils/models.yaml index 2231e6fe60..3bea8777f0 100644 --- a/benchmarks/multi_node/amd_utils/models.yaml +++ b/benchmarks/multi_node/amd_utils/models.yaml @@ -264,6 +264,37 @@ GLM-5-FP8: chunked_prefill_size: 262144 cuda_graph_bs_range: "1-128" +GLM-5.2-MXFP4-AgentX: + base_flags: "--watchdog-timeout 1200 --load-balance-method round_robin --kv-cache-dtype fp8_e4m3 --attention-backend dsa --dsa-prefill-backend tilelang --dsa-decode-backend tilelang --dsa-topk-backend sgl-kernel --disaggregation-transfer-backend mori --tool-call-parser glm47 --reasoning-parser glm45 --trust-remote-code --served-model-name amd/GLM-5.2-MXFP4 --log-level info --log-level-http error --model-loader-extra-config '{\\\"enable_multithread_load\\\": true, \\\"num_threads\\\": 8}'" + # GLM-5.2 built-in MTP via the EAGLE spec path (eagle-topk 1), matching + # srt-slurm-recipes/sglang/glm5/gb200-fp4/glm5-mtp.yaml and DeepSeek-V4 DI. + # build_server_config appends --speculative-num-steps/--speculative-num-draft-tokens + # from DECODE_MTP_SIZE when spec-decoding: mtp. + mtp_flags: "--speculative-algorithm EAGLE --speculative-eagle-topk 1" + dp_flags: "--enable-dp-attention --moe-dense-tp-size 1 --enable-dp-lm-head" + ep_flags: "--moe-a2a-backend mori" + prefill: + mem_fraction_static: 0.8 + disable_radix_cache: false + dp: + max_running_requests: 1024 + chunked_prefill_size: "MORI_MAX_DISPATCH_TOKENS_PREFILL * PREFILL_TP_SIZE" + context_length: 1048576 + no_dp: + max_running_requests: 128 + chunked_prefill_size: 131072 + context_length: 1048576 + decode: + mem_fraction_static: 0.85 + prefill_round_robin_balance: true + disagg_decode_enable_radix_cache: false + dp: + max_running_requests: 1024 + cuda_graph_bs_range: "1-128" + no_dp: + max_running_requests: 128 + cuda_graph_bs_range: "1-128" + DeepSeek-R1-0528-MXFP4-Preview: base_flags: "--decode-log-interval 1000 --log-level warning --watchdog-timeout 3600 --load-balance-method round_robin --kv-cache-dtype fp8_e4m3 --attention-backend aiter --disaggregation-transfer-backend mori" mtp_flags: "--speculative-algorithm NEXTN --speculative-eagle-topk 1" diff --git a/benchmarks/multi_node/amd_utils/server_sglang.sh b/benchmarks/multi_node/amd_utils/server_sglang.sh index aaaca61ef5..3f2241b438 100755 --- a/benchmarks/multi_node/amd_utils/server_sglang.sh +++ b/benchmarks/multi_node/amd_utils/server_sglang.sh @@ -545,14 +545,14 @@ if [[ "$KV_OFFLOADING" != "none" && "$KV_OFFLOAD_BACKEND" == "hicache" ]]; then # ever sending a request. HICACHE_RATIO="${HICACHE_RATIO:-5}" HICACHE_SIZING_FLAGS="--hicache-ratio ${HICACHE_RATIO}" - # DeepSeek V4's hybrid HiCache pool rejects --hicache-size (requires - # --hicache-ratio), so the absolute per-node budget cannot be applied to it. + # DeepSeek V4 / GLM-5.2 hybrid or DSA HiCache pools reject --hicache-size + # (require --hicache-ratio). GB-based sizing also OOMs GLM-5.2 on MI355X. # See sglang _deepseek_v4_num_host_pages() (raises ValueError when # server_args.hicache_size > 0): # https://github.com/sgl-project/sglang/blob/9dd57ef8c48e2cd82292d849f01e2130c5203e67/python/sglang/srt/mem_cache/hybrid_cache/hybrid_pool_assembler.py#L262-L266 # FORCE_HICACHE_RATIO additionally lets a recipe opt into ratio-based sizing # for any other model without unsetting TOTAL_CPU_DRAM_GB (see comment above). - if [[ "${FORCE_HICACHE_RATIO:-0}" != "1" && -n "${TOTAL_CPU_DRAM_GB:-}" && "${TOTAL_CPU_DRAM_GB}" -gt 0 && "${MODEL_NAME}" != *DeepSeek-V4* ]]; then + if [[ "${FORCE_HICACHE_RATIO:-0}" != "1" && -n "${TOTAL_CPU_DRAM_GB:-}" && "${TOTAL_CPU_DRAM_GB}" -gt 0 && "${MODEL_NAME}" != *DeepSeek-V4* && "${MODEL_NAME}" != *GLM-5.2* ]]; then # TOTAL_CPU_DRAM_GB is the prefill worker's per-node budget (only prefill # offloads KV to CPU DRAM today); --hicache-size is per rank per host # pool. A prefill server may span nodes (PREFILL_TP_SIZE is its total diff --git a/benchmarks/multi_node/amd_utils/submit.sh b/benchmarks/multi_node/amd_utils/submit.sh index 6a1598d81d..e22fa722a1 100755 --- a/benchmarks/multi_node/amd_utils/submit.sh +++ b/benchmarks/multi_node/amd_utils/submit.sh @@ -151,6 +151,7 @@ export DRY_RUN="${DRY_RUN:-0}" # Eval-related env vars (threaded from workflow → runner → here → job.slurm → Docker) export RUN_EVAL="${RUN_EVAL:-false}" export EVAL_ONLY="${EVAL_ONLY:-false}" +export ROUTER_READINESS_CANARY="${ROUTER_READINESS_CANARY:-1}" export EVAL_CONC="${EVAL_CONC:-}" export FRAMEWORK="${FRAMEWORK:-}" export PRECISION="${PRECISION:-}" @@ -204,11 +205,30 @@ if [[ -n "${SLURM_REUSE_JOBID:-}" ]]; then # Resolve allocation's nodelist if not already provided. ALLOC_NODELIST="${SLURM_JOB_NODELIST:-$(squeue -h -j "$REUSE_JID" -o '%N' 2>/dev/null)}" + if [[ -z "$ALLOC_NODELIST" ]]; then + ALLOC_NODELIST="$(scontrol show job "$REUSE_JID" 2>/dev/null | awk -F= '/NodeList=/{print $2}' | awk '{print $1}')" + fi + # Expand Slurm bracket form: host-[01,57] -> host-01,host-57 + if [[ "$ALLOC_NODELIST" == *"["* ]]; then + EXPANDED="$(scontrol show hostnames "$ALLOC_NODELIST" 2>/dev/null | paste -sd, -)" + if [[ -n "$EXPANDED" ]]; then + ALLOC_NODELIST="$EXPANDED" + elif [[ -n "${NODE_LIST:-}" ]]; then + ALLOC_NODELIST="$NODE_LIST" + fi + fi if [[ -z "$ALLOC_NODELIST" ]]; then echo "Error: could not resolve nodelist for job ${REUSE_JID}" >&2 exit 1 fi - ALLOC_NNODES=$(scontrol show hostnames "$ALLOC_NODELIST" | wc -l) + if [[ "$ALLOC_NODELIST" == *","* ]]; then + ALLOC_NNODES=$(tr ',' '\n' <<< "$ALLOC_NODELIST" | grep -c .) + else + ALLOC_NNODES=$(scontrol show hostnames "$ALLOC_NODELIST" 2>/dev/null | wc -l) + if [[ "$ALLOC_NNODES" -lt 1 ]]; then + ALLOC_NNODES=1 + fi + fi if [[ "$ALLOC_NNODES" -lt "$NUM_NODES" ]]; then echo "Error: allocation ${REUSE_JID} has ${ALLOC_NNODES} nodes, need ${NUM_NODES}" >&2 exit 1 diff --git a/configs/amd-master.yaml b/configs/amd-master.yaml index cf5cdfeb6d..935048e40e 100644 --- a/configs/amd-master.yaml +++ b/configs/amd-master.yaml @@ -1568,6 +1568,51 @@ dsv4-fp4-mi355x-sglang-disagg-agentic-hicache-mtp: - "DECODE_NODES=1" - "DECODE_MTP_SIZE=3" +# GLM-5.2 MXFP4 MI355X SGLang disaggregated agentic (1P1D TP8, HiCache DRAM KV +# offload, built-in MTP). Mirrors dsv4-fp4-mi355x-sglang-disagg-agentic-hicache; +# L2 HiCache tuning (ratio 1.5, page_first_direct, write_through_selective) +# validated on amd-aim 1P1D — see glm5.2_fp4_mi355x_sglang-disagg.sh. Weights: +# amd/GLM-5.2-MXFP4 under /it-share/data (default MODEL_DIR on mia1*). +glm5.2-fp4-mi355x-sglang-disagg-agentic-hicache: + image: lmsysorg/sglang-rocm:v0.5.16-rocm720-mi35x-20260729 + model: amd/GLM-5.2-MXFP4 + model-prefix: glm5.2 + runner: cluster:mi355x-amds + precision: fp4 + framework: sglang-disagg + kv-p2p-transfer: mori + multinode: true + disagg: true + scenarios: + agentic-coding: + - dram-utilization: 0.80 + search-space: + - spec-decoding: "mtp" + conc-list: [ 2, 4, 8, 16, 32 ] + kv-offloading: dram + kv-offload-backend: { name: hicache } + prefill: + num-worker: 1 + tp: 8 + ep: 1 + dp-attn: false + additional-settings: + - "PREFILL_NODES=1" + - "CLIENT_IMAGE=lmsysorg/sglang-rocm:v0.5.16-rocm720-mi35x-20260729" + - "HICACHE_RATIO=1.5" + - "HICACHE_WRITE_POLICY=write_through_selective" + - "HICACHE_MEM_LAYOUT=page_first_direct" + - "HICACHE_IO_BACKEND=direct" + - "HICACHE_PAGE_SIZE=1" + decode: + num-worker: 1 + tp: 8 + ep: 1 + dp-attn: false + additional-settings: + - "DECODE_NODES=1" + - "DECODE_MTP_SIZE=2" + # DeepSeek-V4-Pro FP8 single-node on MI325X (gfx942) via vLLM. # EXTRAPOLATED bring-up. Same rationale as dsv4-fp8-mi300x-vllm: sglang has no # gfx942 build of the dsv4 nvfp4 MoE / TileLang-MLA kernels, so vLLM runs the diff --git a/perf-changelog.yaml b/perf-changelog.yaml index 52c74bf168..27ad392b61 100644 --- a/perf-changelog.yaml +++ b/perf-changelog.yaml @@ -6195,3 +6195,13 @@ - "Use AITER INT4 quick-reduce, ptpc_fp8 online quantization excluding embeddings, lm_head, gates, and experts, an FP8 KV cache, and three-token MTP with golden synthetic acceptance length 2.99 for benchmark runs." - "Set max-num-seqs to twice the concurrency, select CUDA graph capture sizes by concurrency, and cap max-num-batched-tokens at 16384." pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2576 + +- config-keys: + - glm5.2-fp4-mi355x-sglang-disagg-agentic-hicache + scenario-type: + - agentic-coding + description: + - "Add GLM-5.2 MXFP4 MI355X 1P1D TP8 disaggregated SGLang AgentX with DRAM HiCache KV offload and built-in MTP (spec-decoding: mtp, DECODE_MTP_SIZE=2, EAGLE topk 1)." + - "HiCache L2: ratio 1.5, write_through_selective, page_first_direct, direct IO. Throughput uses synthetic AL 2.50 from golden_al_distribution/glm5.2_mtp.yaml; EVAL_ONLY uses real draft-model MTP." + - "Image: lmsysorg/sglang-rocm:v0.5.16-rocm720-mi35x-20260729; model amd/GLM-5.2-MXFP4; conc-list [2, 4, 8, 16, 32]." + pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2679