[AMD][AgentX] GLM-5.2 FP4 MI355X sglang-disagg HiCache MTP / [AMD][AgentX] MI355X 上 GLM-5.2 FP4 分离式 HiCache MTP

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"

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
+fi

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"
 

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"

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

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

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

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