diff --git a/.github/workflows/benchmark-tmpl.yml b/.github/workflows/benchmark-tmpl.yml index 3b2614e145..8240d00d68 100644 --- a/.github/workflows/benchmark-tmpl.yml +++ b/.github/workflows/benchmark-tmpl.yml @@ -202,6 +202,14 @@ jobs: done fi + # MI325X AgentX artifacts can occupy hundreds of MB per job. They have + # already been uploaded when this anchor runs post-job; removing the + # prior workspace copy pre-job also recovers interrupted-run output. + if [[ "$RUNNER_NAME" == mi325x-amds_* ]]; then + echo "[MI325X] Cleaning benchmark workspace results ..." + rm -rf -- "$GITHUB_WORKSPACE/results" + fi + # Cleanup SLURM resources if command -v squeue >/dev/null 2>&1; then echo "[Slurm] Cleaning up jobs with name: ${{ runner.name }} ..." diff --git a/benchmarks/multi_node/amd_utils/env.sh b/benchmarks/multi_node/amd_utils/env.sh index baa917b449..1b4b4f02c4 100755 --- a/benchmarks/multi_node/amd_utils/env.sh +++ b/benchmarks/multi_node/amd_utils/env.sh @@ -236,12 +236,18 @@ else export SGLANG_HEALTH_CHECK_TIMEOUT=600 # GLM-5: uses NSA (not MLA), needs fused-decode-MLA disabled + fast loading - if [[ "$MODEL_NAME" == "GLM-5-FP8" ]]; then + MODEL_CONFIG_KEY="${MODEL_YAML_KEY:-$MODEL_NAME}" + if [[ "$MODEL_CONFIG_KEY" == "GLM-5-FP8" || "$MODEL_CONFIG_KEY" == "GLM-5.2-FP8" ]]; then export SGLANG_ROCM_FUSED_DECODE_MLA=0 export ROCM_QUICK_REDUCE_QUANTIZATION=INT4 export SAFETENSORS_FAST_GPU=1 fi + if [[ "$MODEL_CONFIG_KEY" == "GLM-5.2-FP8" ]]; then + export SGLANG_DSA_FUSE_TOPK=false + export SGLANG_OPT_USE_TOPK_V2=false + fi + # Disable allocating memory in one pass export MORI_SHMEM_MODE=ISOLATION diff --git a/benchmarks/multi_node/amd_utils/job.slurm b/benchmarks/multi_node/amd_utils/job.slurm index 0ff5a2b169..fb53557143 100755 --- a/benchmarks/multi_node/amd_utils/job.slurm +++ b/benchmarks/multi_node/amd_utils/job.slurm @@ -42,14 +42,14 @@ if [[ -z "${DOCKER_IMAGE_NAME:-}" ]]; then fi MODEL_NAME="${MODEL_NAME:-None}" -# Resolve the models.yaml entry the same way server_sglang.sh does: agentic runs -# (IS_AGENTIC) use the '-AgentX' recipe, non-agentic disaggregated runs use -# '-DI'. Fall back to the bare model name if the variant key is absent. -# MODEL_NAME itself is left unchanged so env.sh/server_sglang.sh still see the base name. -if [[ "${IS_AGENTIC:-0}" == "1" || "${IS_AGENTIC:-}" == "true" ]]; then - MODEL_YAML_KEY="${MODEL_NAME}-AgentX" -else - MODEL_YAML_KEY="${MODEL_NAME}-DI" +# Respect an explicit recipe key (needed when MODEL_NAME is a cache path), then +# follow the standard AgentX/DI naming convention used by current AMD recipes. +if [[ -z "${MODEL_YAML_KEY:-}" ]]; then + if [[ "${IS_AGENTIC:-0}" == "1" || "${IS_AGENTIC:-}" == "true" ]]; then + MODEL_YAML_KEY="${MODEL_NAME}-AgentX" + else + MODEL_YAML_KEY="${MODEL_NAME}-DI" + fi fi if ! grep -q "^${MODEL_YAML_KEY}:" "$MODELS_YAML"; then if grep -q "^${MODEL_NAME}:" "$MODELS_YAML"; then @@ -117,7 +117,7 @@ export MODEL_DIR if [[ "$ENGINE" == "vllm-disagg" ]]; then # vLLM: Extract hf_dir from models.yaml, search multiple paths, resolve HF cache snapshots - DISK_DIR_NAME=$(awk '/^'"$MODEL_NAME"':/{found=1; next} + DISK_DIR_NAME=$(awk -v key="$MODEL_YAML_KEY" '$0 == key ":" {found=1; next} found && /^[^ ]/{exit} found && /hf_dir:/{gsub(/[" ]/, "", $2); print $2; exit}' "$MODELS_YAML") DISK_DIR_NAME="${DISK_DIR_NAME:-$MODEL_NAME}" @@ -193,7 +193,7 @@ else } # Extract hf_dir from models.yaml (same as vllm-disagg path above) - SGL_DISK_DIR_NAME=$(awk '/^'"$MODEL_NAME"':/{found=1; next} + SGL_DISK_DIR_NAME=$(awk -v key="$MODEL_YAML_KEY" '$0 == key ":" {found=1; next} found && /^[^ ]/{exit} found && /hf_dir:/{gsub(/[" ]/, "", $2); print $2; exit}' "$MODELS_YAML") SGL_DISK_DIR_NAME="${SGL_DISK_DIR_NAME:-$MODEL_NAME}" @@ -243,6 +243,11 @@ echo "NUM_NODES: $NUM_NODES (xP=$xP + yD=$yD)" FULL_NODELIST=$(scontrol show hostnames "$SLURM_JOB_NODELIST") SELECTED_NODES=$(echo "$FULL_NODELIST" | head -n $NUM_NODES) SELECTED_NODELIST_STR=$(echo "$SELECTED_NODES" | tr '\n' ',' | sed 's/,$//') +SELECTED_NODE_COUNT=$(echo "$SELECTED_NODES" | sed '/^$/d' | wc -l) +if [[ "$SELECTED_NODE_COUNT" -ne "$NUM_NODES" ]]; then + echo "FATAL ERROR: allocation exposes ${SELECTED_NODE_COUNT} nodes, expected ${NUM_NODES}" >&2 + exit 1 +fi # Docker privilege detection — evaluated per-node since group membership varies. # Exported as a snippet so every srun participant resolves it locally. @@ -267,16 +272,27 @@ echo "Selected nodes: $SELECTED_NODELIST_STR" USER_NAME=$(whoami) MASTER_NODE=$(echo "$SELECTED_NODES" | head -n 1) -NODE0_ADDR=$(srun --nodes=1 --ntasks=1 --time=00:20:00 --nodelist="$MASTER_NODE" bash -c 'ip route get 1.1.1.1') -NODE0_ADDR=$(echo "$NODE0_ADDR" | awk '/src/ {print $7}') - IPS=() for NODE in $SELECTED_NODES; do - IP=$(srun --nodes=1 --ntasks=1 --time=00:20:00 --nodelist="$NODE" bash -c 'ip route get 1.1.1.1') - IP=$(echo "$IP" | awk '/src/ {print $7}') + if ! IP_ROUTE=$(srun --nodes=1 --ntasks=1 --time=00:05:00 --nodelist="$NODE" \ + bash -c 'ip route get 1.1.1.1'); then + echo "FATAL ERROR: unable to launch IP-resolution task on $NODE" >&2 + exit 1 + fi + IP=$(echo "$IP_ROUTE" | awk '/src/ {print $7; exit}') + if [[ -z "$IP" ]]; then + echo "FATAL ERROR: no source IP resolved for $NODE" >&2 + exit 1 + fi IPS+=("$IP") done +if [[ "${#IPS[@]}" -ne "$NUM_NODES" ]]; then + echo "FATAL ERROR: resolved ${#IPS[@]} node IPs, expected ${NUM_NODES}" >&2 + exit 1 +fi +NODE0_ADDR="${IPS[0]}" + echo "Node IPs: ${IPS[*]}" DOCKER_MOUNT_PATH="/workspace" @@ -332,7 +348,8 @@ export SPEC_DECODING="${SPEC_DECODING:-}" export IS_MULTINODE="${IS_MULTINODE:-false}" SANITIZED_USER=$(echo "$USER_NAME" | tr -c 'a-zA-Z0-9_.-' '_') -export DOCKER_CONT_NAME="container_${ENGINE}_${SANITIZED_USER}_${MODEL_NAME}_${SLURM_JOB_ID}" +SANITIZED_MODEL=$(printf '%s' "$MODEL_NAME" | tr -c 'a-zA-Z0-9_.-' '_') +export DOCKER_CONT_NAME="container_${ENGINE}_${SANITIZED_USER}_${SANITIZED_MODEL}_${SLURM_JOB_ID}" # vLLM external router container. # NOTE: vllm/vllm-router only retains ~16 recent nightlies on Docker Hub; older @@ -384,6 +401,7 @@ DOCKER_ENV_COMMON=( -e NODE0_ADDR=\$NODE0_ADDR -e MODEL_DIR=/models -e MODEL_NAME=\$MODEL_NAME + -e MODEL_YAML_KEY=\$MODEL_YAML_KEY -e GPUS_PER_NODE=\$GPUS_PER_NODE -e xP=\$xP -e yD=\$yD @@ -445,6 +463,17 @@ DOCKER_ENV_COMMON=( -e DECODE_ENABLE_DP=\$DECODE_ENABLE_DP -e DECODE_MTP_SIZE=\$DECODE_MTP_SIZE -e IS_MULTINODE=\$IS_MULTINODE + -e SCENARIO_TYPE=\$SCENARIO_TYPE + -e CONC=\$CONC + -e CONC_LIST + -e ROUTER_METADATA + -e KV_P2P_TRANSFER=\$KV_P2P_TRANSFER + -e MODEL=\$MODEL + -e WEKA_LOADER_OVERRIDE=\$WEKA_LOADER_OVERRIDE + -e SYNC_BARRIER_TIMEOUT + -e HF_HUB_CACHE=/models + -e IBDEVICES + -e MORI_RDMA_TC -e DRY_RUN=\${DRY_RUN:-0} ) @@ -539,12 +568,27 @@ if [[ -n "${CLIENT_IMAGE:-}" ]]; then srun --nodelist="$SELECTED_NODELIST_SRUN" bash -c 'eval "$DOCKER_CMD_DETECT"; $DOCKER_CMD pull '"$CLIENT_IMAGE"' >/dev/null 2>&1 || true' 2>/dev/null || true fi +echo "Ensuring Docker image is present on every selected node..." +srun \ + --nodelist="$SELECTED_NODELIST_SRUN" \ + --kill-on-bad-exit=1 \ + --unbuffered \ + bash -c ' +set -euo pipefail +eval "$DOCKER_CMD_DETECT" +if ! $DOCKER_CMD image inspect "$DOCKER_IMAGE_NAME" >/dev/null 2>&1; then + echo "[docker-pull] $(hostname): $DOCKER_IMAGE_NAME" + $DOCKER_CMD pull "$DOCKER_IMAGE_NAME" +fi +' + +SERVER_RUN_RC=0 srun \ --nodelist="$SELECTED_NODELIST_SRUN" \ --kill-on-bad-exit=1 \ --signal=TERM@30 \ --unbuffered \ - bash -lc " + bash -c " set -euo pipefail echo \"Rank \$SLURM_PROCID on \$(hostname)\" @@ -635,6 +679,8 @@ fi # end: if ENGINE == atom-disagg \$DOCKER_CMD ps -aq --filter \"$CONT_FILTER\" | xargs -r \$DOCKER_CMD rm -f || true \$DOCKER_CMD ps -aq | xargs -r \$DOCKER_CMD stop || true +echo \"[docker-run] image=$DOCKER_IMAGE_NAME container=$DOCKER_CONT_NAME\" + # Start vLLM external router container on node 0 if [[ \"$ENGINE\" == \"vllm-disagg\" && \"$ROUTER_TYPE\" == \"vllm-router\" && \"\$SLURM_PROCID\" == \"0\" ]]; then \$DOCKER_CMD rm -f \"$ROUTER_CONT_NAME\" 2>/dev/null || true @@ -694,7 +740,7 @@ fi -v ${MODEL_DIR}:/models \ -v \$HOME/.ssh:/root/.ssh \ --shm-size 128G \ - -v /tmp:/run_logs \ + -v ${BENCHMARK_LOGS_DIR}:/run_logs \ -v ${BENCHMARK_LOGS_DIR}:/benchmark_logs \ -v ${DI_REPO_DIR}:${DOCKER_MOUNT_PATH} \ -v ${HICACHE_MC_CONFIG}:/config/hicache_mc.env:ro \ @@ -717,7 +763,7 @@ DOCKER_EXIT_CODE=\$? echo \"[rank 0] Main container exited (rc=\$DOCKER_EXIT_CODE). Stopping vllm-router...\" \$DOCKER_CMD rm -f \"$ROUTER_CONT_NAME\" 2>/dev/null || true exit \$DOCKER_EXIT_CODE -" +" || SERVER_RUN_RC=$? if [[ "${KEEP_CONTAINERS}" != "1" ]]; then srun --nodelist="$SELECTED_NODELIST_SRUN" bash -c 'eval "$DOCKER_CMD_DETECT"; $DOCKER_CMD rm -f '"$DOCKER_CONT_NAME"' '"$CLIENT_CONT_NAME"' 2>/dev/null || true' @@ -729,3 +775,8 @@ if [[ "${KEEP_CONTAINERS}" != "1" ]]; then ' fi fi + +if [[ "$SERVER_RUN_RC" -ne 0 ]]; then + echo "ERROR: distributed server step failed with exit code $SERVER_RUN_RC" >&2 + exit "$SERVER_RUN_RC" +fi diff --git a/benchmarks/multi_node/amd_utils/models.yaml b/benchmarks/multi_node/amd_utils/models.yaml index a3b3436400..1cc22b3f8c 100644 --- a/benchmarks/multi_node/amd_utils/models.yaml +++ b/benchmarks/multi_node/amd_utils/models.yaml @@ -264,6 +264,38 @@ GLM-5-FP8: chunked_prefill_size: 262144 cuda_graph_bs_range: "1-128" +GLM-5.2-FP8: + base_flags: "--decode-log-interval 1000 --log-level info --watchdog-timeout 3600 --load-balance-method round_robin --served-model-name zai-org/GLM-5.2-FP8 --enable-metrics --disaggregation-transfer-backend mori --tool-call-parser glm47 --reasoning-parser glm45 --dsa-prefill-backend tilelang --dsa-decode-backend tilelang --dsa-topk-backend torch --kv-cache-dtype bfloat16 --context-length 1048576 --max-total-tokens 1048576 --model-loader-extra-config '{\\\"enable_multithread_load\\\": true, \\\"num_threads\\\": 8}'" + mtp_flags: "" + 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.85 + disable_radix_cache: false + dp: + max_running_requests: 16 + chunked_prefill_size: 131072 + cuda_graph_bs: "1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16" + no_dp: + max_running_requests: 16 + chunked_prefill_size: 131072 + cuda_graph_bs_range: "1-16" + decode: + mem_fraction_static: 0.85 + prefill_round_robin_balance: true + dp: + max_running_requests: 16 + chunked_prefill_size: 131072 + cuda_graph_bs_range: "1-16" + ep_only: + max_running_requests: 16 + chunked_prefill_size: 131072 + cuda_graph_bs_range: "1-16" + no_dp: + max_running_requests: 16 + chunked_prefill_size: 131072 + cuda_graph_bs_range: "1-16" + 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 549f83fdd1..cc30a7f94f 100755 --- a/benchmarks/multi_node/amd_utils/server_sglang.sh +++ b/benchmarks/multi_node/amd_utils/server_sglang.sh @@ -81,12 +81,13 @@ import yaml, sys, os config_path = '${MODELS_YAML}' model_name = '${MODEL_NAME}' +explicit_model_key = '${MODEL_YAML_KEY:-}'.strip() # Select the models.yaml recipe variant by run type: agentic runs (IS_AGENTIC) # use the '-AgentX' entry, non-agentic disaggregated runs use '-DI'. # Fall back to the bare model name if the variant-specific key is absent. is_agentic = '${IS_AGENTIC:-0}'.strip().lower() in ('1', 'true') -model_key = f'{model_name}-AgentX' if is_agentic else f'{model_name}-DI' +model_key = explicit_model_key or (f'{model_name}-AgentX' if is_agentic else f'{model_name}-DI') with open(config_path) as f: models = yaml.safe_load(f) @@ -604,7 +605,18 @@ wait_or_die() { # $1 = server pid to watch; rest = blocking command "$@" & local cmd=$! while kill -0 "$cmd" 2>/dev/null; do kill -0 "$watch" 2>/dev/null || { - echo "FATAL: $(hostname) local sglang server (pid $watch) died; tearing down job" >&2 + local server_rc=0 + wait "$watch" || server_rc=$? + sleep 1 + echo "FATAL: $(hostname) local sglang server (pid $watch) died with rc=$server_rc; tearing down job" >&2 + local log_dir="/run_logs/slurm_job-${SLURM_JOB_ID}" + local log_file + for log_file in "$log_dir"/*_"$(hostname)".log; do + [[ -f "$log_file" ]] || continue + echo "===== Local SGLang log: $log_file =====" >&2 + tail -n 200 "$log_file" >&2 || true + echo "===== End local SGLang log =====" >&2 + done kill "$cmd" 2>/dev/null || true return 1 } diff --git a/benchmarks/multi_node/amd_utils/setup_deps.sh b/benchmarks/multi_node/amd_utils/setup_deps.sh index 8fa49dfb27..6c2d6666b9 100644 --- a/benchmarks/multi_node/amd_utils/setup_deps.sh +++ b/benchmarks/multi_node/amd_utils/setup_deps.sh @@ -78,6 +78,72 @@ install_amd_quark() { _SETUP_INSTALLED+=("amd-quark") } +# --------------------------------------------------------------------------- +# SGLang: Patch aiter gluon pa_mqa_logits — fix 2D → 3D instr_shape for +# Triton >= 3.5. The base non-preshuffle variant omitted the conditional used +# by the other variants; GLM-5.2 exercises this path on gfx942. +# --------------------------------------------------------------------------- +patch_gluon_pa_mqa_logits_instr_shape() { + python3 -c ' +import os, sys + +target = "/sgl-workspace/aiter/aiter/ops/triton/gluon/pa_mqa_logits.py" +if not os.path.isfile(target): + print("[SETUP] gluon pa_mqa_logits.py not found, skipping") + sys.exit(0) + +src = open(target).read() +if "[PATCHED] 3D instr_shape for base gluon variant" in src: + print("[SETUP] gluon pa_mqa_logits 3D instr_shape patch already applied") + sys.exit(0) + +old = """\ + mfma_layout: gl.constexpr = gl.amd.AMDMFMALayout( + version=CDNA_VERSION, + instr_shape=[16, 16], + transposed=False, + warps_per_cta=[1, NumWarps], + ) + mfma_layout_a: gl.constexpr = gl.DotOperandLayout( + operand_index=0, parent=mfma_layout, k_width=16 + ) + mfma_layout_b: gl.constexpr = gl.DotOperandLayout( + operand_index=1, parent=mfma_layout, k_width=16 + )""" + +new = """\ + # [PATCHED] 3D instr_shape for base gluon variant + if _Use_2d_instr_shape_mfma_layout: + mfma_layout: gl.constexpr = gl.amd.AMDMFMALayout( + version=CDNA_VERSION, + instr_shape=[16, 16], + transposed=False, + warps_per_cta=[1, NumWarps], + ) + else: + mfma_layout: gl.constexpr = gl.amd.AMDMFMALayout( + version=CDNA_VERSION, + instr_shape=[16, 16, 32], + transposed=False, + warps_per_cta=[1, NumWarps], + ) + mfma_layout_a: gl.constexpr = gl.DotOperandLayout( + operand_index=0, parent=mfma_layout, k_width=16 + ) + mfma_layout_b: gl.constexpr = gl.DotOperandLayout( + operand_index=1, parent=mfma_layout, k_width=16 + )""" + +if old not in src: + print("[SETUP] WARN: gluon pa_mqa_logits pattern not found — aiter version may have changed") + sys.exit(0) + +open(target, "w").write(src.replace(old, new, 1)) +print("[SETUP] Patched: gluon pa_mqa_logits 3D instr_shape for base variant") +' + _SETUP_INSTALLED+=("gluon-instr-shape-fix") +} + # --------------------------------------------------------------------------- # SGLang: Install latest transformers for GLM model type support. # @@ -119,6 +185,7 @@ if [[ "$ENGINE" == "vllm-disagg" ]]; then export PATH="${UCX_HOME}/bin:/usr/local/bin/etcd:/root/.cargo/bin:${PATH}" export LD_LIBRARY_PATH="${UCX_HOME}/lib:${RIXL_HOME}/lib:${RIXL_HOME}/lib/x86_64-linux-gnu:${LD_LIBRARY_PATH:-}" else + patch_gluon_pa_mqa_logits_instr_shape install_transformers_glm5 fi diff --git a/benchmarks/multi_node/amd_utils/submit.sh b/benchmarks/multi_node/amd_utils/submit.sh index 262fd25aa0..99f147f91d 100755 --- a/benchmarks/multi_node/amd_utils/submit.sh +++ b/benchmarks/multi_node/amd_utils/submit.sh @@ -121,6 +121,7 @@ export DECODE_MTP_SIZE=${DECODE_MTP_SIZE} export NUM_NODES=$NUM_NODES export GPUS_PER_NODE=$GPUS_PER_NODE export MODEL_NAME=$MODEL_NAME +export MODEL_YAML_KEY="${MODEL_YAML_KEY:-$MODEL_NAME}" export BENCH_INPUT_LEN=${ISL} export BENCH_OUTPUT_LEN=${OSL} export BENCH_NUM_PROMPTS_MULTIPLIER=${BENCH_NUM_PROMPTS_MULTIPLIER:-10} @@ -145,6 +146,20 @@ export RUNNER_TYPE="${RUNNER_TYPE:-}" export RESULT_FILENAME="${RESULT_FILENAME:-}" export SPEC_DECODING="${SPEC_DECODING:-}" export IS_MULTINODE="${IS_MULTINODE:-false}" +export SCENARIO_TYPE="${SCENARIO_TYPE:-fixed-seq-len}" +export IS_AGENTIC="${IS_AGENTIC:-0}" +export CONC="${CONC:-${CONCURRENCIES%%x*}}" +export CONC_LIST="${CONC_LIST:-${CONCURRENCIES//x/ }}" +export DURATION="${DURATION:-3600}" +export KV_OFFLOADING="${KV_OFFLOADING:-none}" +export KV_OFFLOAD_BACKEND="${KV_OFFLOAD_BACKEND:-}" +export KV_OFFLOAD_BACKEND_METADATA="${KV_OFFLOAD_BACKEND_METADATA:-}" +export ROUTER_METADATA="${ROUTER_METADATA:-}" +export KV_P2P_TRANSFER="${KV_P2P_TRANSFER:-}" +export TOTAL_CPU_DRAM_GB="${TOTAL_CPU_DRAM_GB:-0}" +export MAX_MODEL_LEN="${MAX_MODEL_LEN:-1048576}" +export MODEL="${MODEL:-}" +export WEKA_LOADER_OVERRIDE="${WEKA_LOADER_OVERRIDE:-}" # Log directory: must be on NFS (shared filesystem) so the submit host can read SLURM output. export BENCHMARK_LOGS_DIR="${BENCHMARK_LOGS_DIR:-$(pwd)/benchmark_logs}" @@ -166,7 +181,12 @@ fi # Optional: exclude specific nodes (e.g. nodes with broken Docker sockets). # Set SLURM_EXCLUDE_NODES env var to a comma-separated list of hostnames. EXCLUDE_OPT=() -SLURM_EXCLUDE_NODES="${SLURM_EXCLUDE_NODES:-mia1-p01-g11,mia1-p01-g12,mia1-p01-g15}" +if [[ -z "${SLURM_EXCLUDE_NODES+x}" ]]; then + case "$(hostname -s)" in + mia1*) SLURM_EXCLUDE_NODES="mia1-p01-g11,mia1-p01-g12,mia1-p01-g15" ;; + *) SLURM_EXCLUDE_NODES="" ;; + esac +fi if [[ -n "${SLURM_EXCLUDE_NODES:-}" ]]; then EXCLUDE_OPT=(--exclude "$SLURM_EXCLUDE_NODES") fi diff --git a/benchmarks/multi_node/glm5.2_fp8_mi325x_sglang-disagg.sh b/benchmarks/multi_node/glm5.2_fp8_mi325x_sglang-disagg.sh new file mode 100755 index 0000000000..729cf01423 --- /dev/null +++ b/benchmarks/multi_node/glm5.2_fp8_mi325x_sglang-disagg.sh @@ -0,0 +1,47 @@ +#!/usr/bin/env bash +set -euo pipefail + +source "$(dirname "$0")/../benchmark_lib.sh" + +check_env_vars \ + CONC_LIST \ + ISL \ + OSL \ + IMAGE \ + SPEC_DECODING \ + MODEL_PATH \ + MODEL_NAME \ + 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 \ + FRAMEWORK + +cd "$GITHUB_WORKSPACE/benchmarks/multi_node/amd_utils" + +export TIME_LIMIT="08:00:00" +export CONTAINER_IMAGE="$IMAGE" +export DECODE_MTP_SIZE="${DECODE_MTP_SIZE:-0}" +export MAX_MODEL_LEN=1048576 +export WEKA_LOADER_OVERRIDE="${WEKA_LOADER_OVERRIDE:-semianalysis_cc_traces_weka_062126}" + +[[ "$PREFILL_EP" -eq 1 ]] && export PREFILL_ENABLE_EP=false || export PREFILL_ENABLE_EP=true +[[ "$PREFILL_DP_ATTN" == "true" ]] && export PREFILL_ENABLE_DP=true || export PREFILL_ENABLE_DP=false +[[ "$DECODE_EP" -eq 1 ]] && export DECODE_ENABLE_EP=false || export DECODE_ENABLE_EP=true +[[ "$DECODE_DP_ATTN" == "true" ]] && export DECODE_ENABLE_DP=true || export DECODE_ENABLE_DP=false + +bash ./submit.sh \ + "$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" "${NODELIST:-}" diff --git a/benchmarks/single_node/agentic/glm5.2_fp8_mi325x.sh b/benchmarks/single_node/agentic/glm5.2_fp8_mi325x.sh new file mode 100755 index 0000000000..4fc33b12fc --- /dev/null +++ b/benchmarks/single_node/agentic/glm5.2_fp8_mi325x.sh @@ -0,0 +1,146 @@ +#!/usr/bin/env bash +set -euo pipefail +set -x + +# Full-context AgentX qualification for GLM-5.2 FP8 on one 8xMI325X node. +# The SGLang GLM-5.2 cookbook supports TP8 FP8 on MI325X with the DSA +# TileLang backends. This recipe explicitly requests the native 1M context and +# GPU-resident BF16 KV pool; SGLang startup fails if that allocation cannot be +# satisfied. MTP is intentionally disabled because the AMD path is unvalidated. + +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:-unknown}" +fi +if [[ -n "${ROCR_VISIBLE_DEVICES:-}" ]]; then + export HIP_VISIBLE_DEVICES="$ROCR_VISIBLE_DEVICES" +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 + +# GLM-5.2 has a native 1M context, so replay the complete AgentX corpus rather +# than the repository's default 256K-capped corpus for unrecognized families. +export WEKA_LOADER_OVERRIDE=semianalysis_cc_traces_weka_062126 +resolve_trace_source +install_agentic_deps + +SERVER_LOG="$RESULT_DIR/server.log" +mkdir -p "$RESULT_DIR" +SERVER_PID="" +SERVER_PGID="" + +cleanup() { + if [[ -n "$SERVER_PGID" ]]; then + kill -TERM -- "-$SERVER_PGID" 2>/dev/null || true + elif [[ -n "$SERVER_PID" ]] && kill -0 "$SERVER_PID" 2>/dev/null; then + kill -TERM "$SERVER_PID" 2>/dev/null || true + fi + [[ -z "$SERVER_PID" ]] || wait "$SERVER_PID" 2>/dev/null || true +} +trap cleanup EXIT + +CACHE_ARGS=() +if require_agentic_kv_offload_backend hicache; then + # GLM-5.2's DSA KV pool is replicated across the TP ranks. A 0.75 host + # tier adds roughly 0.75M cold-prefix tokens per rank while staying well + # inside the MI325X runner's measured 3 TB CPU-DRAM budget. + HICACHE_RATIO="${HICACHE_RATIO:-0.75}" + CACHE_ARGS=( + --enable-hierarchical-cache + --hicache-ratio "$HICACHE_RATIO" + --hicache-write-policy "${HICACHE_WRITE_POLICY:-write_back}" + --hicache-io-backend "${HICACHE_IO_BACKEND:-direct}" + --hicache-mem-layout "${HICACHE_MEM_LAYOUT:-page_first_direct}" + ) +fi + +export PYTHONNOUSERSITE=1 +export AIPERF_HTTP_TCP_USER_TIMEOUT=900000 +export SGLANG_TIMEOUT_KEEP_ALIVE=900 +# The mi30x image's sgl-kernel DSA top-k JIT includes CUDA's +# cooperative_groups.h while compiling for gfx942. Use SGLang's portable +# Torch fallback and disable both the fused top-k path and its independently +# gated CUDA-graph planning kernel. +export SGLANG_DSA_FUSE_TOPK=false +export SGLANG_OPT_USE_TOPK_V2=false + +PARALLEL_ARGS=(--tp "$TP" --ep-size "$EP_SIZE") + +# Keep the cookbook profiles as separate topology/cache series so the AgentX +# dashboard can compare their complete curves instead of overlaying identical +# labels. All profiles retain the same 1M request limit and 1M HBM KV pool. +PROFILE=low-latency +CHUNKED_PREFILL_ARGS=(--chunked-prefill-size 131072) +MAX_RUNNING_REQUESTS=$((2 * CONC)) +CUDA_GRAPH_ARGS=(--cuda-graph-max-bs "$MAX_RUNNING_REQUESTS") +if [ "$EP_SIZE" -gt 1 ]; then + PROFILE=balanced + CHUNKED_PREFILL_ARGS=(--chunked-prefill-size 32768) + MAX_RUNNING_REQUESTS=80 + CUDA_GRAPH_ARGS=(--cuda-graph-max-bs 128) +elif [ "$KV_OFFLOADING" != "none" ]; then + PROFILE=hicache + CHUNKED_PREFILL_ARGS=(--chunked-prefill-size 32768) + MAX_RUNNING_REQUESTS=80 + CUDA_GRAPH_ARGS=(--cuda-graph-max-bs 128) +fi +echo "GLM-5.2 MI325X AgentX profile: $PROFILE" + +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[@]}" + --dsa-prefill-backend tilelang + --dsa-decode-backend tilelang + --dsa-topk-backend torch + --kv-cache-dtype bfloat16 + --tool-call-parser glm47 + --reasoning-parser glm45 + --context-length 1048576 + --max-total-tokens 1048576 + "${CHUNKED_PREFILL_ARGS[@]}" + --mem-fraction-static 0.85 + --max-running-requests "$MAX_RUNNING_REQUESTS" + "${CUDA_GRAPH_ARGS[@]}" + "${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 MI325X..." +# Keep the launcher and its renamed sglang::scheduler TP children in one +# process group so teardown cannot leave a 200-GB/GPU child behind. +setsid "${SGLANG_CMD[@]}" > "$SERVER_LOG" 2>&1 & +SERVER_PID=$! +SERVER_PGID=$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 + export SWEBENCH_AGENT_STEP_LIMIT=150 + 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..ce702fb9ed 100644 --- a/configs/amd-master.yaml +++ b/configs/amd-master.yaml @@ -2054,6 +2054,112 @@ minimaxm3-fp8-mi355x-vllm-disagg: dp-attn: false additional-settings: - "DECODE_NODES=1" + +# GLM-5.2 FP8 full-context AgentX curves on one MI325X node. Every arm retains +# context-length=1,048,576 and max-total-tokens=1,048,576. The TP8 baseline +# establishes the low-latency knee; EP8 uses the cookbook balanced batch shape; +# HiCache tests whether spilling cold prefixes to the 3 TB host tier restores +# the cache-hit collapse above conc 3. DP attention is excluded: TileLang's +# DSA kernel exceeds gfx942's per-block shared-memory limit, while AITER has no +# decode kernel for GLM-5.2's bf16 GQA-64 geometry in the current image. +glm5.2-fp8-mi325x-sglang-agentic: + image: lmsysorg/sglang:v0.5.15.post1-rocm720-mi30x + model: zai-org/GLM-5.2-FP8 + model-prefix: glm5.2 + runner: cluster:mi325x-amds + precision: fp8 + framework: sglang + multinode: false + scenarios: + agentic-coding: + - dram-utilization: 0.80 + search-space: + - { tp: 8, ep: 1, kv-offloading: none, conc-list: [1, 2, 3, 4, 5, 6, 8] } + - { tp: 8, ep: 8, kv-offloading: none, conc-list: [1, 2, 3, 4, 5, 6, 8, 10, 12, 16] } + - { tp: 8, ep: 1, kv-offloading: dram, kv-offload-backend: { name: hicache }, conc-list: [3, 4, 5, 6, 8, 10, 12, 16] } + +# GLM-5.2 FP8 full-context AgentX disaggregation on MI325X. Follow the +# DeepSeek-V4-Pro AgentX convention: keep each engine/topology to at most four +# sequential concurrency points and use separate topology arms. MI325X's +# Broadcom bnxt_re MoRI all-to-all path is still unstable, so this sweep varies +# TP8 prefill/decode worker ratios without EP or DP attention. +glm5.2-fp8-mi325x-sglang-disagg-agentic: + image: lmsysorg/sglang:v0.5.15.post1-rocm720-mi30x + model: zai-org/GLM-5.2-FP8 + model-prefix: glm5.2 + runner: cluster:mi325x-amds + precision: fp8 + framework: sglang-disagg + router: { name: sglang-router, version: "0.3.2" } + kv-p2p-transfer: mori + multinode: true + disagg: true + scenarios: + agentic-coding: + - search-space: + # Balanced baseline: one TP8 prefill worker and one TP8 decode worker. + - spec-decoding: none + conc-list: [1, 2, 3, 4] + prefill: + num-worker: 1 + tp: 8 + ep: 1 + dp-attn: false + additional-settings: + - "PREFILL_NODES=1" + - "MODEL_YAML_KEY=GLM-5.2-FP8" + - "SYNC_BARRIER_TIMEOUT=3600" + decode: + num-worker: 1 + tp: 8 + ep: 1 + dp-attn: false + additional-settings: + - "DECODE_NODES=1" + - "DECODE_MTP_SIZE=0" + + # Decode-heavy: split conversations across two independent TP8 decoders. + - spec-decoding: none + conc-list: [3, 4, 6, 8] + prefill: + num-worker: 1 + tp: 8 + ep: 1 + dp-attn: false + additional-settings: + - "PREFILL_NODES=1" + - "MODEL_YAML_KEY=GLM-5.2-FP8" + - "SYNC_BARRIER_TIMEOUT=3600" + decode: + num-worker: 2 + tp: 8 + ep: 1 + dp-attn: false + additional-settings: + - "DECODE_NODES=2" + - "DECODE_MTP_SIZE=0" + + # Prefill-heavy: two TP8 prefill workers feeding one TP8 decoder. + - spec-decoding: none + conc-list: [2, 3, 4, 6] + prefill: + num-worker: 2 + tp: 8 + ep: 1 + dp-attn: false + additional-settings: + - "PREFILL_NODES=2" + - "MODEL_YAML_KEY=GLM-5.2-FP8" + - "SYNC_BARRIER_TIMEOUT=3600" + decode: + num-worker: 1 + tp: 8 + ep: 1 + dp-attn: false + additional-settings: + - "DECODE_NODES=1" + - "DECODE_MTP_SIZE=0" + minimaxm3-fp8-mi300x-vllm-agentic: image: vllm/vllm-openai-rocm:nightly-04c2a8deac44fdb1ca3e2b5ec3e6bf16f3f6a914 model: MiniMaxAI/MiniMax-M3-MXFP8 diff --git a/perf-changelog.yaml b/perf-changelog.yaml index 892e3ee2b6..d274e1727e 100644 --- a/perf-changelog.yaml +++ b/perf-changelog.yaml @@ -4974,3 +4974,13 @@ - "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-mi325x-sglang-agentic + - glm5.2-fp8-mi325x-sglang-disagg-agentic + description: + - "Add four GLM-5.2 FP8 MI325X SGLang single-node AgentX full-context curves: TP8 low-latency, TP8/EP8 balanced, TP8/EP8/DP8 high-throughput, and TP8 HiCache DRAM offload" + - "Keep every arm at an explicit 1,048,576-token request context and 1,048,576-token BF16 GPU KV pool against the complete unfiltered AgentX corpus, with dense concurrency coverage from 1 through 32" + - "Add TP8 SGLang disaggregated AgentX sweeps for balanced 1P1D, decode-heavy 1P2D, and prefill-heavy 2P1D topologies, with at most four concurrency points per server allocation" + - "Keep EP and DP attention disabled because MoRI all-to-all remains unstable on the MI325X Broadcom bnxt_re fleet" + pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2273 diff --git a/runners/launch_mi325x-amds.sh b/runners/launch_mi325x-amds.sh index a2cad5d496..30747b5e70 100644 --- a/runners/launch_mi325x-amds.sh +++ b/runners/launch_mi325x-amds.sh @@ -1,6 +1,10 @@ #!/usr/bin/env bash set -euo pipefail +if [[ "${IS_MULTINODE:-false}" == "true" ]]; then + exec bash "$(dirname "$0")/launch_mi325x-amds_multinode.sh" +fi + export HF_HUB_CACHE_MOUNT="/raid/hf-hub-cache/" PARTITION="compute" @@ -26,7 +30,36 @@ export PORT=$(( 40000 + (JOB_ID % 10000) )) export XDG_CACHE_HOME="/tmp/xdg-cache-$JOB_ID" export TRITON_CACHE_DIR="/tmp/triton-cache-$JOB_ID" -trap 'rc=$?; scancel "$JOB_ID" 2>/dev/null || true; exit "$rc"' EXIT +cleanup_allocation() { + local rc=$? + + # AgentX restores a multi-GB mmap dataset under XDG_CACHE_HOME. Remove this + # allocation's cache and compiled kernels so repeated sweeps cannot fill the + # worker's local disk. The paths include the resolved Slurm job ID and are + # therefore isolated from other allocations. + srun --jobid="$JOB_ID" --job-name="$RUNNER_NAME" \ + bash -c 'rm -rf -- "$XDG_CACHE_HOME" "$TRITON_CACHE_DIR"' \ + 2>/dev/null || true + scancel "$JOB_ID" 2>/dev/null || true + exit "$rc" +} +trap cleanup_allocation EXIT + +# A terminated enroot step can leave its background SGLang process outside the +# completed Slurm step. Because each allocation owns all eight GPUs on its node, +# remove only stale SGLang/router processes before starting the next server. +# This prevents an old 200-GB/GPU model instance from poisoning later jobs. +srun --jobid="$JOB_ID" --job-name="$RUNNER_NAME" bash -c ' + stale_pattern="[s]glang.launch_server|[s]glang_router.launch_router|[s]glang::" + mapfile -t stale_pids < <(pgrep -f "$stale_pattern" || true) + if [ "${#stale_pids[@]}" -gt 0 ]; then + echo "Cleaning stale SGLang processes: ${stale_pids[*]}" + kill -TERM "${stale_pids[@]}" 2>/dev/null || sudo -n kill -TERM "${stale_pids[@]}" 2>/dev/null || true + sleep 5 + mapfile -t stale_pids < <(pgrep -f "$stale_pattern" || true) + [ "${#stale_pids[@]}" -eq 0 ] || sudo -n kill -KILL "${stale_pids[@]}" 2>/dev/null || true + fi +' # Use flock to serialize concurrent imports to the same squash file srun --jobid="$JOB_ID" --job-name="$RUNNER_NAME" bash -c " diff --git a/runners/launch_mi325x-amds_multinode.sh b/runners/launch_mi325x-amds_multinode.sh new file mode 100755 index 0000000000..53ed489f6c --- /dev/null +++ b/runners/launch_mi325x-amds_multinode.sh @@ -0,0 +1,307 @@ +#!/usr/bin/env bash +set -euo pipefail + +ORIGINAL_GITHUB_WORKSPACE="$GITHUB_WORKSPACE" + +scancel_sync() { + local job_id="$1" + local deadline=$((SECONDS + 600)) + scancel "$job_id" 2>/dev/null || true + while squeue -h -j "$job_id" 2>/dev/null | grep -q .; do + if (( SECONDS >= deadline )); then + echo "WARNING: Slurm job $job_id still exists after 600 seconds" >&2 + return 1 + fi + sleep 10 + done +} + +stage_server_logs() { + [[ -d "${BENCHMARK_LOGS_DIR:-}" ]] || return 0 + local archive + archive=$(mktemp /tmp/inferencex-mi325x-server-logs.XXXXXX.tar.gz) + sudo tar -czf "$archive" \ + --exclude='*/agentic' \ + --exclude='*/agentic/*' \ + --exclude='*/agentic-output' \ + --exclude='*/agentic-output/*' \ + -C "$BENCHMARK_LOGS_DIR" . 2>/dev/null || true + sudo chown "$USER" "$archive" 2>/dev/null || true + if [[ -s "$archive" ]]; then + mv -f "$archive" "$ORIGINAL_GITHUB_WORKSPACE/multinode_server_logs.tar.gz" + else + rm -f "$archive" + fi +} + +cleanup() { + local rc=$? + trap - EXIT + if [[ -n "${JOB_ID:-}" ]] && squeue -h -j "$JOB_ID" 2>/dev/null | grep -q .; then + scancel_sync "$JOB_ID" || true + fi + stage_server_logs || true + if [[ -n "${BENCHMARK_LOGS_DIR:-}" && -d "$BENCHMARK_LOGS_DIR" ]]; then + sudo rm -rf -- "$BENCHMARK_LOGS_DIR" 2>/dev/null || true + fi + if [[ -n "${BENCHMARK_LOGS_PARENT:-}" && -d "$BENCHMARK_LOGS_PARENT" ]]; then + rmdir "$BENCHMARK_LOGS_PARENT" 2>/dev/null || true + fi + if [[ -n "${SHARED_WORKSPACE:-}" && -d "$SHARED_WORKSPACE" ]]; then + sudo rm -rf -- "$SHARED_WORKSPACE" 2>/dev/null || true + fi + if [[ -n "${SHARED_WORKSPACES_PARENT:-}" && -d "$SHARED_WORKSPACES_PARENT" ]]; then + rmdir "$SHARED_WORKSPACES_PARENT" 2>/dev/null || true + fi + exit "$rc" +} +trap cleanup EXIT + +export SLURM_ACCOUNT="$USER" +export SLURM_PARTITION="compute" +export SLURM_JOB_NAME="benchmark-sglang-disagg.job" +runner_shared_root="$GITHUB_WORKSPACE" +for _ in 1 2 3 4 5; do + runner_shared_root=$(dirname "$runner_shared_root") +done +if [[ "$(basename "$runner_shared_root")" != "gharunners" ]]; then + echo "ERROR: cannot derive shared runner root from $GITHUB_WORKSPACE" >&2 + exit 1 +fi +export HF_HUB_CACHE_MOUNT="${MI325X_SHARED_HF_CACHE:-$runner_shared_root/.inferencex-hf-cache}" +if [[ "$MODEL" == "zai-org/GLM-5.2-FP8" ]]; then + # Reuse the node-local cache populated and validated by the successful + # single-node MI325X AgentX sweep. Both nodes used by the 1P1D allocation + # have this exact revision; copying the 753 GB checkpoint into the + # controller-visible cache is unnecessary and unreliable through Xet. + export MODEL_PATH="${MI325X_NODE_LOCAL_HF_CACHE:-/raid/hf-hub-cache}" + export MODEL_NAME="models--zai-org--GLM-5.2-FP8/snapshots/ba978f7d347eaf65d22f1a86833408afdb953541" +else + export MODEL_PATH="$HF_HUB_CACHE_MOUNT" +fi +export MODEL_YAML_KEY="${MODEL_YAML_KEY:-${MODEL##*/}}" +export GPUS_PER_NODE=8 +export IBDEVICES="${IBDEVICES:-bnxt_re0,bnxt_re1,bnxt_re2,bnxt_re3,bnxt_re4,bnxt_re5,bnxt_re7,bnxt_re8}" +export MORI_RDMA_TC="${MORI_RDMA_TC:-104}" + +snapshot_has_required_files() { + local snapshot_dir="$1" + local required_file + for required_file in \ + config.json \ + model.safetensors.index.json \ + tokenizer.json \ + tokenizer_config.json; do + if [[ ! -f "$snapshot_dir/$required_file" ]]; then + echo "Model snapshot is incomplete: missing $snapshot_dir/$required_file" >&2 + return 1 + fi + done +} + +resolve_model_name() { + local hf_dir="models--${MODEL//\//--}" + local snapshot_root="$MODEL_PATH/$hf_dir/snapshots" + local snapshot="" + + if [[ -d "$snapshot_root" ]]; then + snapshot=$(find "$snapshot_root" -mindepth 1 -maxdepth 1 -type d -printf '%f\n' | sort | tail -1) + fi + if [[ -n "$snapshot" ]] && snapshot_has_required_files "$snapshot_root/$snapshot"; then + printf '%s\n' "$hf_dir/snapshots/$snapshot" + return 0 + fi + if [[ -d "$MODEL_PATH/${MODEL##*/}" ]]; then + printf '%s\n' "${MODEL##*/}" + return 0 + fi + return 1 +} + +stage_model_to_shared_cache() { + local hf_dir="models--${MODEL//\//--}" + local lock_file="$MODEL_PATH/.${hf_dir}.download.lock" + local python_path="" + local tool_dir="" + + mkdir -p "$MODEL_PATH" + if [[ ! -w "$MODEL_PATH" ]]; then + echo "ERROR: shared model cache '$MODEL_PATH' is not writable by $USER" >&2 + return 1 + fi + exec 9>"$lock_file" + if ! flock -w 18000 9; then + echo "ERROR: timed out waiting to stage '$MODEL' under $MODEL_PATH" >&2 + return 1 + fi + + # Another topology may have completed the download while this launcher + # waited for the shared-cache lock. + if resolve_model_name >/dev/null; then + flock -u 9 + exec 9>&- + return 0 + fi + + echo "Staging '$MODEL' into shared cache $MODEL_PATH" + if ! python3 -c 'import huggingface_hub' >/dev/null 2>&1; then + tool_dir=$(mktemp -d /tmp/inferencex-hf-client.XXXXXX) + python3 -m pip install --quiet --disable-pip-version-check \ + --target "$tool_dir" 'huggingface_hub>=0.30' + python_path="$tool_dir" + fi + + PYTHONPATH="${python_path}${PYTHONPATH:+:$PYTHONPATH}" \ + python3 - "$MODEL" "$MODEL_PATH" <<'PY' +import os +import sys + +from huggingface_hub import snapshot_download + +snapshot_download( + repo_id=sys.argv[1], + cache_dir=sys.argv[2], + token=os.environ.get("HF_TOKEN"), +) +PY + + if [[ -n "$tool_dir" ]]; then + rm -rf -- "$tool_dir" + fi + if ! resolve_model_name >/dev/null; then + echo "ERROR: download completed but '$MODEL' is unresolved under $MODEL_PATH" >&2 + flock -u 9 + exec 9>&- + return 1 + fi + flock -u 9 + exec 9>&- +} + +if [[ -z "${MODEL_NAME:-}" ]]; then + if ! MODEL_NAME=$(resolve_model_name); then + stage_model_to_shared_cache + MODEL_NAME=$(resolve_model_name) + fi + export MODEL_NAME +fi + +runner_tag=${RUNNER_NAME//[^a-zA-Z0-9_.-]/_} +SHARED_WORKSPACES_PARENT="$runner_shared_root/.inferencex-workspaces" +mkdir -p "$SHARED_WORKSPACES_PARENT" +SHARED_WORKSPACE=$(mktemp -d "$SHARED_WORKSPACES_PARENT/mi325x-${runner_tag}.XXXXXX") +rsync -a \ + --exclude='.git' \ + --exclude='LOGS' \ + --exclude='results' \ + "$ORIGINAL_GITHUB_WORKSPACE/" "$SHARED_WORKSPACE/" +export SHARED_WORKSPACES_PARENT +export SHARED_WORKSPACE +export GITHUB_WORKSPACE="$SHARED_WORKSPACE" + +BENCHMARK_LOGS_PARENT="$runner_shared_root/.inferencex-benchmark-logs" +mkdir -p "$BENCHMARK_LOGS_PARENT" +BENCHMARK_LOGS_DIR=$(mktemp -d "$BENCHMARK_LOGS_PARENT/mi325x-${runner_tag}.XXXXXX") +export BENCHMARK_LOGS_PARENT +export BENCHMARK_LOGS_DIR + +script_name="${EXP_NAME%%_*}_${PRECISION}_mi325x_${FRAMEWORK}.sh" +benchmark_script="benchmarks/multi_node/$script_name" +if [[ ! -f "$benchmark_script" ]]; then + echo "ERROR: $benchmark_script does not exist" >&2 + exit 1 +fi + +echo "Submitting $benchmark_script with model path $MODEL_NAME" +JOB_ID=$(bash "$benchmark_script" | tail -1) +if ! [[ "$JOB_ID" =~ ^[0-9]+$ ]]; then + echo "ERROR: benchmark submission returned invalid job ID '$JOB_ID'" >&2 + exit 1 +fi + +log_file="$BENCHMARK_LOGS_DIR/slurm_job-${JOB_ID}.out" +err_file="$BENCHMARK_LOGS_DIR/slurm_job-${JOB_ID}.err" +while [[ ! -f "$log_file" ]]; do + if ! squeue -h -j "$JOB_ID" 2>/dev/null | grep -q .; then + echo "Slurm job $JOB_ID ended before its output log became visible; waiting for NFS metadata" >&2 + for _ in 1 2 3 4 5 6; do + sync + [[ -f "$log_file" || -f "$err_file" ]] && break + sleep 5 + done + if [[ -f "$log_file" ]]; then + break + fi + if [[ -s "$err_file" ]]; then + echo "===== Slurm stderr: job $JOB_ID =====" >&2 + sudo cat "$err_file" >&2 || true + echo "===== End Slurm stderr: job $JOB_ID =====" >&2 + fi + echo "ERROR: Slurm job $JOB_ID ended before creating its output log" >&2 + scontrol show job "$JOB_ID" 2>/dev/null || true + sacct -X -j "$JOB_ID" \ + --format=JobID,State,ExitCode,DerivedExitCode,NodeList,Comment,SystemComment \ + 2>/dev/null || true + exit 1 + fi + sleep 5 +done + +( + while squeue -h -j "$JOB_ID" 2>/dev/null | grep -q .; do + sleep 10 + done +) & +poll_pid=$! +tail -F -s 2 -n+1 "$log_file" --pid="$poll_pid" 2>/dev/null || true +wait "$poll_pid" + +if [[ -s "$err_file" ]]; then + echo "===== Slurm stderr: job $JOB_ID =====" >&2 + sudo cat "$err_file" >&2 || true + echo "===== End Slurm stderr: job $JOB_ID =====" >&2 +fi + +job_state=$(sacct -n -X -j "$JOB_ID" --format=State --parsable2 2>/dev/null | head -1 | cut -d'|' -f1) +job_exit=$(sacct -n -X -j "$JOB_ID" --format=ExitCode --parsable2 2>/dev/null | head -1 | cut -d'|' -f1) +echo "Slurm job $JOB_ID finished: state=${job_state:-unknown}, exit=${job_exit:-unknown}" + +if [[ "${IS_AGENTIC:-0}" == "1" ]]; then + mapfile -d '' aggregate_files < <( + find "$BENCHMARK_LOGS_DIR/logs" -type f \ + -path '*/agentic-output/*.json' -print0 2>/dev/null + ) + for result_file in "${aggregate_files[@]}"; do + staged_result=$(mktemp /tmp/inferencex-mi325x-result.XXXXXX.json) + sudo cp -f "$result_file" "$staged_result" + sudo chown "$USER" "$staged_result" + mv -f "$staged_result" "$ORIGINAL_GITHUB_WORKSPACE/$(basename "$result_file")" + done + + raw_dir=$(find "$BENCHMARK_LOGS_DIR/logs" -type d -path '*/agentic' -print -quit 2>/dev/null || true) + if [[ -n "$raw_dir" ]]; then + staged_raw=$(mktemp -d /tmp/inferencex-mi325x-agentic.XXXXXX) + sudo cp -a "$raw_dir/." "$staged_raw/" + sudo chown -R "$USER" "$staged_raw" + mkdir -p "$ORIGINAL_GITHUB_WORKSPACE/LOGS/agentic" + cp -a "$staged_raw/." "$ORIGINAL_GITHUB_WORKSPACE/LOGS/agentic/" + rm -rf -- "$staged_raw" + fi +else + while IFS= read -r -d '' result_file; do + output="$ORIGINAL_GITHUB_WORKSPACE/${RESULT_FILENAME}_$(basename "$result_file")" + staged_result=$(mktemp /tmp/inferencex-mi325x-result.XXXXXX.json) + sudo cp -f "$result_file" "$staged_result" + sudo chown "$USER" "$staged_result" + mv -f "$staged_result" "$output" + done < <(find "$BENCHMARK_LOGS_DIR/logs" -type f -name '*.json' -print0 2>/dev/null) +fi + +stage_server_logs +scancel_sync "$JOB_ID" || true +JOB_ID="" + +if [[ "${job_state:-}" != COMPLETED* || "${job_exit:-1:0}" != "0:0" ]]; then + echo "ERROR: Slurm benchmark failed: state=${job_state:-unknown}, exit=${job_exit:-unknown}" >&2 + exit 1 +fi