diff --git a/benchmarks/single_node/agentic/glm5.2_fp4_mi355x_sglang.sh b/benchmarks/single_node/agentic/glm5.2_fp4_mi355x_sglang.sh new file mode 100755 index 0000000000..3e020bfca4 --- /dev/null +++ b/benchmarks/single_node/agentic/glm5.2_fp4_mi355x_sglang.sh @@ -0,0 +1,275 @@ +#!/usr/bin/env bash +set -eo pipefail +set -x + +# Agentic trace replay benchmark for GLM-5.2 MXFP4 (amd/GLM-5.2-MXFP4) on +# MI355X using SGLang. +# +# Motivation over the FP8 sibling (glm5.2_fp8_mi355x_sglang.sh): the FP8 +# frontier is prefill-starved (TTFT p50 ~390 s at conc 16) because the ROCm +# DSA path defaults to bf16 KV and TP-mode DSA replicates the KV pool on +# every rank - 1.335M tokens aggregate against a 131k-token mean corpus. +# MXFP4 halves the weight bytes (438 GB vs 756 GB), growing the device pool +# to 1.96M tokens/rank (182.70 GB), and the DP-attention arm un-replicates +# it: each DP rank holds only its own sessions, so aggregate capacity is +# ~15.7M tokens at dp8 - enough to hold the conc-64 working set. +# +# Both arms run HiCache with the Mooncake store as the L3 storage tier +# (wait_complete prefetch), with per-arm L2 ratios (see the CACHE_ARGS +# block). Ratio 2 from the cookbook OOMs the 3 TB hosts outright: the main +# host pool alone would pin 2 x 182.7 GB x 8 = 2.9 TB, and GLM-5.2's DSA +# hicache additionally allocates a per-rank indexer-K sidecar host pool +# (40.33 GB/rank at ratio 2; observed "Requesting 40.33 GB but only have +# 0.87 GB free"). +# +# Required env vars: +# MODEL, TP, CONC, KV_OFFLOADING, TOTAL_CPU_DRAM_GB, RESULT_DIR, DURATION, +# EP_SIZE, DP_ATTENTION +# +# KV_OFFLOADING=dram requires KV_OFFLOAD_BACKEND=mooncake. + +source "$(dirname "$0")/../../benchmark_lib.sh" + +check_env_vars MODEL TP CONC KV_OFFLOADING TOTAL_CPU_DRAM_GB RESULT_DIR DURATION EP_SIZE DP_ATTENTION + +if [[ -n "$SLURM_JOB_ID" ]]; then + echo "JOB $SLURM_JOB_ID running on $SLURMD_NODENAME" +fi + +# ROCR/HIP visibility under slurm cgroups. +if [ -n "$ROCR_VISIBLE_DEVICES" ]; then + export HIP_VISIBLE_DEVICES="$ROCR_VISIBLE_DEVICES" +fi + +# Weights are pre-staged in the NFS HF hub cache (launch_mi355x-amds.sh mounts +# /it-share/hf-hub-cache for GLM-5.2 models); a warm cache makes this a no-op. +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 + + +# A server killed on this node minutes earlier (previous job, crashed run) +# can still be draining its ~1.4 TB of HBM: KFD reclaim takes minutes, and +# booting into a half-drained node fails RCCL init with HIP 'unhandled cuda +# error' / 'invalid argument' (observed as the mooncake-c64 CI failure). +# Wait for the GPUs to come back before launching. +# Per-GPU threshold: idle nodes hold a small driver/firmware VRAM baseline +# (observed up to ~4%/GPU, node-dependent), while a draining or occupied +# GPU sits at 50-90%. Require every GPU <= 10%. +GPU_CLEAN=false +for i in $(seq 1 90); do + VRAM_MAX=$(rocm-smi --showmemuse 2>/dev/null | grep -oE "GPU Memory Allocated \(VRAM%\): [0-9]+" | awk '{if ($NF > m) m = $NF} END {print m+0}') + if [ "${VRAM_MAX:-0}" -le 10 ]; then echo "GPUs clean (vram%max=$VRAM_MAX after $((i*10))s)"; GPU_CLEAN=true; break; fi + echo "waiting for prior-job GPU memory reclaim: vram%max=$VRAM_MAX"; sleep 10 +done +[ "$GPU_CLEAN" = "true" ] || { echo "Error: GPUs still draining prior job's memory after 15min" >&2; exit 1; } + +resolve_trace_source +install_agentic_deps + +SERVER_LOG="$RESULT_DIR/server.log" +ROUTER_LOG="$RESULT_DIR/router.log" +mkdir -p "$RESULT_DIR" + +export PYTHONNOUSERSITE=1 +# Agentic warmup dispatches hundreds of large prompts at once; allow up to +# 15 minutes of TCP progress before AIPerf declares a connection dead. +export AIPERF_HTTP_TCP_USER_TIMEOUT=900000 +# AIPerf pins one pooled keep-alive connection per session (client-side +# keep-alive 300s) while uvicorn's default SGLANG_TIMEOUT_KEEP_ALIVE is 5s; +# inter-turn idle gaps can reuse a socket exactly as the server closes it. +# Outlast the client pool so the race cannot occur. +export SGLANG_TIMEOUT_KEEP_ALIVE=900 +# The DSA indexer's top-k v2 kernel (default since v0.5.14) is JIT-compiled +# from CUDA-only source (cooperative_groups.h) and cannot build for gfx950; +# v1 dispatches to the precompiled HIP op in sgl-kernel (upstream MI355X CI +# runs DSA models the same way). +export SGLANG_OPT_USE_TOPK_V2=false + +# HiCache L2 + Mooncake L3 on every point (sizing rationale in the header). +# Per-arm L2 ratio, both measured on-node. TP arm (182.7 GB/rank device +# pool): the working set oversubscribes the device pool ~3x at conc 32, so +# the host tier is what carries the radix hits - ratio 1.5 (~2.9 TB pinned +# incl. sidecars) validates through the conc-24 long-context storm. The +# DP-attention arm (159.4 GB/rank) only runs at conc >= 32, where each DP +# rank's ~8 sessions nearly fit in its own device pool (~1.5-1.6M of 1.7M +# tokens at conc 64) and the host tier just absorbs overflow - ratio 1.5 +# boots but the host OOM killer takes the server mid-storm at conc 48, so +# it runs ratio 0.5 (~1.2 TB pinned, ~1.8 TB of load headroom) at +# negligible hit-rate cost. +CACHE_ARGS=() +if require_agentic_kv_offload_backend mooncake; then + if [ "$DP_ATTENTION" = "true" ]; then + HICACHE_RATIO="${HICACHE_RATIO:-0.5}" + else + HICACHE_RATIO="${HICACHE_RATIO:-1.5}" + fi + HICACHE_WRITE_POLICY="${HICACHE_WRITE_POLICY:-write_through}" + HICACHE_IO_BACKEND="${HICACHE_IO_BACKEND:-direct}" + HICACHE_MEM_LAYOUT="${HICACHE_MEM_LAYOUT:-page_first_direct}" + L3_PER_RANK_GB="${L3_PER_RANK_GB:-40}" + python3 -c "from mooncake.store import MooncakeDistributedStore" >/dev/null + MOONCAKE_MASTER_PORT=$((PORT + 12000)) + MOONCAKE_MASTER_LOG="$RESULT_DIR/mooncake_master.log" + MOONCAKE_CONFIG_PATH="$RESULT_DIR/mooncake_config.json" + cat > "$MOONCAKE_CONFIG_PATH" < "$MOONCAKE_MASTER_LOG" 2>&1 & + MOONCAKE_MASTER_PID=$! + sleep 2 + kill -0 "$MOONCAKE_MASTER_PID" + echo "HiCache+Mooncake: ratio=$HICACHE_RATIO, l3_per_rank=${L3_PER_RANK_GB} GB, dram_budget=${TOTAL_CPU_DRAM_GB} GB" + CACHE_ARGS=( + --enable-hierarchical-cache + --hicache-ratio "$HICACHE_RATIO" + --hicache-size 0 + --hicache-write-policy "$HICACHE_WRITE_POLICY" + --hicache-io-backend "$HICACHE_IO_BACKEND" + --hicache-mem-layout "$HICACHE_MEM_LAYOUT" + --hicache-storage-backend mooncake + --hicache-storage-prefetch-policy wait_complete + ) +fi + +# Arm selection. TP arm keeps the FP8 sibling's cookbook batch-shaping +# bands. +# +# NOTE: the DP-attention path below is currently DORMANT (no dp-attn arms +# in amd-master.yaml): DSA + dp-attention hangs a collective under +# long-context prefill on ROCm v0.5.14 (watchdog kills the scheduler with +# zero completions; reproduced with and without HiCache, with and without +# the DSv4 DP collective envs; short prompts are fine). Re-enable the +# config arm once upstream fixes the DSA DP prefill path. +# +# When active, the DP-attention (DEP) arm fronts the DP ranks with sglang-router +# using consistent hashing on the AIPerf correlation id so multi-turn +# sessions stay on the DP rank holding their radix/hicache prefix, and +# widens chunked-prefill (whole-engine, /dp ranks) like the B300 sibling. +USE_SGLANG_ROUTER=false +SGLANG_BACKEND_PORT="$PORT" +PARALLEL_ARGS=(--tp "$TP" --ep-size "$EP_SIZE") +MEM_FRACTION_STATIC=0.85 +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) + PARALLEL_ARGS+=(--dp "$TP" --enable-dp-attention) + CHUNKED_PREFILL_SIZE=32768 + export AGENTIC_WARMUP_GRACE_PERIOD=3600 + # Swap the DP gather collectives to gatherv/reduce-scatter on ROCm + # (dsv4_fp4_mi355x_sglang.sh precedent - the only green DP-attention + # config on this cluster/image): with the defaults the DSA DP path + # hangs a collective under long-context prefill load until the + # watchdog kills the scheduler (0/96 storm completions, twice). + export SGLANG_DP_USE_GATHERV=1 + export SGLANG_DP_USE_REDUCE_SCATTER=1 + export GPU_MAX_HW_QUEUES=5 +elif [ "$CONC" -le 16 ]; then + # A full 131072-token prefill chunk needs ~7 GiB/rank of activation + # headroom on top of the static pool; pair it with mem-fraction 0.80 + # like the FP8 sibling's low-conc band (0.85 OOMs the device mid-replay: + # "Tried to allocate 6.86 GiB ... 5.15 GiB is free", run 29751563205). + CHUNKED_PREFILL_SIZE=131072 + MEM_FRACTION_STATIC=0.80 +else + CHUNKED_PREFILL_SIZE=32768 + export AGENTIC_WARMUP_GRACE_PERIOD=3600 +fi +MAX_RUNNING_REQUESTS=$((2 * CONC)) +[ "$MAX_RUNNING_REQUESTS" -gt 256 ] && MAX_RUNNING_REQUESTS=256 +CUDA_GRAPH_MAX_BS=$MAX_RUNNING_REQUESTS + +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[@]}" + --dsa-prefill-backend tilelang + --dsa-decode-backend tilelang + # GLM-5.2 emits the GLM-4.7-style tool-call format; glm47 is required for + # structured message.tool_calls (SWE-bench agentic evals die without it). + # The glm45 reasoning parser keeps hybrid thinking in reasoning_content. + --tool-call-parser glm47 + --reasoning-parser glm45 + --chunked-prefill-size "$CHUNKED_PREFILL_SIZE" + --mem-fraction-static "$MEM_FRACTION_STATIC" + --max-running-requests "$MAX_RUNNING_REQUESTS" + --cuda-graph-max-bs "$CUDA_GRAPH_MAX_BS" + "${CACHE_ARGS[@]}" + --watchdog-timeout 1800 + --enable-metrics +) + +printf '%q ' "${SGLANG_CMD[@]}" | tee "$RESULT_DIR/sglang_command.txt" +printf '\n' | tee -a "$RESULT_DIR/sglang_command.txt" + +echo "Starting SGLang server for MI355X..." +"${SGLANG_CMD[@]}" > "$SERVER_LOG" 2>&1 & +SERVER_PID=$! +echo "Server PID: $SERVER_PID" + +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 [ "${EVAL_ONLY}" = "true" ]; then + # GLM-5.2's chat template defaults to reasoning_effort=Max when the + # client passes no chat_template_kwargs (mini-swe-agent doesn't), and the + # heavy thinking burns the default 75-step budget before submission. + # Double the step budget for this recipe; others keep the shared default. + export SWEBENCH_AGENT_STEP_LIMIT=150 + # Pin eval agent parallelism to the proven-green level: workers default + # to CONC, and at 64 concurrent Modal sandboxes the cluster's egress + # collapses (18k "Cannot connect to *.modal.host" errors crippled the + # trajectories in run 29764760177) while 32 ran clean. The serving + # config is unchanged - only the agent's session fan-out is capped. + export SWEBENCH_AGENT_WORKERS="${SWEBENCH_AGENT_WORKERS:-32}" + run_eval --port "$PORT" +else + build_replay_cmd "$RESULT_DIR" + REPLAY_CMD+=" --server-metrics http://localhost:$SGLANG_BACKEND_PORT/metrics" + run_agentic_replay_and_write_outputs "$RESULT_DIR" +fi diff --git a/benchmarks/single_node/agentic/glm5.2_fp8_mi355x_sglang.sh b/benchmarks/single_node/agentic/glm5.2_fp8_mi355x_sglang.sh new file mode 100755 index 0000000000..88fadf1fcf --- /dev/null +++ b/benchmarks/single_node/agentic/glm5.2_fp8_mi355x_sglang.sh @@ -0,0 +1,261 @@ +#!/usr/bin/env bash +set -eo pipefail +set -x + +# Agentic trace replay benchmark for GLM-5.2 FP8 on MI355X using SGLang. +# +# Server flags follow the SGLang cookbook MI355X FP8 single-node recipes +# (https://docs.sglang.io/cookbook/autoregressive/GLM/GLM-5.2): TP8 with the +# DSA tilelang prefill/decode backends, no MTP (spec decoding is not yet +# validated on ROCm gfx950). The cookbook's low-latency / balanced / +# high-throughput strategies differ only in batch-shaping levers +# (chunked-prefill / mem-fraction / graph bs / max-running), which this +# script derives from CONC so one search-space arm traces the full frontier. +# +# Required env vars: +# MODEL, TP, CONC, KV_OFFLOADING, TOTAL_CPU_DRAM_GB, RESULT_DIR, DURATION, +# EP_SIZE, DP_ATTENTION +# +# KV_OFFLOADING=dram requires KV_OFFLOAD_BACKEND=hicache. + +source "$(dirname "$0")/../../benchmark_lib.sh" + +check_env_vars MODEL TP CONC KV_OFFLOADING TOTAL_CPU_DRAM_GB RESULT_DIR DURATION EP_SIZE DP_ATTENTION + +if [[ "$DP_ATTENTION" = "true" ]]; then + echo "Error: DP-attention is not part of the GLM-5.2 MI355X cookbook recipe" >&2 + exit 1 +fi + +if [[ -n "$SLURM_JOB_ID" ]]; then + echo "JOB $SLURM_JOB_ID running on $SLURMD_NODENAME" +fi + +# ROCR/HIP visibility under slurm cgroups. +if [ -n "$ROCR_VISIBLE_DEVICES" ]; then + export HIP_VISIBLE_DEVICES="$ROCR_VISIBLE_DEVICES" +fi + +# Weights are pre-staged in the NFS HF hub cache (launch_mi355x-amds.sh mounts +# /it-share/hf-hub-cache for this model); a warm cache makes this a no-op. +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 + + +# A server killed on this node minutes earlier (previous job, crashed run) +# can still be draining its ~1.4 TB of HBM: KFD reclaim takes minutes, and +# booting into a half-drained node fails RCCL init with HIP 'unhandled cuda +# error' / 'invalid argument' (observed as the mooncake-c64 CI failure). +# Wait for the GPUs to come back before launching. +# Per-GPU threshold: idle nodes hold a small driver/firmware VRAM baseline +# (observed up to ~4%/GPU, node-dependent), while a draining or occupied +# GPU sits at 50-90%. Require every GPU <= 10%. +GPU_CLEAN=false +for i in $(seq 1 90); do + VRAM_MAX=$(rocm-smi --showmemuse 2>/dev/null | grep -oE "GPU Memory Allocated \(VRAM%\): [0-9]+" | awk '{if ($NF > m) m = $NF} END {print m+0}') + if [ "${VRAM_MAX:-0}" -le 10 ]; then echo "GPUs clean (vram%max=$VRAM_MAX after $((i*10))s)"; GPU_CLEAN=true; break; fi + echo "waiting for prior-job GPU memory reclaim: vram%max=$VRAM_MAX"; sleep 10 +done +[ "$GPU_CLEAN" = "true" ] || { echo "Error: GPUs still draining prior job's memory after 15min" >&2; exit 1; } + +resolve_trace_source +install_agentic_deps + +SERVER_LOG="$RESULT_DIR/server.log" +mkdir -p "$RESULT_DIR" + +# KV offloading: all backends layer on HiCache, which extends RadixAttention +# by spilling prefixes evicted from the HBM KV pool to a pinned host pool +# instead of recomputing them. GLM-5.2 is DSA-family, so every rank holds +# complete per-token KV (124.54 GB / 1.335M tokens per rank at mem-fraction +# 0.85, replicated on all 8 ranks) and the L2 host pool is sized through the +# host/device token-capacity ratio like the B300 sibling and DSv4 recipes, +# NOT a GB-based --hicache-size (which pins TOTAL_CPU_DRAM_GB/TP per rank +# and OOMs the node - see the B300 #2279 notes). Write policy / IO backend / +# memory layout follow the same-cluster dsv4-fp4-mi355x-sglang-agentic-hicache +# recipe (the ROCm-proven combo for DSA KV pools on this image). +# +# hicache - L2 host pool only, ratio 1.5 (~1.5 TB pinned on the 3 TB +# nodes; the B300 sibling caps at 0.75 against its 170 GB +# device pool, and the MI355X pool is a third smaller, so 1.5 +# doubles host token capacity in the same headroom class). +# mooncake - L2 at ratio 0.5 (~0.5 TB staging tier) + Mooncake store as +# the HiCache L3 storage backend: local mooncake_master, +# embedded per-rank segments over TCP (single node - transfers +# are host-local), 40% of the DRAM budget (~1 TB), keeping +# total pinned memory in the same ~1.5 TB envelope. +# mori - L2 at ratio 0.5 + AMD MORI UMBP (Unified Memory Buffer Pool) +# as the HiCache L3 storage backend: standalone in-process +# per-rank DRAM pools (BUILD_UMBP=ON ships in this image), +# same 40%-of-budget sizing, SSD tier disabled so the arm +# measures the DRAM path against mooncake like-for-like. +CACHE_ARGS=() +if agentic_kv_offload_enabled; then + HICACHE_WRITE_POLICY="${HICACHE_WRITE_POLICY:-write_through}" + HICACHE_IO_BACKEND="${HICACHE_IO_BACKEND:-direct}" + HICACHE_MEM_LAYOUT="${HICACHE_MEM_LAYOUT:-page_first_direct}" + # 40% of the node DRAM budget for the L3 tier, split per rank. + L3_PER_RANK_GB=$((TOTAL_CPU_DRAM_GB * 2 / 5 / TP)) + case "$KV_OFFLOAD_BACKEND" in + hicache) + HICACHE_RATIO="${HICACHE_RATIO:-1.5}" + ;; + mooncake) + HICACHE_RATIO="${HICACHE_RATIO:-0.5}" + python3 -c "from mooncake.store import MooncakeDistributedStore" >/dev/null + MOONCAKE_MASTER_PORT=$((PORT + 12000)) + MOONCAKE_MASTER_LOG="$RESULT_DIR/mooncake_master.log" + MOONCAKE_CONFIG_PATH="$RESULT_DIR/mooncake_config.json" + cat > "$MOONCAKE_CONFIG_PATH" < "$MOONCAKE_MASTER_LOG" 2>&1 & + MOONCAKE_MASTER_PID=$! + sleep 2 + kill -0 "$MOONCAKE_MASTER_PID" + ;; + mori) + HICACHE_RATIO="${HICACHE_RATIO:-0.5}" + python3 -c "import mori.umbp" >/dev/null + ;; + *) + echo "Error: unsupported KV_OFFLOAD_BACKEND '$KV_OFFLOAD_BACKEND' (expected: hicache, mooncake, mori)" >&2 + exit 1 + ;; + esac + echo "HiCache tier: backend=$KV_OFFLOAD_BACKEND, ratio=$HICACHE_RATIO, dram_budget=${TOTAL_CPU_DRAM_GB} GB, l3_per_rank=${L3_PER_RANK_GB} GB, write_policy=$HICACHE_WRITE_POLICY, io_backend=$HICACHE_IO_BACKEND, mem_layout=$HICACHE_MEM_LAYOUT" + 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" + ) + if [[ "$KV_OFFLOAD_BACKEND" == "mooncake" ]]; then + CACHE_ARGS+=(--hicache-storage-backend mooncake) + elif [[ "$KV_OFFLOAD_BACKEND" == "mori" ]]; then + CACHE_ARGS+=( + --hicache-storage-backend mori + --hicache-storage-backend-extra-config + "{\"dram_capacity_bytes\": $((L3_PER_RANK_GB * 1024 * 1024 * 1024)), \"ssd_enabled\": false}" + ) + fi +fi + +# Cookbook batch-shaping by concurrency band: +# low-latency (conc <= 16): chunked-prefill 131072, mem-fraction 0.80 +# balanced/high-throughput (conc >= 32): chunked-prefill 32768, mem 0.85 +# AgentX concurrency counts live session trees, not individual requests, so +# max-running-requests is 2*CONC (subagent fan-out headroom); the CUDA-graph +# range covers it up to the cookbook high-throughput cap of 256. +if [ "$CONC" -le 16 ]; then + # 65536, not the cookbook's 131072: FP8 weights leave only ~12 GiB of + # dynamic headroom above the 0.80 static pool, and a 131k-token chunk's + # ~10 GiB activation burst OOMs on allocator fragmentation about every + # other run ("Tried to allocate 9.91 GiB ... 11.93 GiB is free", 2/3 CI + # failures at conc 16). Halving the chunk halves the burst; the MXFP4 + # sibling keeps 131072 (its smaller weights leave ~40 GiB more slack). + CHUNKED_PREFILL_SIZE=65536 + MEM_FRACTION_STATIC=0.80 +else + CHUNKED_PREFILL_SIZE=32768 + MEM_FRACTION_STATIC=0.85 + # MI355X prefill is slow relative to the 1M-context agentic corpus; give + # the warmup drain the same extended grace as the B300 saturation arm. + export AGENTIC_WARMUP_GRACE_PERIOD=3600 +fi +MAX_RUNNING_REQUESTS=$((2 * CONC)) +CUDA_GRAPH_MAX_BS=$MAX_RUNNING_REQUESTS +[ "$CUDA_GRAPH_MAX_BS" -gt 256 ] && CUDA_GRAPH_MAX_BS=256 + +export PYTHONNOUSERSITE=1 +# Agentic warmup dispatches hundreds of large prompts at once; allow up to +# 15 minutes of TCP progress before AIPerf declares a connection dead. +export AIPERF_HTTP_TCP_USER_TIMEOUT=900000 +# AIPerf pins one pooled keep-alive connection per session (client-side +# keep-alive 300s) while uvicorn's default SGLANG_TIMEOUT_KEEP_ALIVE is 5s; +# inter-turn idle gaps can reuse a socket exactly as the server closes it. +# Outlast the client pool so the race cannot occur. +export SGLANG_TIMEOUT_KEEP_ALIVE=900 +# The DSA indexer's top-k v2 kernel (default since v0.5.14) is JIT-compiled +# from CUDA-only source (cooperative_groups.h) and cannot build for gfx950, +# killing decode CUDA-graph capture. v1 dispatches to the precompiled HIP op +# in sgl-kernel instead; upstream's own MI355X CI runs DSA models this way +# (scripts/ci/slurm/launch_mi355x.sh). +export SGLANG_OPT_USE_TOPK_V2=false + +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 + --tp "$TP" + --ep-size "$EP_SIZE" + --dsa-prefill-backend tilelang + --dsa-decode-backend tilelang + # GLM-5.2 emits the GLM-4.7-style tool-call format; glm47 is required for + # structured message.tool_calls (SWE-bench agentic evals die without it). + # The glm45 reasoning parser keeps hybrid thinking in reasoning_content. + --tool-call-parser glm47 + --reasoning-parser glm45 + --chunked-prefill-size "$CHUNKED_PREFILL_SIZE" + --mem-fraction-static "$MEM_FRACTION_STATIC" + --max-running-requests "$MAX_RUNNING_REQUESTS" + --cuda-graph-max-bs "$CUDA_GRAPH_MAX_BS" + "${CACHE_ARGS[@]}" + --watchdog-timeout 1800 + --enable-metrics +) + +printf '%q ' "${SGLANG_CMD[@]}" | tee "$RESULT_DIR/sglang_command.txt" +printf '\n' | tee -a "$RESULT_DIR/sglang_command.txt" + +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" + +if [ "${EVAL_ONLY}" = "true" ]; then + # GLM-5.2's chat template defaults to reasoning_effort=Max when the + # client passes no chat_template_kwargs (mini-swe-agent doesn't), and the + # heavy thinking burns the default 75-step budget before submission + # (observed on the B300 sibling recipe: 12/23 trajectories exited + # LimitsExceeded unsubmitted while 10 of the 11 that submitted resolved). + # Double the step budget for this recipe; others keep the shared default. + export SWEBENCH_AGENT_STEP_LIMIT=150 + # Pin eval agent parallelism to the proven-green level: workers default + # to CONC, and at 64 concurrent Modal sandboxes the cluster's egress + # collapses (18k "Cannot connect to *.modal.host" errors crippled the + # trajectories in run 29764760177) while 32 ran clean. The serving + # config is unchanged - only the agent's session fan-out is capped. + export SWEBENCH_AGENT_WORKERS="${SWEBENCH_AGENT_WORKERS:-32}" + run_eval --port "$PORT" +else + build_replay_cmd "$RESULT_DIR" + REPLAY_CMD+=" --server-metrics http://localhost:$PORT/metrics" + run_agentic_replay_and_write_outputs "$RESULT_DIR" +fi diff --git a/configs/amd-master.yaml b/configs/amd-master.yaml index ac4e276906..be05a1a89f 100644 --- a/configs/amd-master.yaml +++ b/configs/amd-master.yaml @@ -1551,6 +1551,81 @@ dsv4-fp4-mi355x-sglang-agentic-hicache: - { 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: [64] } +# GLM-5.2 MI355X FP8 AgentX frontier from the SGLang cookbook MI355X FP8 +# single-node recipes (https://docs.sglang.io/cookbook/autoregressive/GLM/GLM-5.2): +# TP8 with DSA tilelang prefill/decode backends, STP only (MTP/EAGLE is not +# yet validated on ROCm gfx950). The cookbook's low-latency / balanced / +# high-throughput strategies differ only in batch-shaping levers, which the +# benchmark script derives from CONC, so a single TP8 arm traces the frontier. +glm5.2-fp8-mi355x-sglang-agentic: + # Image: pinned to the build the DSv4 MI355X agentic recipe runs green on + # (same DSA architecture family, same harness/corpus/cluster). The + # cookbook-verified v0.5.13.post1-...-20260618 build segfaults in AITER's + # gemm_a8w8_blockscale_bpreshuffle under the agentic corpus (GLM-5.2 gfx950 + # tuned configs are not upstream yet - ROCm/aiter#4243). Builds >= v0.5.14 + # default the DSA indexer to the topk v2 kernel, whose JIT source includes + # the CUDA-only cooperative_groups.h and cannot build for gfx950; the + # benchmark script forces the precompiled-HIP v1 path with + # SGLANG_OPT_USE_TOPK_V2=false (mirroring upstream's own MI355X CI). + # KV offloading (all arms layer on HiCache; sizing rationale in the + # benchmark script): the device pool holds only 1.335M tokens against + # 131k-token mean prompts, so the radix-cache cliff bites around conc + # 10-16 - much earlier than the B300 sibling - and host tiers are what + # extend the frontier past it. Three backends are traced: hicache + # (L2 host pool only), mooncake (L2 + Mooncake store as HiCache L3), + # and mori (L2 + AMD MORI UMBP as HiCache L3). The primary hicache arm + # runs the full conc list; the storage arms sample the post-cliff band + # where the L3 tier actually engages. + image: lmsysorg/sglang-rocm:v0.5.14-rocm720-mi35x-20260710 + model: zai-org/GLM-5.2-FP8 + model-prefix: glm5.2 + runner: cluster:mi355x-amds + precision: fp8 + framework: sglang + multinode: false + scenarios: + agentic-coding: + - dram-utilization: 0.80 + search-space: + # TEMPORARY single-point matrix while cluster capacity is constrained + # (2026-07-21): full lists were hicache [1-64], mooncake/mori [16,32,64]. + # Restore before merge. + - { tp: 8, ep: 1, kv-offloading: dram, kv-offload-backend: { name: hicache }, conc-list: [16] } + +# GLM-5.2 MXFP4 (amd/GLM-5.2-MXFP4) MI355X frontier: the FP8 sibling above is +# prefill-starved (TTFT p50 ~390 s at conc 16) because ROCm DSA defaults to +# bf16 KV and TP-mode DSA replicates the pool on every rank (1.335M tokens +# aggregate vs 131k-token mean prompts). MXFP4 halves the weight bytes, +# growing the pool to 1.96M tokens/rank, and the DP-attention arm at conc +# >= 32 un-replicates it (~15.7M tokens aggregate at dp8). All points run +# HiCache + Mooncake-store L3 at ratio 1 (cookbook's ratio 2 pins 2.9 TB of +# host KV before the 40.33 GB/rank DSA indexer sidecars and OOMs the 3 TB +# nodes; rationale in the benchmark script). +glm5.2-fp4-mi355x-sglang-agentic: + image: lmsysorg/sglang-rocm:v0.5.14-rocm720-mi35x-20260710 + model: amd/GLM-5.2-MXFP4 + model-prefix: glm5.2 + runner: cluster:mi355x-amds + precision: fp4 + framework: sglang + multinode: false + scenarios: + agentic-coding: + - dram-utilization: 0.80 + search-space: + # DP-attention arm (tp8/dp8/ep8, conc >= 32) is BLOCKED upstream: the + # DSA DP path hangs a collective under long-context prefill on ROCm + # v0.5.14 - 0/96 storm completions until the watchdog kills the + # scheduler, reproduced on two nodes with hicache ratios 1.5/1/0.5, + # with the DSv4 DP collective envs, and with hicache fully disabled + # (short prompts work; the hang is prefill-load-specific). The + # benchmark script keeps dormant DEP support for when upstream fixes + # it; until then the TP arm extends to conc 64 for high-conc coverage. + # TEMPORARY single-point matrix while cluster capacity is constrained + # (2026-07-21): full list was [1, 2, 4, 8, 16, 24, 32, 48, 64]. + # Restore before merge. + - { tp: 8, ep: 1, kv-offloading: dram, kv-offload-backend: { name: mooncake }, conc-list: [8] } + # MiniMax-M3 MXFP8 MI355X recipe: # https://github.com/vllm-project/recipes/commit/2a3728ed9892debfd767a72a58ebc90b33f186e5 # MXFP8 runs from TP=4 on gfx950; block size 128 is mandatory for MSA. diff --git a/perf-changelog.yaml b/perf-changelog.yaml index 892e3ee2b6..442b654d51 100644 --- a/perf-changelog.yaml +++ b/perf-changelog.yaml @@ -4974,3 +4974,23 @@ - "Run 29651235293 showed the 1M-context corpus working set outgrowing the HBM KV pool past conc 8 (TP8) / conc 64 (DP8): gpu_kv_cache_usage pinned at 1.0 and the radix hit rate collapsed from a ~0.97 theoretical ceiling to 0.04-0.06, so every post-knee turn re-prefilled its full history and throughput fell together with interactivity" - "HiCache spills evicted prefixes to host DRAM and restores them at C2C bandwidth instead of recomputing; sizing follows the qwen3.5-fp8-b300-sglang-agentic-hicache recipe (GLM-5.2 is plain GQA: one host pool per rank, GB-based --hicache-size)" pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2280 + +- config-keys: + - glm5.2-fp8-mi355x-sglang-agentic + description: + - "Add GLM-5.2 FP8 MI355X SGLang single-node agentic benchmarks from the SGLang cookbook MI355X FP8 recipes (TP8, DSA tilelang backends, STP only - MTP not validated on ROCm)" + - "Image: lmsysorg/sglang-rocm:v0.5.14-rocm720-mi35x-20260710, the build the DSv4 MI355X agentic recipe runs green on (the cookbook-verified 20260618 build segfaults in AITER gemm_a8w8_blockscale_bpreshuffle under the agentic corpus - GLM-5.2 gfx950 tuned configs not upstream yet, ROCm/aiter#4243); builds >= v0.5.14 default the DSA indexer topk to a CUDA-only JIT kernel (cooperative_groups.h, cannot build for gfx950), so the recipe forces the precompiled-HIP v1 path with SGLANG_OPT_USE_TOPK_V2=false, mirroring upstream MI355X CI; model zai-org/GLM-5.2-FP8" + - "TP8 arms with DRAM KV offloading on all points, three backends: hicache (HiCache L2 host pool only, ratio 1.5 = ~1.5 TB pinned against the 124.54 GB/rank DSA device pool on 3 TB nodes), mooncake (L2 ratio 0.5 + Mooncake store as HiCache L3: local mooncake_master, embedded per-rank TCP segments, 40% of the DRAM budget), and mori (L2 ratio 0.5 + AMD MORI UMBP standalone per-rank DRAM pools as HiCache L3, SSD tier disabled, same 40% budget) - both storage backends ship in the pinned ROCm image" + - "hicache arm traces the full frontier at conc [1-64]; mooncake/mori arms sample the post-cliff band [16, 32, 64] where the L3 tier engages (the 1.335M-token device pool cliffs around conc 10-16 on the 131k-token mean corpus); L2 write policy/IO/mem-layout follow the ROCm-proven dsv4-fp4-mi355x-sglang-agentic-hicache combo (write_through, direct, page_first_direct); cookbook low-latency vs balanced/high-throughput batch-shaping derived from CONC in the benchmark script" + - "launch_mi355x-amds.sh: GLM-5.2 weights served from the NFS hf-hub-cache (MiniMax-M3 pattern); exclude broken-but-idle mia1-p01-g09/g11 from single-node salloc" + pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2272 + +- config-keys: + - glm5.2-fp4-mi355x-sglang-agentic + description: + - "Add GLM-5.2 MXFP4 (amd/GLM-5.2-MXFP4) MI355X SGLang agentic benchmarks attacking the FP8 frontier's prefill starvation: MXFP4 halves weight bytes (438 GB vs 756 GB), growing the bf16-KV device pool from 1.335M to 1.96M tokens/rank (182.70 GB)" + - "TP arm at conc [1-64] keeps the cookbook batch-shaping bands; every point runs HiCache with the Mooncake store as L3 (write_through, io direct, page_first_direct, wait_complete prefetch, local mooncake_master with 40 GB/rank embedded TCP segments). The intended DP-attention arm (TP8/DP8/EP8, conc >= 32, which un-replicates the DSA KV pool to ~13.7M aggregate tokens) is BLOCKED upstream: the DSA DP path hangs a collective under long-context prefill on ROCm v0.5.14 (watchdog kills the scheduler with 0 storm completions; reproduced on two nodes, with hicache ratios 1.5/1/0.5, with the DSv4 DP collective envs, and with hicache disabled entirely; short prompts work) - the benchmark script keeps dormant DEP support for when upstream fixes it" + - "hicache-ratio measured on-node per arm: the cookbook's ratio 2 pins 2 x 182.7 GB x 8 = 2.9 TB of host KV before GLM-5.2's DSA hicache allocates its additional 40.33 GB/rank indexer-K sidecar host pools and OOMs the 3 TB nodes; the TP arm ships at ratio 1.5 (validated boot + conc-24 long-context storm; its working set oversubscribes the device pool ~3x so the host tier carries the hits) and the DP-attention arm at ratio 0.5 (each DP rank's sessions nearly fit in its own device pool at conc >= 32 so L2 only absorbs overflow; 1.5 boots but the host OOM killer takes the server mid-storm at conc 48)" + - "Both MI355X recipe scripts gain a boot-time GPU-reclaim guard: killing a ~1.4 TB-HBM server takes minutes of KFD drain, and booting within that window fails RCCL init with HIP 'unhandled cuda error' (the mooncake-c64 failure signature in run 29694796593); the guard waits for rocm-smi to report the GPUs clean before launching" + - "launch_mi355x-amds.sh: widen the NFS hf-hub-cache mount match to amd/GLM-5.2* checkpoints" + pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2272 diff --git a/runners/launch_mi355x-amds.sh b/runners/launch_mi355x-amds.sh index 75a67c1c58..bfd6b91ebb 100644 --- a/runners/launch_mi355x-amds.sh +++ b/runners/launch_mi355x-amds.sh @@ -268,7 +268,10 @@ else export GPU_COUNT="${GPU_COUNT:-${TP:?TP must be set}}" set -x - salloc --partition=$PARTITION --gres=gpu:$GPU_COUNT --exclusive --cpus-per-task=128 --time=500 --no-shell --job-name="$RUNNER_NAME" + # mia1-p01-g09 and mia1-p01-g11 are broken but sit idle in slurm (g09 was + # already excluded on the multi-node disagg path); keep single-node jobs + # off them too. + 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" JOB_ID=$(squeue --name="$RUNNER_NAME" -h -o %A | head -n1) srun --jobid=$JOB_ID bash -c "docker stop \$(docker ps -a -q)" @@ -306,6 +309,13 @@ else export HF_HUB_CACHE_MOUNT="/it-share/hf-hub-cache/" fi + # GLM-5.2-FP8 (~756 GB) is pre-staged once on the NFS share rather than + # re-downloaded onto each node's local NVMe cache (same pattern as + # MiniMax-M3 above). + if [[ "$MODEL" == zai-org/GLM-5.2* || "$MODEL" == amd/GLM-5.2* ]]; then + export HF_HUB_CACHE_MOUNT="/it-share/hf-hub-cache/" + fi + SCRIPT_BASE="${EXP_NAME%%_*}_${PRECISION}_mi355x" SCRIPT_FW="benchmarks/single_node/${SCENARIO_SUBDIR:-fixed_seq_len/}${SCRIPT_BASE}_${FRAMEWORK}${SPEC_SUFFIX}.sh" SCRIPT_FALLBACK="benchmarks/single_node/${SCENARIO_SUBDIR:-fixed_seq_len/}${SCRIPT_BASE}${FRAMEWORK_SUFFIX}${SPEC_SUFFIX}.sh"