From c1e0cce7036833c9464a9ae1ac6c05523b868dc2 Mon Sep 17 00:00:00 2001 From: Cam Quilici Date: Mon, 10 Aug 2026 20:47:51 -0500 Subject: [PATCH 01/18] feat(amd): port MI355X DSV4 AgentX to srt-slurm --- .../disagg-1p1d-dep8-mtp-hicache-c128.yaml | 148 ++++++++++++++++++ .../disagg-1p1d-dep8-mtp-hicache-c96.yaml | 148 ++++++++++++++++++ .../agentic/disagg-1p1d-tp8-mtp-hicache.yaml | 142 +++++++++++++++++ configs/amd-master.yaml | 28 +++- runners/launch_mi355x-amds-srt.sh | 130 +++++++++++---- 5 files changed, 561 insertions(+), 35 deletions(-) create mode 100644 benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c128.yaml create mode 100644 benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c96.yaml create mode 100644 benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-tp8-mtp-hicache.yaml diff --git a/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c128.yaml b/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c128.yaml new file mode 100644 index 000000000..ef864c559 --- /dev/null +++ b/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c128.yaml @@ -0,0 +1,148 @@ +name: "mi355x-dsv4-agentic-disagg-1p1d-dep8-mtp-hicache-c128" + +model: + path: "hf:deepseek-ai/DeepSeek-V4-Pro" + container: "sglang-rocm-dsv4-mi355x" + precision: "fp4" + +identity: + model: + repo: "deepseek-ai/DeepSeek-V4-Pro" + container: + image: "lmsysorg/sglang-rocm:v0.5.15.post1-rocm720-mi35x-20260719" + frameworks: + sglang: "0.5.15.post1" + sglang-router: "0.3.2" + amd-mori: "bundled" + +slurm: + time_limit: "04:00:00" + +resources: + gpu_type: "mi355x" + gpus_per_node: 8 + prefill_nodes: 1 + decode_nodes: 1 + prefill_workers: 1 + decode_workers: 1 + gpus_per_prefill: 8 + gpus_per_decode: 8 + +frontend: + type: sglang + enable_multiple_frontends: false + args: + policy: consistent_hashing + dp-aware: true + cache-threshold: 0.3 + balance-abs-threshold: 2 + balance-rel-threshold: 1.1 + disable-circuit-breaker: true + health-failure-threshold: 100 + health-check-timeout-secs: 600 + health-check-interval-secs: 30 + +backend: + type: sglang + prefill_environment: &common_environment + HF_HOME: "/hf_hub_cache" + HF_HUB_CACHE: "/hf_hub_cache/hub" + HUGGINGFACE_HUB_CACHE: "/hf_hub_cache/hub" + PYTHONUNBUFFERED: "1" + IBDEVICES: "rdma0,rdma1,rdma2,rdma3,rdma4,rdma5,rdma6,rdma7" + MORI_RDMA_TC: "104" + MORI_IO_SQ_BACKOFF_TIMEOUT_US: "500000" + MORI_IO_QP_MAX_SEND_WR: "32767" + SGLANG_MORI_COMBINE_DTYPE: "auto" + SGLANG_MORI_NUM_MAX_DISPATCH_TOKENS_PER_RANK: "16384" + SGLANG_DISAGGREGATION_BOOTSTRAP_TIMEOUT: "1800" + SGLANG_DISAGGREGATION_WAITING_TIMEOUT: "1800" + decode_environment: + <<: *common_environment + SGLANG_MORI_NUM_MAX_DISPATCH_TOKENS_PER_RANK: "64" + SGLANG_SIMULATE_ACC_LEN: "2.49" + SGLANG_SIMULATE_ACC_METHOD: "match-expected" + SGLANG_SIMULATE_ACC_TOKEN_MODE: "real-draft-token" + sglang_config: + prefill: &common_server + served-model-name: "deepseek-ai/DeepSeek-V4-Pro" + trust-remote-code: true + tp-size: 8 + ep-size: 8 + dp-size: 8 + watchdog-timeout: 3600 + load-balance-method: round_robin + kv-cache-dtype: fp8_e4m3 + attention-backend: dsv4 + page-size: 256 + swa-full-tokens-ratio: 0.1 + disable-shared-experts-fusion: true + tool-call-parser: deepseekv4 + reasoning-parser: deepseek-v4 + disaggregation-transfer-backend: mori + log-level: info + log-level-http: error + ep-dispatch-algorithm: fake + moe-a2a-backend: mori + deepep-mode: normal + speculative-algorithm: EAGLE + speculative-eagle-topk: 1 + speculative-num-steps: 3 + speculative-num-draft-tokens: 4 + enable-dp-attention: true + enable-prefill-delayer: true + context-length: 1048576 + enable-metrics: true + enable-cache-report: true + mem-fraction-static: 0.72 + max-running-requests: 128 + chunked-prefill-size: 65536 + disable-cuda-graph: true + enable-hierarchical-cache: true + hicache-ratio: 4 + hicache-io-backend: direct + hicache-mem-layout: page_first + hicache-write-policy: write_through + hicache-storage-prefetch-policy: best_effort + decode: + <<: *common_server + mem-fraction-static: 0.85 + chunked-prefill-size: null + disable-cuda-graph: false + enable-hierarchical-cache: false + hicache-ratio: null + hicache-io-backend: null + hicache-mem-layout: null + hicache-write-policy: null + hicache-storage-prefetch-policy: null + prefill-round-robin-balance: true + cuda-graph-bs: [1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31, 32, 33, 34, 35, 36, 37, 38, 39, 40, 41, 42, 43, 44, 45, 46, 47, 48, 49, 50, 51, 52, 53, 54, 55, 56, 57, 58, 59, 60, 61, 62, 63, 64, 65, 66, 67, 68, 69, 70, 71, 72, 73, 74, 75, 76, 77, 78, 79, 80, 81, 82, 83, 84, 85, 86, 87, 88, 89, 90, 91, 92, 93, 94, 95, 96, 97, 98, 99, 100, 101, 102, 103, 104, 105, 106, 107, 108, 109, 110, 111, 112, 113, 114, 115, 116, 117, 118, 119, 120, 121, 122, 123, 124, 125, 126, 127, 128] + +sbatch_directives: + mem: "0" + +srun_options: + mem: "0" + container-writable: "" + container-remap-root: "" + +health_check: + max_attempts: 720 + interval_seconds: 5 + +benchmark: + type: custom + command: | + set -euo pipefail + mkdir -p /infmax-workspace/LOGS/agentic + export RESULT_DIR=/infmax-workspace/LOGS/agentic + export AGENTIC_OUTPUT_DIR=/infmax-workspace + exec bash /infmax-workspace/benchmarks/multi_node/agentic_srt.sh + env: + INFMAX_CONTAINER_WORKSPACE: /infmax-workspace + PORT: "8000" + IS_MULTINODE: "true" + AIPERF_REQUIRED_SERVER_METRIC_PREFIX: "sglang:" + AIPERF_DATASET_MMAP_CACHE_DIR: "/aiperf_mmap_cache" + HF_HUB_CACHE: "/hf_hub_cache/hub" + WEKA_LOADER_OVERRIDE: "semianalysis_cc_traces_weka_062126" diff --git a/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c96.yaml b/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c96.yaml new file mode 100644 index 000000000..8bc6785c3 --- /dev/null +++ b/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c96.yaml @@ -0,0 +1,148 @@ +name: "mi355x-dsv4-agentic-disagg-1p1d-dep8-mtp-hicache-c96" + +model: + path: "hf:deepseek-ai/DeepSeek-V4-Pro" + container: "sglang-rocm-dsv4-mi355x" + precision: "fp4" + +identity: + model: + repo: "deepseek-ai/DeepSeek-V4-Pro" + container: + image: "lmsysorg/sglang-rocm:v0.5.15.post1-rocm720-mi35x-20260719" + frameworks: + sglang: "0.5.15.post1" + sglang-router: "0.3.2" + amd-mori: "bundled" + +slurm: + time_limit: "04:00:00" + +resources: + gpu_type: "mi355x" + gpus_per_node: 8 + prefill_nodes: 1 + decode_nodes: 1 + prefill_workers: 1 + decode_workers: 1 + gpus_per_prefill: 8 + gpus_per_decode: 8 + +frontend: + type: sglang + enable_multiple_frontends: false + args: + policy: consistent_hashing + dp-aware: true + cache-threshold: 0.3 + balance-abs-threshold: 2 + balance-rel-threshold: 1.1 + disable-circuit-breaker: true + health-failure-threshold: 100 + health-check-timeout-secs: 600 + health-check-interval-secs: 30 + +backend: + type: sglang + prefill_environment: &common_environment + HF_HOME: "/hf_hub_cache" + HF_HUB_CACHE: "/hf_hub_cache/hub" + HUGGINGFACE_HUB_CACHE: "/hf_hub_cache/hub" + PYTHONUNBUFFERED: "1" + IBDEVICES: "rdma0,rdma1,rdma2,rdma3,rdma4,rdma5,rdma6,rdma7" + MORI_RDMA_TC: "104" + MORI_IO_SQ_BACKOFF_TIMEOUT_US: "500000" + MORI_IO_QP_MAX_SEND_WR: "32767" + SGLANG_MORI_COMBINE_DTYPE: "auto" + SGLANG_MORI_NUM_MAX_DISPATCH_TOKENS_PER_RANK: "16384" + SGLANG_DISAGGREGATION_BOOTSTRAP_TIMEOUT: "1800" + SGLANG_DISAGGREGATION_WAITING_TIMEOUT: "1800" + decode_environment: + <<: *common_environment + SGLANG_MORI_NUM_MAX_DISPATCH_TOKENS_PER_RANK: "48" + SGLANG_SIMULATE_ACC_LEN: "2.49" + SGLANG_SIMULATE_ACC_METHOD: "match-expected" + SGLANG_SIMULATE_ACC_TOKEN_MODE: "real-draft-token" + sglang_config: + prefill: &common_server + served-model-name: "deepseek-ai/DeepSeek-V4-Pro" + trust-remote-code: true + tp-size: 8 + ep-size: 8 + dp-size: 8 + watchdog-timeout: 3600 + load-balance-method: round_robin + kv-cache-dtype: fp8_e4m3 + attention-backend: dsv4 + page-size: 256 + swa-full-tokens-ratio: 0.1 + disable-shared-experts-fusion: true + tool-call-parser: deepseekv4 + reasoning-parser: deepseek-v4 + disaggregation-transfer-backend: mori + log-level: info + log-level-http: error + ep-dispatch-algorithm: fake + moe-a2a-backend: mori + deepep-mode: normal + speculative-algorithm: EAGLE + speculative-eagle-topk: 1 + speculative-num-steps: 3 + speculative-num-draft-tokens: 4 + enable-dp-attention: true + enable-prefill-delayer: true + context-length: 1048576 + enable-metrics: true + enable-cache-report: true + mem-fraction-static: 0.72 + max-running-requests: 96 + chunked-prefill-size: 65536 + disable-cuda-graph: true + enable-hierarchical-cache: true + hicache-ratio: 4 + hicache-io-backend: direct + hicache-mem-layout: page_first + hicache-write-policy: write_through + hicache-storage-prefetch-policy: best_effort + decode: + <<: *common_server + mem-fraction-static: 0.85 + chunked-prefill-size: null + disable-cuda-graph: false + enable-hierarchical-cache: false + hicache-ratio: null + hicache-io-backend: null + hicache-mem-layout: null + hicache-write-policy: null + hicache-storage-prefetch-policy: null + prefill-round-robin-balance: true + cuda-graph-bs: [1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31, 32, 33, 34, 35, 36, 37, 38, 39, 40, 41, 42, 43, 44, 45, 46, 47, 48, 49, 50, 51, 52, 53, 54, 55, 56, 57, 58, 59, 60, 61, 62, 63, 64, 65, 66, 67, 68, 69, 70, 71, 72, 73, 74, 75, 76, 77, 78, 79, 80, 81, 82, 83, 84, 85, 86, 87, 88, 89, 90, 91, 92, 93, 94, 95, 96, 97, 98, 99, 100, 101, 102, 103, 104, 105, 106, 107, 108, 109, 110, 111, 112, 113, 114, 115, 116, 117, 118, 119, 120, 121, 122, 123, 124, 125, 126, 127, 128] + +sbatch_directives: + mem: "0" + +srun_options: + mem: "0" + container-writable: "" + container-remap-root: "" + +health_check: + max_attempts: 720 + interval_seconds: 5 + +benchmark: + type: custom + command: | + set -euo pipefail + mkdir -p /infmax-workspace/LOGS/agentic + export RESULT_DIR=/infmax-workspace/LOGS/agentic + export AGENTIC_OUTPUT_DIR=/infmax-workspace + exec bash /infmax-workspace/benchmarks/multi_node/agentic_srt.sh + env: + INFMAX_CONTAINER_WORKSPACE: /infmax-workspace + PORT: "8000" + IS_MULTINODE: "true" + AIPERF_REQUIRED_SERVER_METRIC_PREFIX: "sglang:" + AIPERF_DATASET_MMAP_CACHE_DIR: "/aiperf_mmap_cache" + HF_HUB_CACHE: "/hf_hub_cache/hub" + WEKA_LOADER_OVERRIDE: "semianalysis_cc_traces_weka_062126" diff --git a/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-tp8-mtp-hicache.yaml b/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-tp8-mtp-hicache.yaml new file mode 100644 index 000000000..815ec2868 --- /dev/null +++ b/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-tp8-mtp-hicache.yaml @@ -0,0 +1,142 @@ +name: "mi355x-dsv4-agentic-disagg-1p1d-tp8-mtp-hicache" + +model: + path: "hf:deepseek-ai/DeepSeek-V4-Pro" + container: "sglang-rocm-dsv4-mi355x" + precision: "fp4" + +identity: + model: + repo: "deepseek-ai/DeepSeek-V4-Pro" + container: + image: "lmsysorg/sglang-rocm:v0.5.15.post1-rocm720-mi35x-20260719" + frameworks: + sglang: "0.5.15.post1" + sglang-router: "0.3.2" + amd-mori: "bundled" + +slurm: + time_limit: "04:00:00" + +resources: + gpu_type: "mi355x" + gpus_per_node: 8 + prefill_nodes: 1 + decode_nodes: 1 + prefill_workers: 1 + decode_workers: 1 + gpus_per_prefill: 8 + gpus_per_decode: 8 + +frontend: + type: sglang + enable_multiple_frontends: false + args: + policy: consistent_hashing + dp-aware: true + cache-threshold: 0.3 + balance-abs-threshold: 2 + balance-rel-threshold: 1.1 + disable-circuit-breaker: true + health-failure-threshold: 100 + health-check-timeout-secs: 600 + health-check-interval-secs: 30 + +backend: + type: sglang + prefill_environment: &common_environment + HF_HOME: "/hf_hub_cache" + HF_HUB_CACHE: "/hf_hub_cache/hub" + HUGGINGFACE_HUB_CACHE: "/hf_hub_cache/hub" + PYTHONUNBUFFERED: "1" + IBDEVICES: "rdma0,rdma1,rdma2,rdma3,rdma4,rdma5,rdma6,rdma7" + MORI_RDMA_TC: "104" + MORI_IO_SQ_BACKOFF_TIMEOUT_US: "500000" + MORI_IO_QP_MAX_SEND_WR: "32767" + SGLANG_MORI_COMBINE_DTYPE: "auto" + SGLANG_MORI_NUM_MAX_DISPATCH_TOKENS_PER_RANK: "16384" + SGLANG_DISAGGREGATION_BOOTSTRAP_TIMEOUT: "1800" + SGLANG_DISAGGREGATION_WAITING_TIMEOUT: "1800" + decode_environment: + <<: *common_environment + SGLANG_MORI_NUM_MAX_DISPATCH_TOKENS_PER_RANK: "1024" + SGLANG_SIMULATE_ACC_LEN: "2.49" + SGLANG_SIMULATE_ACC_METHOD: "match-expected" + SGLANG_SIMULATE_ACC_TOKEN_MODE: "real-draft-token" + sglang_config: + prefill: &common_server + served-model-name: "deepseek-ai/DeepSeek-V4-Pro" + trust-remote-code: true + tp-size: 8 + watchdog-timeout: 3600 + load-balance-method: round_robin + kv-cache-dtype: fp8_e4m3 + attention-backend: dsv4 + page-size: 256 + swa-full-tokens-ratio: 0.1 + disable-shared-experts-fusion: true + tool-call-parser: deepseekv4 + reasoning-parser: deepseek-v4 + disaggregation-transfer-backend: mori + log-level: info + log-level-http: error + speculative-algorithm: EAGLE + speculative-eagle-topk: 1 + speculative-num-steps: 3 + speculative-num-draft-tokens: 4 + context-length: 1048576 + enable-metrics: true + enable-cache-report: true + mem-fraction-static: 0.72 + max-running-requests: 64 + chunked-prefill-size: 8192 + disable-cuda-graph: true + enable-hierarchical-cache: true + hicache-ratio: 4 + hicache-io-backend: direct + hicache-mem-layout: page_first + hicache-write-policy: write_through + hicache-storage-prefetch-policy: best_effort + decode: + <<: *common_server + mem-fraction-static: 0.85 + max-running-requests: 128 + chunked-prefill-size: null + disable-cuda-graph: false + enable-hierarchical-cache: false + hicache-ratio: null + hicache-io-backend: null + hicache-mem-layout: null + hicache-write-policy: null + hicache-storage-prefetch-policy: null + prefill-round-robin-balance: true + cuda-graph-bs: [1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31, 32, 33, 34, 35, 36, 37, 38, 39, 40, 41, 42, 43, 44, 45, 46, 47, 48, 49, 50, 51, 52, 53, 54, 55, 56, 57, 58, 59, 60, 61, 62, 63, 64, 65, 66, 67, 68, 69, 70, 71, 72, 73, 74, 75, 76, 77, 78, 79, 80, 81, 82, 83, 84, 85, 86, 87, 88, 89, 90, 91, 92, 93, 94, 95, 96, 97, 98, 99, 100, 101, 102, 103, 104, 105, 106, 107, 108, 109, 110, 111, 112, 113, 114, 115, 116, 117, 118, 119, 120, 121, 122, 123, 124, 125, 126, 127, 128] + +sbatch_directives: + mem: "0" + +srun_options: + mem: "0" + container-writable: "" + container-remap-root: "" + +health_check: + max_attempts: 720 + interval_seconds: 5 + +benchmark: + type: custom + command: | + set -euo pipefail + mkdir -p /infmax-workspace/LOGS/agentic + export RESULT_DIR=/infmax-workspace/LOGS/agentic + export AGENTIC_OUTPUT_DIR=/infmax-workspace + exec bash /infmax-workspace/benchmarks/multi_node/agentic_srt.sh + env: + INFMAX_CONTAINER_WORKSPACE: /infmax-workspace + PORT: "8000" + IS_MULTINODE: "true" + AIPERF_REQUIRED_SERVER_METRIC_PREFIX: "sglang:" + AIPERF_DATASET_MMAP_CACHE_DIR: "/aiperf_mmap_cache" + HF_HUB_CACHE: "/hf_hub_cache/hub" + WEKA_LOADER_OVERRIDE: "semianalysis_cc_traces_weka_062126" diff --git a/configs/amd-master.yaml b/configs/amd-master.yaml index 39b9e2d5d..543960d0e 100644 --- a/configs/amd-master.yaml +++ b/configs/amd-master.yaml @@ -1696,18 +1696,16 @@ dsv4-fp4-mi355x-sglang-disagg-agentic-hicache-mtp: ep: 1 dp-attn: false additional-settings: - - "PREFILL_NODES=1" - - "CLIENT_IMAGE=lmsysorg/sglang-rocm:v0.5.15.post1-rocm720-mi35x-20260719" + - "CONFIG_FILE=recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-tp8-mtp-hicache.yaml" decode: num-worker: 1 tp: 8 ep: 1 dp-attn: false additional-settings: - - "DECODE_NODES=1" - "DECODE_MTP_SIZE=3" - spec-decoding: "mtp" - conc-list: [ 96, 128 ] + conc-list: [ 96 ] kv-offloading: dram kv-offload-backend: { name: hicache } prefill: @@ -1716,15 +1714,31 @@ dsv4-fp4-mi355x-sglang-disagg-agentic-hicache-mtp: ep: 8 dp-attn: true additional-settings: - - "PREFILL_NODES=1" - - "CLIENT_IMAGE=lmsysorg/sglang-rocm:v0.5.15.post1-rocm720-mi35x-20260719" + - "CONFIG_FILE=recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c96.yaml" + decode: + num-worker: 1 + tp: 8 + ep: 8 + dp-attn: true + additional-settings: + - "DECODE_MTP_SIZE=3" + - spec-decoding: "mtp" + conc-list: [ 128 ] + kv-offloading: dram + kv-offload-backend: { name: hicache } + prefill: + num-worker: 1 + tp: 8 + ep: 8 + dp-attn: true + additional-settings: + - "CONFIG_FILE=recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c128.yaml" decode: num-worker: 1 tp: 8 ep: 8 dp-attn: true additional-settings: - - "DECODE_NODES=1" - "DECODE_MTP_SIZE=3" # DeepSeek-V4-Pro FP8 single-node on MI325X (gfx942) via vLLM. diff --git a/runners/launch_mi355x-amds-srt.sh b/runners/launch_mi355x-amds-srt.sh index a5e94c6a0..67388fb35 100755 --- a/runners/launch_mi355x-amds-srt.sh +++ b/runners/launch_mi355x-amds-srt.sh @@ -7,17 +7,23 @@ set -euo pipefail SRT_SLURM_REPOSITORY="https://github.com/SemiAnalysisAI/srt-slurm.git" SRT_SLURM_COMMIT="31e72da43ed21fe941c039be51b2cad1a3cf428a" SLURM_PARTITION="compute" -SGLANG_IMAGE="lmsysorg/sglang-rocm:v0.5.17-rocm720-mi35x-20260809" SHARED_BASE="/it-share/gharunners2/srt-slurm" -SHARED_IMAGE="${SHARED_BASE}/containers/sglang-rocm-v0.5.17-mi35x-20260809.sqsh" SHARED_HF_CACHE="/it-share/hf-hub-cache" +SHARED_AIPERF_CACHE="/it-share/aiperf-cache" SHARED_RESULTS="${SHARED_BASE}/results" : "${GITHUB_WORKSPACE:?GITHUB_WORKSPACE must be set by Actions}" : "${RESULT_FILENAME:?RESULT_FILENAME must be set by the benchmark workflow}" : "${CONFIG_FILE:?CONFIG_FILE must name an srt-slurm recipe}" +: "${IMAGE:?IMAGE must identify the SGLang container image}" : "${MODEL:?MODEL must identify the Hugging Face model}" +SGLANG_IMAGE="$IMAGE" +IMAGE_KEY="${SGLANG_IMAGE//\//_}" +IMAGE_KEY="${IMAGE_KEY//:/_}" +SHARED_IMAGE="${SHARED_BASE}/containers/${IMAGE_KEY}.sqsh" +LOCAL_IMAGE="/var/lib/squash/${IMAGE_KEY}.sqsh" + CONFIG_PATH="${CONFIG_FILE%%:*}" LOCAL_RECIPE="${GITHUB_WORKSPACE}/benchmarks/multi_node/srt-slurm-recipes/${CONFIG_PATH#recipes/}" CLUSTER_PROFILE="${GITHUB_WORKSPACE}/benchmarks/multi_node/srt-slurm-recipes/cluster-configs/mi355x-amds.yaml" @@ -27,7 +33,7 @@ CLUSTER_PROFILE="${GITHUB_WORKSPACE}/benchmarks/multi_node/srt-slurm-recipes/clu RUN_KEY="${GITHUB_RUN_ID:-manual}-${GITHUB_RUN_ATTEMPT:-0}-${RUNNER_NAME:-runner}" WORK_DIR="${GITHUB_WORKSPACE}/.srt-slurm-${RUN_KEY}" SRT_REPO_DIR="${WORK_DIR}/srt-slurm" -mkdir -p "$WORK_DIR" "$SHARED_RESULTS" +mkdir -p "$WORK_DIR" "$SHARED_RESULTS" "$SHARED_AIPERF_CACHE" # Materialize one immutable, shared squashfs and the small public validation # model. This job exits normally and never cancels or preempts another job. @@ -49,7 +55,7 @@ flock -w 2400 9 if ! unsquashfs -s "$SHARED_IMAGE" >/dev/null 2>&1; then tmp="${SHARED_IMAGE}.tmp.\${SLURM_JOB_ID}" rm -f "\$tmp" - local_image="/var/lib/squash/lmsysorg_sglang-rocm_v0.5.17-rocm720-mi35x-20260809.sqsh" + local_image="${LOCAL_IMAGE}" if unsquashfs -s "\$local_image" >/dev/null 2>&1; then cp --sparse=always "\$local_image" "\$tmp" else @@ -81,22 +87,76 @@ ACTUAL_SRT_COMMIT=$(git -C "$SRT_REPO_DIR" rev-parse HEAD) mkdir -p "${SRT_REPO_DIR}/$(dirname "$CONFIG_PATH")" cp "$LOCAL_RECIPE" "${SRT_REPO_DIR}/${CONFIG_PATH}" cp "$CLUSTER_PROFILE" "${WORK_DIR}/srtslurm.yaml" -python3 - "${WORK_DIR}/srtslurm.yaml" "$GITHUB_WORKSPACE" "$SHARED_RESULTS" <<'PY' +python3 - "${WORK_DIR}/srtslurm.yaml" "${SRT_REPO_DIR}/${CONFIG_PATH}" \ + "$GITHUB_WORKSPACE" "$SHARED_RESULTS" "$SHARED_AIPERF_CACHE" "$SHARED_IMAGE" <<'PY' +import os import sys from pathlib import Path -path = Path(sys.argv[1]) -workspace, results = sys.argv[2:] +import yaml + +profile_path = Path(sys.argv[1]) +recipe_path = Path(sys.argv[2]) +workspace, results, aiperf_cache, image_path = sys.argv[3:] needle = " /it-share/hf-hub-cache: /hf_hub_cache\n" -text = path.read_text() +text = profile_path.read_text() if text.count(needle) != 1: raise SystemExit("expected exactly one Hugging Face cache mount") -path.write_text( +profile_path.write_text( text.replace( needle, - needle + f" {workspace}: /infmax-workspace\n {results}: /results\n", + needle + + f" {aiperf_cache}: /aiperf_mmap_cache\n" + + f" {workspace}: /infmax-workspace\n" + + f" {results}: /results\n", ) ) + +recipe = yaml.safe_load(recipe_path.read_text()) +container_alias = recipe["model"]["container"] +profile = yaml.safe_load(profile_path.read_text()) +profile.setdefault("containers", {})[container_alias] = image_path +profile_path.write_text(yaml.safe_dump(profile, sort_keys=False)) + +benchmark_env = recipe.setdefault("benchmark", {}).setdefault("env", {}) +forwarded = ( + "AIPERF_EXPERIMENTAL_FAST", + "CONC", + "CONC_LIST", + "DURATION", + "EVAL_CONC", + "EVAL_LIMIT", + "EVAL_ONLY", + "FRAMEWORK", + "IS_AGENTIC", + "KV_OFFLOADING", + "MAX_MODEL_LEN", + "MODEL", + "MODEL_PREFIX", + "PRECISION", + "RESULT_FILENAME", + "RUN_EVAL", + "RUNNER_TYPE", + "SPEC_DECODING", + "TOTAL_CPU_DRAM_GB", +) +for key in forwarded: + value = os.environ.get(key) + if value: + benchmark_env[key] = value + +if os.environ.get("EVAL_ONLY", "false").lower() == "true" or os.environ.get("RUN_EVAL", "false").lower() == "true": + decode_env = recipe.get("backend", {}).get("decode_environment", {}) + for key in ( + "SGLANG_SIMULATE_ACC_LEN", + "SGLANG_SIMULATE_ACC_METHOD", + "SGLANG_SIMULATE_ACC_TOKEN_MODE", + ): + decode_env.pop(key, None) + server_config = recipe.get("backend", {}).get("sglang_config", {}) + for mode in ("prefill", "decode"): + server_config.get(mode, {}).pop("ep-dispatch-algorithm", None) +recipe_path.write_text(yaml.safe_dump(recipe, sort_keys=False)) PY export PATH="$HOME/.local/bin:$PATH" @@ -135,11 +195,14 @@ read -r JOB_STATE JOB_EXIT JOB_NODELIST < <( echo "srt-slurm job ${JOB_ID}: state=${JOB_STATE} exit=${JOB_EXIT} nodes=${JOB_NODELIST}" RESULT_DIR="${SHARED_RESULTS}/${JOB_ID}" +OUTPUT_LOG_DIR="${SHARED_BASE}/outputs/${JOB_ID}/logs" mkdir -p "$GITHUB_WORKSPACE/LOGS" -if [[ -f "$RESULT_DIR/runtime-logs.tar.gz" ]]; then - cp "$RESULT_DIR/runtime-logs.tar.gz" "$GITHUB_WORKSPACE/multinode_server_logs.tar.gz" +if [[ -d "$OUTPUT_LOG_DIR" ]]; then + tar -C "$OUTPUT_LOG_DIR" -czf "$GITHUB_WORKSPACE/multinode_server_logs.tar.gz" . +fi +if [[ -d "$RESULT_DIR" ]]; then + cp -R "$RESULT_DIR/." "$GITHUB_WORKSPACE/LOGS/" fi -cp -R "$RESULT_DIR/." "$GITHUB_WORKSPACE/LOGS/" if [[ "${DISAGG:-false}" == "true" ]]; then PREFILL_GPUS=$((PREFILL_NUM_WORKERS * PREFILL_TP)) @@ -149,21 +212,32 @@ else TOTAL_GPUS=$((PREFILL_NUM_WORKERS * PREFILL_TP * ${PREFILL_PP_SIZE:-1} * ${PREFILL_PCP_SIZE:-1})) fi -shopt -s nullglob -RESULTS=("$RESULT_DIR"/fixed-seq/*.json) -shopt -u nullglob -[[ ${#RESULTS[@]} -gt 0 ]] || { echo "No fixed-sequence results found in $RESULT_DIR" >&2; exit 1; } -for result in "${RESULTS[@]}"; do - concurrency=$(basename "$result" | sed -n 's/.*-c\([0-9][0-9]*\)\.json/\1/p') - [[ -n "$concurrency" ]] || { echo "Cannot parse concurrency from $result" >&2; exit 1; } - if [[ "${DISAGG:-false}" == "true" ]]; then - output="${GITHUB_WORKSPACE}/${RESULT_FILENAME}_srt-${JOB_ID}_conc${concurrency}_gpus_${TOTAL_GPUS}_ctx_${PREFILL_GPUS}_gen_${DECODE_GPUS}.json" - else - output="${GITHUB_WORKSPACE}/${RESULT_FILENAME}_srt-${JOB_ID}_conc${concurrency}_gpus_${TOTAL_GPUS}.json" - fi - cp "$result" "$output" - echo "Collected $output" -done +if [[ "${IS_AGENTIC:-0}" == "1" ]]; then + shopt -s nullglob + RESULTS=("$GITHUB_WORKSPACE/${RESULT_FILENAME}"_conc*.json) + shopt -u nullglob + [[ ${#RESULTS[@]} -gt 0 ]] || { + echo "No AgentX aggregate results found for ${RESULT_FILENAME}" >&2 + exit 1 + } + printf 'Collected %s\n' "${RESULTS[@]}" +else + shopt -s nullglob + RESULTS=("$RESULT_DIR"/fixed-seq/*.json) + shopt -u nullglob + [[ ${#RESULTS[@]} -gt 0 ]] || { echo "No fixed-sequence results found in $RESULT_DIR" >&2; exit 1; } + for result in "${RESULTS[@]}"; do + concurrency=$(basename "$result" | sed -n 's/.*-c\([0-9][0-9]*\)\.json/\1/p') + [[ -n "$concurrency" ]] || { echo "Cannot parse concurrency from $result" >&2; exit 1; } + if [[ "${DISAGG:-false}" == "true" ]]; then + output="${GITHUB_WORKSPACE}/${RESULT_FILENAME}_srt-${JOB_ID}_conc${concurrency}_gpus_${TOTAL_GPUS}_ctx_${PREFILL_GPUS}_gen_${DECODE_GPUS}.json" + else + output="${GITHUB_WORKSPACE}/${RESULT_FILENAME}_srt-${JOB_ID}_conc${concurrency}_gpus_${TOTAL_GPUS}.json" + fi + cp "$result" "$output" + echo "Collected $output" + done +fi if [[ "$JOB_STATE" != COMPLETED || "$JOB_EXIT" != 0:0 ]]; then echo "srt-slurm validation failed: ${JOB_STATE} (${JOB_EXIT})" >&2 From 2c094312e50f1ba8bb8a214fb0ce6ad247a4c700 Mon Sep 17 00:00:00 2001 From: Cam Quilici Date: Mon, 10 Aug 2026 20:50:07 -0500 Subject: [PATCH 02/18] docs: register MI355X AgentX srt-slurm port --- perf-changelog.yaml | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/perf-changelog.yaml b/perf-changelog.yaml index efd8c51dd..598f9346e 100644 --- a/perf-changelog.yaml +++ b/perf-changelog.yaml @@ -5788,3 +5788,11 @@ - "models.yaml: add mtp_flags (--speculative-algorithm EAGLE --speculative-eagle-topk 1) to DeepSeek-V4-Pro-AgentX; num-steps/num-draft-tokens are derived from DECODE_MTP_SIZE in server_sglang.sh. DeepSeek-V4's MTP head is a native NextN/EAGLE draft shipped with the model, so no --speculative-draft-model-path is needed." - "Two search-space arms: TP8/EP1 no-DP at conc-list [2,4,8,16,32], and TP8/EP8/DPA (ep=8, dp-attn=true) at conc-list [64,96,128], both on image lmsysorg/sglang-rocm:v0.5.15.post1-rocm720-mi35x-20260719." pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2309 + +- config-keys: + - dsv4-fp4-mi355x-sglang-disagg-agentic-hicache-mtp + description: + - "Port the merged MI355X DeepSeek V4 AgentX MTP 1P1D search space from the legacy AMD multinode launcher to srt-slurm." + - "Use srt-slurm's native SGLang Router frontend and AMD MoRI transport without Dynamo, NATS, etcd, or NIXL." + - "Preserve TP8/EP1 c1,c2,c4,c8,c16,c32 and TP8/EP8/DPA c96,c128, EAGLE MTP, the golden simulated average length, HiCache settings, and logical prefill/decode server-metrics collection through the custom AgentX benchmark." + pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2555 From a3c0ad18e90d530feeef9a66e5945b1c3e973530 Mon Sep 17 00:00:00 2001 From: Cam Quilici Date: Mon, 10 Aug 2026 21:05:37 -0500 Subject: [PATCH 03/18] refactor(amd): remove superseded DSV4 launcher path --- .../agentic/dsv4_fp4_mi355x_sglang-disagg.sh | 184 ------------------ benchmarks/multi_node/amd_utils/job.slurm | 4 - benchmarks/multi_node/amd_utils/models.yaml | 35 ---- .../multi_node/amd_utils/server_sglang.sh | 48 +---- perf-changelog.yaml | 1 + 5 files changed, 4 insertions(+), 268 deletions(-) delete mode 100755 benchmarks/multi_node/agentic/dsv4_fp4_mi355x_sglang-disagg.sh diff --git a/benchmarks/multi_node/agentic/dsv4_fp4_mi355x_sglang-disagg.sh b/benchmarks/multi_node/agentic/dsv4_fp4_mi355x_sglang-disagg.sh deleted file mode 100755 index f43c81e1c..000000000 --- a/benchmarks/multi_node/agentic/dsv4_fp4_mi355x_sglang-disagg.sh +++ /dev/null @@ -1,184 +0,0 @@ -#!/usr/bin/env bash - -# Agentic trace-replay recipe for a disaggregated SGLang server on MI355X -# (DeepSeek-V4-Pro FP4, 1P1D TP8). -# -# CI-style sibling of dsr1_fp4_mi355x_sglang-disagg.sh: driven entirely by -# environment variables and submits a SLURM job via submit.sh. The agentic / -# HiCache-offload configuration mirrors the DSR1 recipe but uses DSV4-Pro -# specific flags (dsv4 attention backend, page-size 256, SWA settings). - -SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" -source "$SCRIPT_DIR/../../benchmark_lib.sh" - -check_env_vars \ - CONC_LIST \ - ISL \ - OSL \ - IMAGE \ - SPEC_DECODING \ - MODEL_PATH \ - 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 - -# Use upstreamed multi_node scripts (no external clone needed) -cd "$GITHUB_WORKSPACE/benchmarks/multi_node/amd_utils" || exit 1 - -# Set up SGL launch script-specific environment variables -export TIME_LIMIT="${TIME_LIMIT:-08:00:00}" -export MODEL_PATH=$MODEL_PATH -export MODEL_NAME=$MODEL_NAME -export CONTAINER_IMAGE=$IMAGE - -# ── Identity / result naming ── -export MODEL_PREFIX="${MODEL_PREFIX:-dsv4}" -export PRECISION="${PRECISION:-fp4}" -export RESULT_FILENAME="${RESULT_FILENAME:-${RUNNER_NAME:-dsv4-fp4-agentic}}" - -# ── Agentic benchmark params ── -export DURATION="${DURATION:-1800}" -# DSV4-Pro max model len for agentic traces (matches single-node recipe). -export MAX_MODEL_LEN="${MAX_MODEL_LEN:-1000000}" - -# ── In-tree sglang patches ── -# mori_conn.py targets hybrid-state bugs (GLM-5, Qwen3.5). DSV4-Pro uses a -# pure MoE/DSA architecture without hybrid state; skip to avoid interference. -export MORI_CONN_PATCH="${MORI_CONN_PATCH:-skip}" - -# ── Aiter fault mitigation ── -# --disable-custom-all-reduce avoids a known aiter fault on MI355X. -export DISABLE_CUSTOM_ALL_REDUCE="${DISABLE_CUSTOM_ALL_REDUCE:-0}" - -# ── KV cache offloading (HiCache) ── -# KV_OFFLOADING=none | dram (passed from YAML; default none for disagg). -# KV_OFFLOAD_BACKEND selects the backend when offloading is on; this recipe -# only implements HiCache, so "hicache" is the only supported value. -# HICACHE_TIER: L2 -> GPU + CPU-DRAM host pool. L3 -> + Mooncake store. -export KV_OFFLOADING="${KV_OFFLOADING:-none}" -if [[ "$KV_OFFLOADING" != "none" ]]; then - export KV_OFFLOAD_BACKEND="${KV_OFFLOAD_BACKEND:-hicache}" -fi -# HiCache/Mooncake tunables only matter when KV offloading is enabled. -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}" - # DSV4 uses page-size 256 (set in models.yaml); HiCache must match. - export HICACHE_PAGE_SIZE="${HICACHE_PAGE_SIZE:-256}" - # HiCache ratio (host pool = ratio * GPU KV pool). - export HICACHE_RATIO="${HICACHE_RATIO:-4}" - # server_sglang.sh prefers an absolute --hicache-size (derived from - # TOTAL_CPU_DRAM_GB, the sweep generator's per-node DRAM budget) over - # --hicache-ratio whenever TOTAL_CPU_DRAM_GB is set. DSv4 wants the - # ratio-based pool instead. Use FORCE_HICACHE_RATIO to opt out of the - # --hicache-size path rather than unsetting TOTAL_CPU_DRAM_GB itself: - # that var is also the shared client-side gate (benchmark_lib.sh requires - # it to be a positive integer whenever KV_OFFLOADING=dram) and gets - # forwarded into the aiperf sibling container's client.env, so unsetting - # it here made the client fail its own env validation before benchmarking - # ("DRAM KV offloading requires a positive configured TOTAL_CPU_DRAM_GB - # capacity") even though the servers came up fine. - export FORCE_HICACHE_RATIO=1 - - # ── HiCache layout/backend by tier ── - # L3 (Mooncake): page_first + direct + write_through + storage=mooncake - # L2 (CPU DRAM): layer_first + direct + write_through_selective + storage=none - # NOTE: write_through_selective evicts only under GPU memory pressure, avoiding - # the mori RDMA race that causes GPU memory access faults with write_through. - 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 - 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:-}" - fi - export HICACHE_PREFETCH_POLICY="${HICACHE_PREFETCH_POLICY:-best_effort}" - # Shared nodes: use non-default Mooncake ports to avoid collisions. - 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 - -# ── MoRIIO RDMA Send Queue tuning ── -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:-32768}" - -# ── SGLang PD router policy + server metrics ── -export PREFILL_ROUTER_POLICY="${PREFILL_ROUTER_POLICY:-consistent_hashing}" -export ENABLE_METRICS="${ENABLE_METRICS:-1}" - -# ── MTP ── -export DECODE_MTP_SIZE="${DECODE_MTP_SIZE:-0}" - -# Derive EP/DP enable flags from the topology inputs. -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 - -# Launch the job. CONC_LIST is space-delimited in YAML; submit.sh wants 'x'. -JOB_ID=$(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}) - -if [[ $? -ne 0 ]]; then - echo "Failed to submit job" >&2 - exit 1 -fi - -echo "$JOB_ID" diff --git a/benchmarks/multi_node/amd_utils/job.slurm b/benchmarks/multi_node/amd_utils/job.slurm index 3c8a6f16a..e92a7e6a9 100755 --- a/benchmarks/multi_node/amd_utils/job.slurm +++ b/benchmarks/multi_node/amd_utils/job.slurm @@ -458,7 +458,6 @@ DOCKER_ENV_COMMON=( -e PREFILL_DP_ATTN=\${PREFILL_DP_ATTN:-false} -e PREFILL_CONTEXT_LENGTH=\${PREFILL_CONTEXT_LENGTH:-} -e PREFILL_CHUNKED_PREFILL_SIZE=\${PREFILL_CHUNKED_PREFILL_SIZE:-} - -e DISABLE_CUSTOM_ALL_REDUCE=\${DISABLE_CUSTOM_ALL_REDUCE:-} -e MAX_MODEL_LEN=\${MAX_MODEL_LEN:-} -e DURATION=\${DURATION:-1800} -e IS_AGENTIC=\${IS_AGENTIC:-0} @@ -467,8 +466,6 @@ DOCKER_ENV_COMMON=( -e KV_OFFLOAD_BACKEND_METADATA=\"\${KV_OFFLOAD_BACKEND_METADATA:-}\" -e TOTAL_CPU_DRAM_GB=\${TOTAL_CPU_DRAM_GB:-} -e ENABLE_METRICS=\${ENABLE_METRICS:-0} - -e PREFILL_ROUTER_POLICY=\${PREFILL_ROUTER_POLICY:-random} - -e DECODE_ROUTER_POLICY=\${DECODE_ROUTER_POLICY:-random} -e MORI_IO_SQ_BACKOFF_TIMEOUT_US=\${MORI_IO_SQ_BACKOFF_TIMEOUT_US:-} -e MORI_IO_QP_MAX_SEND_WR=\${MORI_IO_QP_MAX_SEND_WR:-} -e DECODE_TP_SIZE=\$DECODE_TP_SIZE @@ -538,7 +535,6 @@ fi HICACHE_MC_CONFIG="${BENCHMARK_LOGS_DIR}/hicache_mc_${SLURM_JOB_ID}.env" cat > "$HICACHE_MC_CONFIG" <=16 queue-saturation / decode-stall failure mode). - # Mirrors the single-node DSv4 agentic recipe (dsv4_fp4_mi355x.sh=8192). - # Was 65280 (255*256, the dsv4 compressor kernel uint16 token cap); 8192 - # (32*256) stays a page-size multiple well under that cap. - chunked_prefill_size: 8192 - 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-V4-Pro-DI: base_flags: "--decode-log-interval 100 --log-level info --watchdog-timeout 3600 --load-balance-method round_robin --kv-cache-dtype fp8_e4m3 --attention-backend dsv4 --page-size 256 --swa-full-tokens-ratio 0.1 --disable-shared-experts-fusion --tool-call-parser deepseekv4 --reasoning-parser deepseek-v4 --disaggregation-transfer-backend mori" # DeepSeek-V4 built-in MTP: load via the EAGLE spec path with eagle-topk 1 (a single diff --git a/benchmarks/multi_node/amd_utils/server_sglang.sh b/benchmarks/multi_node/amd_utils/server_sglang.sh index aaaca61ef..45716c03a 100755 --- a/benchmarks/multi_node/amd_utils/server_sglang.sh +++ b/benchmarks/multi_node/amd_utils/server_sglang.sh @@ -537,12 +537,6 @@ if [[ "$KV_OFFLOADING" != "none" && "$KV_OFFLOAD_BACKEND" == "hicache" ]]; then # per-node DRAM budget computed by the sweep generator (enforcement); fall # back to --hicache-ratio (relative to the GPU KV pool) when no budget is # provided, keeping configs that predate the budget unchanged. - # FORCE_HICACHE_RATIO lets a recipe opt into ratio-based sizing without - # unsetting TOTAL_CPU_DRAM_GB — that var is also the shared client-side - # gate (benchmark_lib.sh requires it whenever KV_OFFLOADING=dram) and is - # forwarded verbatim into client.env below, so unsetting it here would - # make the aiperf client container fail its own env validation before - # 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 @@ -550,9 +544,7 @@ if [[ "$KV_OFFLOADING" != "none" && "$KV_OFFLOAD_BACKEND" == "hicache" ]]; then # 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 [[ -n "${TOTAL_CPU_DRAM_GB:-}" && "${TOTAL_CPU_DRAM_GB}" -gt 0 && "${MODEL_NAME}" != *DeepSeek-V4* ]]; 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 @@ -791,13 +783,7 @@ if [ "$NODE_RANK" -eq 0 ]; then # across the agentic trace; round_robin decode keeps the single decode worker # fed evenly. Override via ROUTER_RESILIENCE_FLAGS / ROUTER_POLICY_FLAGS. ROUTER_RESILIENCE_FLAGS="${ROUTER_RESILIENCE_FLAGS:---disable-circuit-breaker --health-failure-threshold 100 --health-check-timeout-secs 600 --health-check-interval-secs 30}" - # server_sglang.sh previously read ROUTER_PREFILL_POLICY, but the recipe - # scripts export PREFILL_ROUTER_POLICY, so the recipe's policy override was - # silently ignored and the router always fell back to this hardcoded - # default. Also comment out ROUTER_DECODE_POLICY for now (superseded by - # --dp-aware below). - ROUTER_PREFILL_POLICY="${PREFILL_ROUTER_POLICY:-consistent_hashing}" - # ROUTER_DECODE_POLICY="${ROUTER_DECODE_POLICY:-round_robin}" + ROUTER_PREFILL_POLICY="consistent_hashing" ROUTER_CACHE_THRESHOLD="${ROUTER_CACHE_THRESHOLD:-0.3}" ROUTER_BALANCE_ABS_THRESHOLD="${ROUTER_BALANCE_ABS_THRESHOLD:-2}" ROUTER_BALANCE_REL_THRESHOLD="${ROUTER_BALANCE_REL_THRESHOLD:-1.1}" @@ -1275,35 +1261,7 @@ else fi set +x - # Agentic trace replay doesn't reproduce real token-by-token traffic, so - # measured MTP/EAGLE acceptance there isn't representative (PR #2309 - # review: https://github.com/SemiAnalysisAI/InferenceX/pull/2309#pullrequestreview-4778348624). - # Per the AgentX fairness guidelines (golden_al_distribution/README.md), - # agentic throughput benchmarks simulate acceptance at the model's - # committed golden AL instead of measuring real (non-representative) - # acceptance. Eval runs (RUN_EVAL / EVAL_ONLY) need real acceptance so - # GSM8K scores reflect actual MTP behavior. Golden curve source: - # golden_al_distribution/dsv4_mtp.yaml (thinking_on). - DECODE_SIM_ACC_ENV="" - if [[ "$DECODE_MTP_SIZE" -gt 0 ]] && { [[ "${IS_AGENTIC:-0}" == "1" ]] || [[ "${IS_AGENTIC:-}" == "true" ]]; }; then - if [[ "${EVAL_ONLY:-false}" == "true" ]] || [[ "${RUN_EVAL:-false}" == "true" ]]; then - echo "[INFO] Eval mode: synthetic MTP disabled (using real acceptance)" - else - DSV4_GOLDEN_AL="" - case "${MODEL_NAME}:${DECODE_MTP_SIZE}" in - *DeepSeek-V4*:1) DSV4_GOLDEN_AL=1.79 ;; - *DeepSeek-V4*:2) DSV4_GOLDEN_AL=2.27 ;; - *DeepSeek-V4*:3) DSV4_GOLDEN_AL=2.49 ;; - esac - if [[ -n "$DSV4_GOLDEN_AL" ]]; then - DECODE_SIM_ACC_ENV="SGLANG_SIMULATE_ACC_LEN=${DSV4_GOLDEN_AL} SGLANG_SIMULATE_ACC_METHOD=match-expected SGLANG_SIMULATE_ACC_TOKEN_MODE=real-draft-token" - else - echo "WARNING: agentic MTP run (model=${MODEL_NAME}, DECODE_MTP_SIZE=${DECODE_MTP_SIZE}) has no golden AL wired in server_sglang.sh -- falling back to real (unsimulated, non-representative) acceptance. Add a case in server_sglang.sh and golden_al_distribution/ before shipping this arm. See golden_al_distribution/README.md." >&2 - fi - fi - fi - - DECODE_CMD="SGLANG_MORI_COMBINE_DTYPE=${MORI_COMBINE_DTYPE_DECODE} ${DECODE_MORI_MOE_ENV} SGLANG_MORI_NUM_MAX_DISPATCH_TOKENS_PER_RANK=${MORI_NUM_MAX_DISPATCH_TOKENS_PER_RANK_DECODE:-${MORI_MAX_DISPATCH_TOKENS_DECODE}} MORI_IO_SQ_BACKOFF_TIMEOUT_US=${MORI_IO_SQ_BACKOFF_TIMEOUT_US} MORI_IO_QP_MAX_SEND_WR=${MORI_IO_QP_MAX_SEND_WR} ${DECODE_SIM_ACC_ENV} ${LAUNCH_PREFIX:-} python3 -m sglang.launch_server \ + DECODE_CMD="SGLANG_MORI_COMBINE_DTYPE=${MORI_COMBINE_DTYPE_DECODE} ${DECODE_MORI_MOE_ENV} SGLANG_MORI_NUM_MAX_DISPATCH_TOKENS_PER_RANK=${MORI_NUM_MAX_DISPATCH_TOKENS_PER_RANK_DECODE:-${MORI_MAX_DISPATCH_TOKENS_DECODE}} MORI_IO_SQ_BACKOFF_TIMEOUT_US=${MORI_IO_SQ_BACKOFF_TIMEOUT_US} MORI_IO_QP_MAX_SEND_WR=${MORI_IO_QP_MAX_SEND_WR} ${LAUNCH_PREFIX:-} python3 -m sglang.launch_server \ --model-path ${MODEL_DIR}/${MODEL_NAME} \ --disaggregation-mode decode \ --disaggregation-ib-device ${IBDEVICES} \ diff --git a/perf-changelog.yaml b/perf-changelog.yaml index 598f9346e..7d1a56859 100644 --- a/perf-changelog.yaml +++ b/perf-changelog.yaml @@ -5795,4 +5795,5 @@ - "Port the merged MI355X DeepSeek V4 AgentX MTP 1P1D search space from the legacy AMD multinode launcher to srt-slurm." - "Use srt-slurm's native SGLang Router frontend and AMD MoRI transport without Dynamo, NATS, etcd, or NIXL." - "Preserve TP8/EP1 c1,c2,c4,c8,c16,c32 and TP8/EP8/DPA c96,c128, EAGLE MTP, the golden simulated average length, HiCache settings, and logical prefill/decode server-metrics collection through the custom AgentX benchmark." + - "Remove the superseded DSV4-specific amd_utils benchmark wrapper, model stanza, HiCache override plumbing, and hard-coded synthetic-acceptance injection; the srt-slurm recipes now own those settings directly." pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2555 From 2f6d590d7cfe022746f2d963c7dd712b59fda6ff Mon Sep 17 00:00:00 2001 From: Cam Quilici Date: Mon, 10 Aug 2026 21:12:06 -0500 Subject: [PATCH 04/18] fix(amd): respect Slurm CPU affinity for SGLang workers --- .../mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c128.yaml | 3 +++ .../mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c96.yaml | 3 +++ .../mi355x/agentic/disagg-1p1d-tp8-mtp-hicache.yaml | 3 +++ 3 files changed, 9 insertions(+) diff --git a/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c128.yaml b/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c128.yaml index ef864c559..ca5eb5b93 100644 --- a/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c128.yaml +++ b/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c128.yaml @@ -49,6 +49,9 @@ backend: HF_HUB_CACHE: "/hf_hub_cache/hub" HUGGINGFACE_HUB_CACHE: "/hf_hub_cache/hub" PYTHONUNBUFFERED: "1" + # Slurm owns CPU placement for each worker task. The ROCm image may inherit + # SGLang affinity settings that target CPUs outside the task cpuset. + SGLANG_SET_CPU_AFFINITY: "0" IBDEVICES: "rdma0,rdma1,rdma2,rdma3,rdma4,rdma5,rdma6,rdma7" MORI_RDMA_TC: "104" MORI_IO_SQ_BACKOFF_TIMEOUT_US: "500000" diff --git a/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c96.yaml b/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c96.yaml index 8bc6785c3..4cc33205d 100644 --- a/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c96.yaml +++ b/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c96.yaml @@ -49,6 +49,9 @@ backend: HF_HUB_CACHE: "/hf_hub_cache/hub" HUGGINGFACE_HUB_CACHE: "/hf_hub_cache/hub" PYTHONUNBUFFERED: "1" + # Slurm owns CPU placement for each worker task. The ROCm image may inherit + # SGLang affinity settings that target CPUs outside the task cpuset. + SGLANG_SET_CPU_AFFINITY: "0" IBDEVICES: "rdma0,rdma1,rdma2,rdma3,rdma4,rdma5,rdma6,rdma7" MORI_RDMA_TC: "104" MORI_IO_SQ_BACKOFF_TIMEOUT_US: "500000" diff --git a/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-tp8-mtp-hicache.yaml b/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-tp8-mtp-hicache.yaml index 815ec2868..f0905f004 100644 --- a/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-tp8-mtp-hicache.yaml +++ b/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-tp8-mtp-hicache.yaml @@ -49,6 +49,9 @@ backend: HF_HUB_CACHE: "/hf_hub_cache/hub" HUGGINGFACE_HUB_CACHE: "/hf_hub_cache/hub" PYTHONUNBUFFERED: "1" + # Slurm owns CPU placement for each worker task. The ROCm image may inherit + # SGLang affinity settings that target CPUs outside the task cpuset. + SGLANG_SET_CPU_AFFINITY: "0" IBDEVICES: "rdma0,rdma1,rdma2,rdma3,rdma4,rdma5,rdma6,rdma7" MORI_RDMA_TC: "104" MORI_IO_SQ_BACKOFF_TIMEOUT_US: "500000" From 4e77ff0b0a14a8e29de488304756b6266d719c43 Mon Sep 17 00:00:00 2001 From: Cam Quilici Date: Mon, 10 Aug 2026 21:37:56 -0500 Subject: [PATCH 05/18] fix(amd): allocate full MI355X CPU topology for DSV4 --- .../mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c128.yaml | 7 ++++--- .../mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c96.yaml | 7 ++++--- .../mi355x/agentic/disagg-1p1d-tp8-mtp-hicache.yaml | 7 ++++--- 3 files changed, 12 insertions(+), 9 deletions(-) diff --git a/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c128.yaml b/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c128.yaml index ca5eb5b93..7dfcc43c7 100644 --- a/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c128.yaml +++ b/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c128.yaml @@ -49,9 +49,6 @@ backend: HF_HUB_CACHE: "/hf_hub_cache/hub" HUGGINGFACE_HUB_CACHE: "/hf_hub_cache/hub" PYTHONUNBUFFERED: "1" - # Slurm owns CPU placement for each worker task. The ROCm image may inherit - # SGLang affinity settings that target CPUs outside the task cpuset. - SGLANG_SET_CPU_AFFINITY: "0" IBDEVICES: "rdma0,rdma1,rdma2,rdma3,rdma4,rdma5,rdma6,rdma7" MORI_RDMA_TC: "104" MORI_IO_SQ_BACKOFF_TIMEOUT_US: "500000" @@ -122,6 +119,10 @@ backend: cuda-graph-bs: [1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31, 32, 33, 34, 35, 36, 37, 38, 39, 40, 41, 42, 43, 44, 45, 46, 47, 48, 49, 50, 51, 52, 53, 54, 55, 56, 57, 58, 59, 60, 61, 62, 63, 64, 65, 66, 67, 68, 69, 70, 71, 72, 73, 74, 75, 76, 77, 78, 79, 80, 81, 82, 83, 84, 85, 86, 87, 88, 89, 90, 91, 92, 93, 94, 95, 96, 97, 98, 99, 100, 101, 102, 103, 104, 105, 106, 107, 108, 109, 110, 111, 112, 113, 114, 115, 116, 117, 118, 119, 120, 121, 122, 123, 124, 125, 126, 127, 128] sbatch_directives: + # SGLang maps the eight TP ranks across the MI355X host CPU topology. Give + # each backend task the complete node cpuset so every placement stays within + # its Slurm allocation. + cpus-per-task: "128" mem: "0" srun_options: diff --git a/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c96.yaml b/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c96.yaml index 4cc33205d..fbe0c99df 100644 --- a/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c96.yaml +++ b/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c96.yaml @@ -49,9 +49,6 @@ backend: HF_HUB_CACHE: "/hf_hub_cache/hub" HUGGINGFACE_HUB_CACHE: "/hf_hub_cache/hub" PYTHONUNBUFFERED: "1" - # Slurm owns CPU placement for each worker task. The ROCm image may inherit - # SGLang affinity settings that target CPUs outside the task cpuset. - SGLANG_SET_CPU_AFFINITY: "0" IBDEVICES: "rdma0,rdma1,rdma2,rdma3,rdma4,rdma5,rdma6,rdma7" MORI_RDMA_TC: "104" MORI_IO_SQ_BACKOFF_TIMEOUT_US: "500000" @@ -122,6 +119,10 @@ backend: cuda-graph-bs: [1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31, 32, 33, 34, 35, 36, 37, 38, 39, 40, 41, 42, 43, 44, 45, 46, 47, 48, 49, 50, 51, 52, 53, 54, 55, 56, 57, 58, 59, 60, 61, 62, 63, 64, 65, 66, 67, 68, 69, 70, 71, 72, 73, 74, 75, 76, 77, 78, 79, 80, 81, 82, 83, 84, 85, 86, 87, 88, 89, 90, 91, 92, 93, 94, 95, 96, 97, 98, 99, 100, 101, 102, 103, 104, 105, 106, 107, 108, 109, 110, 111, 112, 113, 114, 115, 116, 117, 118, 119, 120, 121, 122, 123, 124, 125, 126, 127, 128] sbatch_directives: + # SGLang maps the eight TP ranks across the MI355X host CPU topology. Give + # each backend task the complete node cpuset so every placement stays within + # its Slurm allocation. + cpus-per-task: "128" mem: "0" srun_options: diff --git a/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-tp8-mtp-hicache.yaml b/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-tp8-mtp-hicache.yaml index f0905f004..490286d09 100644 --- a/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-tp8-mtp-hicache.yaml +++ b/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-tp8-mtp-hicache.yaml @@ -49,9 +49,6 @@ backend: HF_HUB_CACHE: "/hf_hub_cache/hub" HUGGINGFACE_HUB_CACHE: "/hf_hub_cache/hub" PYTHONUNBUFFERED: "1" - # Slurm owns CPU placement for each worker task. The ROCm image may inherit - # SGLang affinity settings that target CPUs outside the task cpuset. - SGLANG_SET_CPU_AFFINITY: "0" IBDEVICES: "rdma0,rdma1,rdma2,rdma3,rdma4,rdma5,rdma6,rdma7" MORI_RDMA_TC: "104" MORI_IO_SQ_BACKOFF_TIMEOUT_US: "500000" @@ -116,6 +113,10 @@ backend: cuda-graph-bs: [1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31, 32, 33, 34, 35, 36, 37, 38, 39, 40, 41, 42, 43, 44, 45, 46, 47, 48, 49, 50, 51, 52, 53, 54, 55, 56, 57, 58, 59, 60, 61, 62, 63, 64, 65, 66, 67, 68, 69, 70, 71, 72, 73, 74, 75, 76, 77, 78, 79, 80, 81, 82, 83, 84, 85, 86, 87, 88, 89, 90, 91, 92, 93, 94, 95, 96, 97, 98, 99, 100, 101, 102, 103, 104, 105, 106, 107, 108, 109, 110, 111, 112, 113, 114, 115, 116, 117, 118, 119, 120, 121, 122, 123, 124, 125, 126, 127, 128] sbatch_directives: + # SGLang maps the eight TP ranks across the MI355X host CPU topology. Give + # each backend task the complete node cpuset so every placement stays within + # its Slurm allocation. + cpus-per-task: "128" mem: "0" srun_options: From 7fd1a27082847e9e2493b4c927921d249bd9a250 Mon Sep 17 00:00:00 2001 From: Cam Quilici Date: Mon, 10 Aug 2026 21:46:00 -0500 Subject: [PATCH 06/18] chore(amd): pin resilient srt-slurm diagnostics --- perf-changelog.yaml | 1 + runners/launch_mi355x-amds-srt.sh | 2 +- 2 files changed, 2 insertions(+), 1 deletion(-) diff --git a/perf-changelog.yaml b/perf-changelog.yaml index 7d1a56859..0db759f60 100644 --- a/perf-changelog.yaml +++ b/perf-changelog.yaml @@ -5796,4 +5796,5 @@ - "Use srt-slurm's native SGLang Router frontend and AMD MoRI transport without Dynamo, NATS, etcd, or NIXL." - "Preserve TP8/EP1 c1,c2,c4,c8,c16,c32 and TP8/EP8/DPA c96,c128, EAGLE MTP, the golden simulated average length, HiCache settings, and logical prefill/decode server-metrics collection through the custom AgentX benchmark." - "Remove the superseded DSV4-specific amd_utils benchmark wrapper, model stanza, HiCache override plumbing, and hard-coded synthetic-acceptance injection; the srt-slurm recipes now own those settings directly." + - "Allocate all 128 logical CPUs per MI355X backend node so SGLang TP8 affinity follows the complete host topology, and pin srt-slurm diagnostics that preserve failure-log tails containing non-UTF-8 ROCm compiler output." pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2555 diff --git a/runners/launch_mi355x-amds-srt.sh b/runners/launch_mi355x-amds-srt.sh index 67388fb35..007a67351 100755 --- a/runners/launch_mi355x-amds-srt.sh +++ b/runners/launch_mi355x-amds-srt.sh @@ -5,7 +5,7 @@ set -euo pipefail # in explicitly with CONFIG_FILE; all existing MI355X launch behavior remains # unchanged for every other row. SRT_SLURM_REPOSITORY="https://github.com/SemiAnalysisAI/srt-slurm.git" -SRT_SLURM_COMMIT="31e72da43ed21fe941c039be51b2cad1a3cf428a" +SRT_SLURM_COMMIT="8bd8aef4089174d9e5acdab6e99184b63615255e" SLURM_PARTITION="compute" SHARED_BASE="/it-share/gharunners2/srt-slurm" SHARED_HF_CACHE="/it-share/hf-hub-cache" From 8851bb6f4ad12f7202138bdbede217ed11a98467 Mon Sep 17 00:00:00 2001 From: Cam Quilici Date: Mon, 10 Aug 2026 22:13:40 -0500 Subject: [PATCH 07/18] fix(amd): restore DSV4 ROCm kernel routing --- .../disagg-1p1d-dep8-mtp-hicache-c128.yaml | 35 +++++++++++++++++++ .../disagg-1p1d-dep8-mtp-hicache-c96.yaml | 35 +++++++++++++++++++ .../agentic/disagg-1p1d-tp8-mtp-hicache.yaml | 35 +++++++++++++++++++ perf-changelog.yaml | 1 + 4 files changed, 106 insertions(+) diff --git a/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c128.yaml b/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c128.yaml index 7dfcc43c7..48af0094b 100644 --- a/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c128.yaml +++ b/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c128.yaml @@ -49,6 +49,41 @@ backend: HF_HUB_CACHE: "/hf_hub_cache/hub" HUGGINGFACE_HUB_CACHE: "/hf_hub_cache/hub" PYTHONUNBUFFERED: "1" + # Validated DeepSeek V4 ROCm kernel routing. These settings keep FP4 MoE, + # indexer, MLA, and fused-compress dispatch on the supported MI355X paths. + TORCH_NCCL_BLOCKING_WAIT: "1" + NCCL_BLOCKING_WAIT: "1" + HSA_NO_SCRATCH_RECLAIM: "0" + GPU_MAX_HW_QUEUES: "5" + SGLANG_AITER_MLA_PERSIST: "0" + SGLANG_ENABLE_UNIFIED_RADIX_TREE: "1" + SGLANG_OPT_UNIFIED_CACHE_FREE_OUT_OF_WINDOW_SLOTS: "1" + SGLANG_MORI_DISPATCH_DTYPE: "auto" + SGLANG_MORI_DISPATCH_INTER_KERNEL_SWITCH_THRESHOLD: "4096" + SGLANG_DEFAULT_THINKING: "1" + SGLANG_DSV4_REASONING_EFFORT: "high" + SGLANG_OPT_DEEPGEMM_HC_PRENORM: "false" + SGLANG_USE_AITER: "1" + SGLANG_USE_ROCM700A: "0" + SGLANG_OPT_USE_FUSED_COMPRESS: "true" + SGLANG_HACK_FLASHMLA_BACKEND: "unified_kv_triton" + SGLANG_OPT_FP8_WO_A_GEMM: "false" + SGLANG_OPT_USE_JIT_INDEXER_METADATA: "false" + SGLANG_OPT_USE_TOPK_V2: "false" + SGLANG_OPT_USE_AITER_INDEXER: "true" + SGLANG_OPT_USE_TILELANG_INDEXER: "false" + SGLANG_OPT_USE_TILELANG_MHC_PRE: "false" + SGLANG_OPT_USE_TILELANG_MHC_POST: "false" + SGLANG_FP8_PAGED_MQA_LOGITS_TORCH: "1" + SGLANG_OPT_USE_FUSED_COMPRESS_TRITON: "true" + SGLANG_OPT_USE_MULTI_STREAM_OVERLAP: "false" + SGLANG_ROCM_USE_MULTI_STREAM: "false" + AITER_BF16_FP8_MOE_BOUND: "0" + SGLANG_EAGER_INPUT_NO_COPY: "true" + SGLANG_SHARED_EXPERT_TP1: "1" + SGLANG_DP_SHARED_EXPERT_LOCAL: "1" + SGLANG_DP_USE_GATHERV: "1" + SGLANG_DP_USE_REDUCE_SCATTER: "1" IBDEVICES: "rdma0,rdma1,rdma2,rdma3,rdma4,rdma5,rdma6,rdma7" MORI_RDMA_TC: "104" MORI_IO_SQ_BACKOFF_TIMEOUT_US: "500000" diff --git a/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c96.yaml b/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c96.yaml index fbe0c99df..aa73a641a 100644 --- a/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c96.yaml +++ b/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c96.yaml @@ -49,6 +49,41 @@ backend: HF_HUB_CACHE: "/hf_hub_cache/hub" HUGGINGFACE_HUB_CACHE: "/hf_hub_cache/hub" PYTHONUNBUFFERED: "1" + # Validated DeepSeek V4 ROCm kernel routing. These settings keep FP4 MoE, + # indexer, MLA, and fused-compress dispatch on the supported MI355X paths. + TORCH_NCCL_BLOCKING_WAIT: "1" + NCCL_BLOCKING_WAIT: "1" + HSA_NO_SCRATCH_RECLAIM: "0" + GPU_MAX_HW_QUEUES: "5" + SGLANG_AITER_MLA_PERSIST: "0" + SGLANG_ENABLE_UNIFIED_RADIX_TREE: "1" + SGLANG_OPT_UNIFIED_CACHE_FREE_OUT_OF_WINDOW_SLOTS: "1" + SGLANG_MORI_DISPATCH_DTYPE: "auto" + SGLANG_MORI_DISPATCH_INTER_KERNEL_SWITCH_THRESHOLD: "4096" + SGLANG_DEFAULT_THINKING: "1" + SGLANG_DSV4_REASONING_EFFORT: "high" + SGLANG_OPT_DEEPGEMM_HC_PRENORM: "false" + SGLANG_USE_AITER: "1" + SGLANG_USE_ROCM700A: "0" + SGLANG_OPT_USE_FUSED_COMPRESS: "true" + SGLANG_HACK_FLASHMLA_BACKEND: "unified_kv_triton" + SGLANG_OPT_FP8_WO_A_GEMM: "false" + SGLANG_OPT_USE_JIT_INDEXER_METADATA: "false" + SGLANG_OPT_USE_TOPK_V2: "false" + SGLANG_OPT_USE_AITER_INDEXER: "true" + SGLANG_OPT_USE_TILELANG_INDEXER: "false" + SGLANG_OPT_USE_TILELANG_MHC_PRE: "false" + SGLANG_OPT_USE_TILELANG_MHC_POST: "false" + SGLANG_FP8_PAGED_MQA_LOGITS_TORCH: "1" + SGLANG_OPT_USE_FUSED_COMPRESS_TRITON: "true" + SGLANG_OPT_USE_MULTI_STREAM_OVERLAP: "false" + SGLANG_ROCM_USE_MULTI_STREAM: "false" + AITER_BF16_FP8_MOE_BOUND: "0" + SGLANG_EAGER_INPUT_NO_COPY: "true" + SGLANG_SHARED_EXPERT_TP1: "1" + SGLANG_DP_SHARED_EXPERT_LOCAL: "1" + SGLANG_DP_USE_GATHERV: "1" + SGLANG_DP_USE_REDUCE_SCATTER: "1" IBDEVICES: "rdma0,rdma1,rdma2,rdma3,rdma4,rdma5,rdma6,rdma7" MORI_RDMA_TC: "104" MORI_IO_SQ_BACKOFF_TIMEOUT_US: "500000" diff --git a/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-tp8-mtp-hicache.yaml b/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-tp8-mtp-hicache.yaml index 490286d09..f61e40841 100644 --- a/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-tp8-mtp-hicache.yaml +++ b/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-tp8-mtp-hicache.yaml @@ -49,6 +49,41 @@ backend: HF_HUB_CACHE: "/hf_hub_cache/hub" HUGGINGFACE_HUB_CACHE: "/hf_hub_cache/hub" PYTHONUNBUFFERED: "1" + # Validated DeepSeek V4 ROCm kernel routing. These settings keep FP4 MoE, + # indexer, MLA, and fused-compress dispatch on the supported MI355X paths. + TORCH_NCCL_BLOCKING_WAIT: "1" + NCCL_BLOCKING_WAIT: "1" + HSA_NO_SCRATCH_RECLAIM: "0" + GPU_MAX_HW_QUEUES: "5" + SGLANG_AITER_MLA_PERSIST: "0" + SGLANG_ENABLE_UNIFIED_RADIX_TREE: "1" + SGLANG_OPT_UNIFIED_CACHE_FREE_OUT_OF_WINDOW_SLOTS: "1" + SGLANG_MORI_DISPATCH_DTYPE: "auto" + SGLANG_MORI_DISPATCH_INTER_KERNEL_SWITCH_THRESHOLD: "4096" + SGLANG_DEFAULT_THINKING: "1" + SGLANG_DSV4_REASONING_EFFORT: "high" + SGLANG_OPT_DEEPGEMM_HC_PRENORM: "false" + SGLANG_USE_AITER: "1" + SGLANG_USE_ROCM700A: "0" + SGLANG_OPT_USE_FUSED_COMPRESS: "true" + SGLANG_HACK_FLASHMLA_BACKEND: "unified_kv_triton" + SGLANG_OPT_FP8_WO_A_GEMM: "false" + SGLANG_OPT_USE_JIT_INDEXER_METADATA: "false" + SGLANG_OPT_USE_TOPK_V2: "false" + SGLANG_OPT_USE_AITER_INDEXER: "true" + SGLANG_OPT_USE_TILELANG_INDEXER: "false" + SGLANG_OPT_USE_TILELANG_MHC_PRE: "false" + SGLANG_OPT_USE_TILELANG_MHC_POST: "false" + SGLANG_FP8_PAGED_MQA_LOGITS_TORCH: "1" + SGLANG_OPT_USE_FUSED_COMPRESS_TRITON: "true" + SGLANG_OPT_USE_MULTI_STREAM_OVERLAP: "false" + SGLANG_ROCM_USE_MULTI_STREAM: "false" + AITER_BF16_FP8_MOE_BOUND: "0" + SGLANG_EAGER_INPUT_NO_COPY: "true" + SGLANG_SHARED_EXPERT_TP1: "1" + SGLANG_DP_SHARED_EXPERT_LOCAL: "1" + SGLANG_DP_USE_GATHERV: "1" + SGLANG_DP_USE_REDUCE_SCATTER: "1" IBDEVICES: "rdma0,rdma1,rdma2,rdma3,rdma4,rdma5,rdma6,rdma7" MORI_RDMA_TC: "104" MORI_IO_SQ_BACKOFF_TIMEOUT_US: "500000" diff --git a/perf-changelog.yaml b/perf-changelog.yaml index 0db759f60..44ad9d824 100644 --- a/perf-changelog.yaml +++ b/perf-changelog.yaml @@ -5797,4 +5797,5 @@ - "Preserve TP8/EP1 c1,c2,c4,c8,c16,c32 and TP8/EP8/DPA c96,c128, EAGLE MTP, the golden simulated average length, HiCache settings, and logical prefill/decode server-metrics collection through the custom AgentX benchmark." - "Remove the superseded DSV4-specific amd_utils benchmark wrapper, model stanza, HiCache override plumbing, and hard-coded synthetic-acceptance injection; the srt-slurm recipes now own those settings directly." - "Allocate all 128 logical CPUs per MI355X backend node so SGLang TP8 affinity follows the complete host topology, and pin srt-slurm diagnostics that preserve failure-log tails containing non-UTF-8 ROCm compiler output." + - "Move the known-green DeepSeek V4 ROCm kernel-routing contract into each self-contained srt-slurm recipe; this avoids the unsupported AITer fused-MoE heuristic dispatch path while keeping legacy harness-only environment indirection removed." pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2555 From 7d53ce265200105bedc548445a1472d579a37cd7 Mon Sep 17 00:00:00 2001 From: Cam Quilici Date: Mon, 10 Aug 2026 23:15:22 -0500 Subject: [PATCH 08/18] fix(amd): pin DSV4 runtime identity --- .../mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c128.yaml | 4 ++-- .../mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c96.yaml | 4 ++-- .../mi355x/agentic/disagg-1p1d-tp8-mtp-hicache.yaml | 4 ++-- perf-changelog.yaml | 1 + 4 files changed, 7 insertions(+), 6 deletions(-) diff --git a/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c128.yaml b/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c128.yaml index 48af0094b..0be10e096 100644 --- a/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c128.yaml +++ b/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c128.yaml @@ -11,9 +11,9 @@ identity: container: image: "lmsysorg/sglang-rocm:v0.5.15.post1-rocm720-mi35x-20260719" frameworks: - sglang: "0.5.15.post1" + sglang: "0.5.15.post1.dev20260719+ga03ca46a28" sglang-router: "0.3.2" - amd-mori: "bundled" + amd-mori: "0.5.15.post1.dev20260719+ga03ca46a28" slurm: time_limit: "04:00:00" diff --git a/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c96.yaml b/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c96.yaml index aa73a641a..ce6e63657 100644 --- a/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c96.yaml +++ b/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-dep8-mtp-hicache-c96.yaml @@ -11,9 +11,9 @@ identity: container: image: "lmsysorg/sglang-rocm:v0.5.15.post1-rocm720-mi35x-20260719" frameworks: - sglang: "0.5.15.post1" + sglang: "0.5.15.post1.dev20260719+ga03ca46a28" sglang-router: "0.3.2" - amd-mori: "bundled" + amd-mori: "0.5.15.post1.dev20260719+ga03ca46a28" slurm: time_limit: "04:00:00" diff --git a/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-tp8-mtp-hicache.yaml b/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-tp8-mtp-hicache.yaml index f61e40841..c9a7fd646 100644 --- a/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-tp8-mtp-hicache.yaml +++ b/benchmarks/multi_node/srt-slurm-recipes/sglang/deepseek-v4/mi355x/agentic/disagg-1p1d-tp8-mtp-hicache.yaml @@ -11,9 +11,9 @@ identity: container: image: "lmsysorg/sglang-rocm:v0.5.15.post1-rocm720-mi35x-20260719" frameworks: - sglang: "0.5.15.post1" + sglang: "0.5.15.post1.dev20260719+ga03ca46a28" sglang-router: "0.3.2" - amd-mori: "bundled" + amd-mori: "0.5.15.post1.dev20260719+ga03ca46a28" slurm: time_limit: "04:00:00" diff --git a/perf-changelog.yaml b/perf-changelog.yaml index 44ad9d824..d8b316165 100644 --- a/perf-changelog.yaml +++ b/perf-changelog.yaml @@ -5798,4 +5798,5 @@ - "Remove the superseded DSV4-specific amd_utils benchmark wrapper, model stanza, HiCache override plumbing, and hard-coded synthetic-acceptance injection; the srt-slurm recipes now own those settings directly." - "Allocate all 128 logical CPUs per MI355X backend node so SGLang TP8 affinity follows the complete host topology, and pin srt-slurm diagnostics that preserve failure-log tails containing non-UTF-8 ROCm compiler output." - "Move the known-green DeepSeek V4 ROCm kernel-routing contract into each self-contained srt-slurm recipe; this avoids the unsupported AITer fused-MoE heuristic dispatch path while keeping legacy harness-only environment indirection removed." + - "Declare the exact SGLang and AMD MoRI runtime fingerprint reported by the pinned ROCm image so srt-slurm identity verification is reproducible." pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2555 From ab17f87169b20fd43f7f3e587fe60e3b650654d3 Mon Sep 17 00:00:00 2001 From: Cam Quilici Date: Mon, 10 Aug 2026 23:36:33 -0500 Subject: [PATCH 09/18] test(amd): drop redundant launcher contract suite --- utils/test_mi300x_srt_slurm_contract.py | 225 ------------------------ 1 file changed, 225 deletions(-) delete mode 100644 utils/test_mi300x_srt_slurm_contract.py diff --git a/utils/test_mi300x_srt_slurm_contract.py b/utils/test_mi300x_srt_slurm_contract.py deleted file mode 100644 index d558d7afb..000000000 --- a/utils/test_mi300x_srt_slurm_contract.py +++ /dev/null @@ -1,225 +0,0 @@ -"""High-signal contract checks for the MI300X srt-slurm bring-up lane.""" - -import os -import subprocess -from pathlib import Path - -import yaml - - -REPO_ROOT = Path(__file__).resolve().parents[1] -RECIPE_PATH = ( - REPO_ROOT - / "benchmarks/multi_node/srt-slurm-recipes/vllm/qwen3-0.6b/mi300x/agg-fixed-seq.yaml" -) -DISAGG_RECIPE_PATH = RECIPE_PATH.with_name("disagg-1p1d-fixed-seq.yaml") -CLUSTER_PATH = ( - REPO_ROOT - / "benchmarks/multi_node/srt-slurm-recipes/cluster-configs/mi300x-amds.yaml" -) -MASTER_CONFIG_PATH = REPO_ROOT / "configs/amd-master.yaml" -SRT_LAUNCHER_PATH = REPO_ROOT / "runners/launch_mi300x-amds-srt.sh" - - -def test_mi300x_cluster_uses_the_rocm_slurm_contract(): - cluster = yaml.safe_load(CLUSTER_PATH.read_text()) - - assert cluster["accelerator_vendor"] == "amd" - assert cluster["network_interface"] is None - assert cluster["gpu_sbatch_directive"] == "gres" - assert cluster["use_segment_sbatch_directive"] is False - assert cluster["runtime_config_transport"] == "embedded" - assert cluster["default_sbatch_directives"]["exclude"] == ( - "chi-mi300x-049,chi-mi300x-121" - ) - assert cluster["default_mounts"]["/dev/kfd"] == "/dev/kfd" - assert cluster["default_mounts"]["/dev/dri"] == "/dev/dri" - image_path = cluster["containers"]["vllm-rocm-v0.26.0"] - assert image_path.endswith("/vllm-openai-rocm-v0.26.0.sqsh") - router_path = cluster["containers"]["vllm-router-20260809"] - assert router_path.endswith("/vllm-router-nightly-20260809-d2ba586.sqsh") - - for recipe_path in (RECIPE_PATH, DISAGG_RECIPE_PATH): - recipe = yaml.safe_load(recipe_path.read_text()) - assert recipe["model"]["container"] == "vllm-rocm-v0.26.0" - assert recipe["identity"]["container"]["image"] == ( - "vllm/vllm-openai-rocm:v0.26.0" - ) - - -def test_disaggregated_recipe_uses_native_router_and_moriio(): - recipe = yaml.safe_load(DISAGG_RECIPE_PATH.read_text()) - - assert recipe["frontend"] == { - "type": "vllm-router", - "enable_multiple_frontends": False, - "container_image": "vllm-router-20260809", - "args": { - "policy": "consistent_hash", - "prefill-policy": "consistent_hash", - "decode-policy": "consistent_hash", - }, - } - assert recipe["backend"]["connector"] == "moriio" - for role in ("prefill", "decode"): - assert recipe["backend"][f"{role}_environment"][ - "VLLM_ROCM_USE_AITER" - ] == "1" - assert recipe["backend"]["vllm_config"][role][ - "attention-backend" - ] == "ROCM_AITER_FA" - assert "dynamo" not in recipe - serialized = DISAGG_RECIPE_PATH.read_text().lower() - assert "nixl" not in serialized - assert "nats" not in serialized - assert "etcd" not in serialized - - -def test_official_matrix_routes_disagg_through_the_pinned_srt_launcher(): - config = yaml.safe_load(MASTER_CONFIG_PATH.read_text())[ - "qwen3-0.6b-fp16-mi300x-vllm-srt-disagg" - ] - search = config["scenarios"]["fixed-seq-len"][0]["search-space"][0] - launcher = SRT_LAUNCHER_PATH.read_text() - - assert config["runner"] == "mi300x-disagg" - assert config["router"] == { - "name": "vllm-router", - "version": "nightly-20260809-d2ba586", - } - assert config["kv-p2p-transfer"] == "moriio" - assert search["prefill"]["additional-settings"] == [ - "CONFIG_FILE=recipes/vllm/qwen3-0.6b/mi300x/" - "disagg-1p1d-fixed-seq.yaml" - ] - assert "315e4b06a7e0806194a646ea21832e750e896a46" in launcher - assert "make setup-compute ARCH=x86_64" in launcher - assert "--no-preflight" in launcher - assert 'VLLM_IMAGE="vllm/vllm-openai-rocm:v0.26.0"' in launcher - assert ( - 'VLLM_ROUTER_IMAGE="vllm/vllm-router:nightly-20260809-d2ba586"' - in launcher - ) - assert 'enroot import -o "\\$tmp" "docker://\\${image}"' in launcher - assert 'exec {lock_fd}>"\\${target}.lock"' in launcher - assert 'flock -w 2400 "\\$lock_fd"' in launcher - assert 'unsquashfs -s "\\$tmp"' in launcher - assert 'mv "\\$tmp" "\\$target"' in launcher - assert 'REMOTE_SRT_RUNTIME="${REMOTE_BASE}/runtime/srt-slurm-${SRT_SLURM_COMMIT}"' in launcher - assert 'git -C "\\$srt_runtime" checkout --quiet --detach "${SRT_SLURM_COMMIT}"' in launcher - assert 'make -C "\\$srt_runtime" --no-print-directory setup-compute ARCH=x86_64' in launcher - assert 'export SRTCTL_RUNTIME_SOURCE_DIR="$REMOTE_SRT_RUNTIME"' in launcher - assert "scancel" not in launcher - - -def test_official_matrix_routes_aggregate_through_the_pinned_srt_launcher(): - config = yaml.safe_load(MASTER_CONFIG_PATH.read_text())[ - "qwen3-0.6b-fp16-mi300x-vllm-srt-agg" - ] - search = config["scenarios"]["fixed-seq-len"][0]["search-space"] - launcher = SRT_LAUNCHER_PATH.read_text() - - assert config["runner"] == "mi300x-disagg" - assert config["multinode"] is True - assert config["disagg"] is False - assert search == [ - { - "conc-list": [1], - "prefill": { - "num-worker": 1, - "tp": 1, - "ep": 1, - "dp-attn": False, - "additional-settings": [ - "CONFIG_FILE=recipes/vllm/qwen3-0.6b/mi300x/" - "agg-fixed-seq.yaml" - ], - }, - "decode": { - "num-worker": 0, - "tp": 1, - "ep": 1, - "dp-attn": False, - }, - } - ] - assert ': "${CONFIG_FILE:?CONFIG_FILE must name an srt-slurm recipe}"' in launcher - assert 'JOB_BATCH_HOST=$(scontrol show job "$JOB_ID" -dd' in launcher - assert '--nodelist="$JOB_BATCH_HOST"' in launcher - assert "TOTAL_GPUS=$((PREFILL_NUM_WORKERS * PREFILL_TP" in launcher - - -def test_aggregate_recipe_uses_direct_vllm_without_dynamo_or_a_router(): - recipe = yaml.safe_load(RECIPE_PATH.read_text()) - - assert recipe["resources"] == { - "gpu_type": "mi300x", - "gpus_per_node": 1, - "agg_nodes": 1, - "agg_workers": 1, - "gpus_per_agg": 1, - } - assert recipe["frontend"] == { - "type": "vllm", - "enable_multiple_frontends": False, - } - assert recipe["backend"]["connector"] is None - serialized = RECIPE_PATH.read_text().lower() - assert "dynamo" not in serialized - assert "nixl" not in serialized - assert "moriio" not in serialized - - -def test_fixed_sequence_recipe_uses_inferencex_custom_benchmark(): - recipe = yaml.safe_load(RECIPE_PATH.read_text()) - benchmark = recipe["benchmark"] - command = benchmark["command"] - - assert benchmark["type"] == "custom" - assert "/infmax-workspace/utils/bench_serving/benchmark_serving.py" in command - assert 'result_root="/results/${SLURM_JOB_ID}"' in command - assert "--backend openai-chat" in command - assert "--endpoint /v1/chat/completions" in command - assert "--random-input-len 128" in command - assert "--random-output-len 32" in command - assert "--random-range-ratio 1.0" in command - assert "best-of" not in command - assert "sa-bench" not in command - - -def test_fixed_sequence_commands_keep_all_arguments_attached(tmp_path): - fake_bin = tmp_path / "bin" - fake_bin.mkdir() - fake_python = fake_bin / "python3" - fake_python.write_text( - "#!/bin/bash\n" - 'printf "%s\\n" "$@" >> "$FAKE_ARGS_LOG"\n' - 'printf "%s\\n" --CALL-END-- >> "$FAKE_ARGS_LOG"\n' - ) - fake_python.chmod(0o755) - - for recipe_path in (RECIPE_PATH, DISAGG_RECIPE_PATH): - command = yaml.safe_load(recipe_path.read_text())["benchmark"]["command"] - result_dir = tmp_path / recipe_path.stem - command = command.replace( - 'result_root="/results/${SLURM_JOB_ID}"', - f'result_root="{result_dir}"', - ) - args_log = tmp_path / f"{recipe_path.stem}.args" - env = { - **os.environ, - "PATH": f"{fake_bin}:{os.environ['PATH']}", - "FAKE_ARGS_LOG": str(args_log), - "SRT_FRONTEND_HOST": "127.0.0.1", - "SRT_FRONTEND_PORT": "8000", - "SLURM_JOB_ID": "123", - } - - subprocess.run(["bash", "-n"], input=command, text=True, check=True) - subprocess.run(["bash", "-c", command], env=env, check=True) - - calls = args_log.read_text().split("--CALL-END--\n") - calls = [[arg for arg in call.splitlines() if arg] for call in calls if call] - assert len(calls) == 2 - assert all("--model" in call and "Qwen/Qwen3-0.6B" in call for call in calls) - assert [call[call.index("--num-prompts") + 1] for call in calls] == ["4", "16"] From 3ad7e2d64dce89daabd7d819643f246291991fc8 Mon Sep 17 00:00:00 2001 From: Cam Quilici Date: Tue, 11 Aug 2026 08:40:16 -0500 Subject: [PATCH 10/18] fix(ci): restore e2e eval matrix output --- .github/workflows/e2e-tests.yml | 1 + utils/changelog_gate_tests/test_run_sweep_gating.py | 8 ++++++++ 2 files changed, 9 insertions(+) diff --git a/.github/workflows/e2e-tests.yml b/.github/workflows/e2e-tests.yml index c973df94e..d90f3fc54 100644 --- a/.github/workflows/e2e-tests.yml +++ b/.github/workflows/e2e-tests.yml @@ -261,6 +261,7 @@ jobs: MULTI_AGENTIC_EVAL=$(echo "$CONFIG_JSON" | python3 -c "import sys,json; d=json.load(sys.stdin); print(json.dumps([x for x in d if x.get('scenario-type') == 'agentic-coding' and 'prefill' in x and x.get('run-eval', False)]))" | score_matrix multi-agentic-eval) SINGLE=$(echo "$CONFIG_JSON" | python3 -c "import sys,json; d=json.load(sys.stdin); print(json.dumps([x for x in d if 'prefill' not in x and x.get('scenario-type') != 'agentic-coding' and not x.get('eval-only', False)]))" | score_matrix single) MULTI=$(echo "$CONFIG_JSON" | python3 -c "import sys,json; d=json.load(sys.stdin); print(json.dumps([x for x in d if 'prefill' in x and x.get('scenario-type') != 'agentic-coding' and not x.get('eval-only', False)]))" | score_matrix multi) + EVALS=$(echo "$CONFIG_JSON" | python3 -c "import sys,json; d=json.load(sys.stdin); print(json.dumps([x for x in d if 'prefill' not in x and x.get('scenario-type') != 'agentic-coding' and x.get('run-eval', False)]))" | score_matrix eval) MULTI_EVAL=$(echo "$CONFIG_JSON" | python3 -c "import sys,json; d=json.load(sys.stdin); print(json.dumps([x for x in d if 'prefill' in x and x.get('scenario-type') != 'agentic-coding' and x.get('run-eval', False)]))" | score_matrix multi-eval) { echo "agentic-config=$AGENTIC" diff --git a/utils/changelog_gate_tests/test_run_sweep_gating.py b/utils/changelog_gate_tests/test_run_sweep_gating.py index ed30ec5ad..a2ecad93d 100644 --- a/utils/changelog_gate_tests/test_run_sweep_gating.py +++ b/utils/changelog_gate_tests/test_run_sweep_gating.py @@ -417,6 +417,14 @@ def test_e2e_workflow_cannot_dispatch_database_ingest() -> None: assert "INFX_FRONTEND_PAT" not in workflow +def test_e2e_workflow_defines_every_emitted_matrix_variable() -> None: + workflow = (REPO_ROOT / ".github/workflows/e2e-tests.yml").read_text() + assigned = set(re.findall(r"^\s+([A-Z_]+)=", workflow, re.MULTILINE)) + emitted = set(re.findall(r'echo "[a-z-]+=\$([A-Z_]+)"', workflow)) + + assert emitted <= assigned + + def test_priority_classifier_runs_for_enabled_actions() -> None: scenario = { **_PR, From d0376da7b1fb46ff62ef27f03740bc803037fd98 Mon Sep 17 00:00:00 2001 From: Cam Quilici Date: Tue, 11 Aug 2026 10:15:48 -0500 Subject: [PATCH 11/18] fix(router): gate SGLang startup on backend readiness --- perf-changelog.yaml | 1 + runners/launch_mi355x-amds-srt.sh | 2 +- 2 files changed, 2 insertions(+), 1 deletion(-) diff --git a/perf-changelog.yaml b/perf-changelog.yaml index d8b316165..0974b7b8d 100644 --- a/perf-changelog.yaml +++ b/perf-changelog.yaml @@ -5799,4 +5799,5 @@ - "Allocate all 128 logical CPUs per MI355X backend node so SGLang TP8 affinity follows the complete host topology, and pin srt-slurm diagnostics that preserve failure-log tails containing non-UTF-8 ROCm compiler output." - "Move the known-green DeepSeek V4 ROCm kernel-routing contract into each self-contained srt-slurm recipe; this avoids the unsupported AITer fused-MoE heuristic dispatch path while keeping legacy harness-only environment indirection removed." - "Declare the exact SGLang and AMD MoRI runtime fingerprint reported by the pinned ROCm image so srt-slurm identity verification is reproducible." + - "Rebase the AMD runtime on the active native-router branch and wait for every advertised SGLang HTTP worker to become healthy before launching Model Gateway, preventing slow ROCm startup from permanently dropping a static prefill worker after the router's one-shot registration timeout." pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2555 diff --git a/runners/launch_mi355x-amds-srt.sh b/runners/launch_mi355x-amds-srt.sh index 007a67351..4669cc0cf 100755 --- a/runners/launch_mi355x-amds-srt.sh +++ b/runners/launch_mi355x-amds-srt.sh @@ -5,7 +5,7 @@ set -euo pipefail # in explicitly with CONFIG_FILE; all existing MI355X launch behavior remains # unchanged for every other row. SRT_SLURM_REPOSITORY="https://github.com/SemiAnalysisAI/srt-slurm.git" -SRT_SLURM_COMMIT="8bd8aef4089174d9e5acdab6e99184b63615255e" +SRT_SLURM_COMMIT="5c0d5a718194198447e8ceea5c37048c2fb8eaf5" SLURM_PARTITION="compute" SHARED_BASE="/it-share/gharunners2/srt-slurm" SHARED_HF_CACHE="/it-share/hf-hub-cache" From 4ef136887fa4510df2131aea54d5919a41b81593 Mon Sep 17 00:00:00 2001 From: Wenyao Gao <105094497+edwingao28@users.noreply.github.com> Date: Tue, 11 Aug 2026 09:33:39 -0700 Subject: [PATCH 12/18] =?UTF-8?q?Emit=20per-role=20average=20power=20on=20?= =?UTF-8?q?multinode=20disaggregated=20runs=20/=20=E5=A4=9A=E8=8A=82?= =?UTF-8?q?=E7=82=B9=E5=88=86=E7=A6=BB=E5=BC=8F=E8=BF=90=E8=A1=8C=E8=BE=93?= =?UTF-8?q?=E5=87=BA=E6=AF=8F=E8=A7=92=E8=89=B2=E5=B9=B3=E5=9D=87=E5=8A=9F?= =?UTF-8?q?=E7=8E=87=20(#2553)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- .github/AGENT_OPERATIONS.md | 2 +- utils/aggregate_power_multinode.py | 7 +++++++ utils/process_result.py | 2 ++ utils/test_aggregate_power_multinode.py | 21 +++++++++++++++++++++ utils/test_process_result.py | 2 ++ 5 files changed, 33 insertions(+), 1 deletion(-) diff --git a/.github/AGENT_OPERATIONS.md b/.github/AGENT_OPERATIONS.md index 1652ac87f..0ef9acd1e 100644 --- a/.github/AGENT_OPERATIONS.md +++ b/.github/AGENT_OPERATIONS.md @@ -92,7 +92,7 @@ For multi-node configurations, `--all-evals` creates one eval job per engine top Single-node fixed-sequence results may include `power_valid`, `avg_power_w`, `avg_total_gpu_power_w`, `total_gpu_energy_j`, and joules per query/input/output/total token. Invalid telemetry records `power_valid: 0` without energy metrics and fails only with `REQUIRE_POWER=1`. -Multinode disaggregated results add `prefill_gpu_energy_j`, `decode_gpu_energy_j`, `prefill_joules_per_input_token`, and `decode_joules_per_output_token`. Role energy covers the full formal benchmark window, not kernel-level phases. +Multinode disaggregated results add `prefill_gpu_energy_j`, `decode_gpu_energy_j`, `prefill_avg_power_w`, `decode_avg_power_w`, `prefill_joules_per_input_token`, and `decode_joules_per_output_token`. Role energy covers the full formal benchmark window, not kernel-level phases, and the role watts are that energy divided by the same window and by the role's GPU count. For srt-slurm recipes, `telemetry: {provider: dcgm-power}` enables official energy collection. `runners/launch_gb200-nv.sh` and `runners/launch_gb300-nv.sh` are the source of truth for `POWER_SRT_SLURM_PIN`. CI derives `POWER_PRODUCER_SHA` from the launcher stamp. `utils/test_gb200_power_official_contract.py` and `utils/test_gb300_power_official_contract.py` enforce the recipe/launcher contract. Only `PRECISION=fp8` dcgm-power lanes are validated. diff --git a/utils/aggregate_power_multinode.py b/utils/aggregate_power_multinode.py index 7460eecd2..cabed59d6 100644 --- a/utils/aggregate_power_multinode.py +++ b/utils/aggregate_power_multinode.py @@ -15,6 +15,9 @@ Role energy semantics: ``prefill_gpu_energy_j`` / ``decode_gpu_energy_j`` are the board-level energy of that role's GPUs integrated over the FULL formal serving window. They are not kernel-level prefill/decode phase energies. +``prefill_avg_power_w`` / ``decode_avg_power_w`` divide that role energy by the +same full window and by the role's GPU count, so they are the mean board draw +of that role's GPUs across the whole serving window, not a phase power. Ordinary benchmark runs are best-effort: invalid telemetry records ``power_valid=0`` (and no energy metrics) in the aggregate plus a validation @@ -108,6 +111,8 @@ ROLE_METRIC_KEYS = ( "prefill_gpu_energy_j", "decode_gpu_energy_j", + "prefill_avg_power_w", + "decode_avg_power_w", "prefill_joules_per_input_token", "decode_joules_per_output_token", ) @@ -1062,11 +1067,13 @@ def validate_and_integrate( } if prefill_gpus > 0: metrics["prefill_gpu_energy_j"] = role_energy["prefill"] + metrics["prefill_avg_power_w"] = role_energy["prefill"] / duration_s / prefill_gpus metrics["prefill_joules_per_input_token"] = ( role_energy["prefill"] / benchmark.total_input_tokens ) if decode_gpus > 0: metrics["decode_gpu_energy_j"] = role_energy["decode"] + metrics["decode_avg_power_w"] = role_energy["decode"] / duration_s / decode_gpus metrics["decode_joules_per_output_token"] = ( role_energy["decode"] / benchmark.total_output_tokens ) diff --git a/utils/process_result.py b/utils/process_result.py index 79e378cd0..a8bdc8cca 100644 --- a/utils/process_result.py +++ b/utils/process_result.py @@ -46,6 +46,8 @@ def get_optional_component_metadata(env_var): _MULTINODE_ROLE_METRIC_KEYS = ( "prefill_gpu_energy_j", "decode_gpu_energy_j", + "prefill_avg_power_w", + "decode_avg_power_w", "prefill_joules_per_input_token", "decode_joules_per_output_token", ) diff --git a/utils/test_aggregate_power_multinode.py b/utils/test_aggregate_power_multinode.py index fdc279643..e4c78a115 100644 --- a/utils/test_aggregate_power_multinode.py +++ b/utils/test_aggregate_power_multinode.py @@ -239,6 +239,8 @@ def test_emits_all_metrics_exactly(self, tmp_path): assert agg["joules_per_total_token"] == round(84000 / 36864, 6) assert agg["prefill_gpu_energy_j"] == 48000.0 assert agg["decode_gpu_energy_j"] == 36000.0 + assert agg["prefill_avg_power_w"] == 400.0 + assert agg["decode_avg_power_w"] == 300.0 assert agg["prefill_joules_per_input_token"] == round(48000 / 32768, 6) assert agg["decode_joules_per_output_token"] == round(36000 / 4096, 6) @@ -258,6 +260,23 @@ def test_emits_all_metrics_exactly(self, tmp_path): } assert set(sidecar["per_gpu_energy_j"]) == set(sidecar["per_gpu_role"]) + def test_role_watts_close_over_the_whole_deployment(self, tmp_path): + """Role watts weighted by their GPU counts must reproduce the whole- + deployment watts -- the topology gate makes the role partition + exhaustive. Catches dividing by the total device count.""" + + def ramp(host, idx, ts): + if (host, idx) == ("node-d", 0): + return 300.0 + (ts - FIRST_TS) + return dict(((h, i), w) for h, i, _r, _g, w in DEVICES)[(host, idx)] + + pkg = build_package(tmp_path, power_fn=ramp) + assert pkg.run() == 0 + agg = pkg.agg() + weighted = 2 * agg["prefill_avg_power_w"] + 2 * agg["decode_avg_power_w"] + assert weighted == pytest.approx(agg["avg_total_gpu_power_w"], abs=1e-3) + assert weighted / 4 == pytest.approx(agg["avg_power_w"], abs=1e-3) + def test_strict_mode_passes_on_valid_package(self, tmp_path): pkg = build_package(tmp_path) assert pkg.run(require_power=True) == 0 @@ -444,6 +463,8 @@ def test_none_het_groups_valid_for_non_het_deployments(self, tmp_path): assert agg["power_valid"] == 1 assert agg["prefill_gpu_energy_j"] == 48000.0 assert agg["decode_gpu_energy_j"] == 36000.0 + assert agg["prefill_avg_power_w"] == 400.0 + assert agg["decode_avg_power_w"] == 300.0 def test_mixed_none_and_real_het_groups_rejected(self, tmp_path): pkg = build_package(tmp_path) diff --git a/utils/test_process_result.py b/utils/test_process_result.py index 1142a117f..4d5219010 100644 --- a/utils/test_process_result.py +++ b/utils/test_process_result.py @@ -1224,6 +1224,8 @@ def test_valid_package_patches_role_energy(self, tmp_path, power_env): assert agg["power_valid"] == 1 assert agg["prefill_gpu_energy_j"] == 48000.0 assert agg["decode_gpu_energy_j"] == 36000.0 + assert agg["prefill_avg_power_w"] == 400.0 + assert agg["decode_avg_power_w"] == 300.0 assert (tmp_path / "power_validation_benchmark_result.json").is_file() def test_missing_package_is_best_effort(self, tmp_path, power_env): From c62f6e335bbe041a862670ff1d1fa747cdcecab9 Mon Sep 17 00:00:00 2001 From: Yifan Qiao Date: Tue, 11 Aug 2026 09:58:00 -0700 Subject: [PATCH 13/18] =?UTF-8?q?Retune=20DSV4=20B300=20AgentX=20MTP=20swe?= =?UTF-8?q?ep=20/=20=E8=B0=83=E4=BC=98=20DSV4=20B300=20AgentX=20MTP=20?= =?UTF-8?q?=E6=89=AB=E6=8F=8F=20(#2536)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * perf(agentx): retune DSV4 B300 MTP sweep Update the vLLM image and concurrency grid, raise the DEP4 host-memory budget, and set long-prefill-token-threshold to 512 for DEP4 and DEP8. 中文：更新 vLLM 镜像与并发配置，提高 DEP4 主机内存预算，并为 DEP4 和 DEP8 将 long-prefill-token-threshold 设为 512。 * docs(changelog): link DSV4 B300 MTP PR Replace the temporary changelog placeholder with draft PR #2536. 中文：将变更日志中的临时占位符替换为草稿 PR #2536 的链接。 * perf(agentx): extend DEP8 SimpleCPU sweep Move every DEP8 point to vLLM SimpleCPU offload, widen the concurrency grid through c512, and drop DEP4 c40. 中文：将全部 DEP8 配置切换到 vLLM SimpleCPU 卸载，将并发扫描扩展至 c512，并移除 DEP4 c40。 * perf(agentx): finalize SimpleCPU sweep grid Move TP4 c16 to SimpleCPU offload, skip TP4 c12 and DEP8 c64, and extend the DEP8 sweep through c576. 中文：将 TP4 c16 切换到 SimpleCPU 卸载，移除 TP4 c12 与 DEP8 c64，并将 DEP8 扫描扩展至 c576。 * Update perf-changelog.yaml --------- Co-authored-by: Cameron Quilici --- .../agentic/dsv4_fp4_b300_vllm_mtp.sh | 28 +++++++++---------- configs/nvidia-master.yaml | 14 +++++----- perf-changelog.yaml | 7 +++++ 3 files changed, 28 insertions(+), 21 deletions(-) diff --git a/benchmarks/single_node/agentic/dsv4_fp4_b300_vllm_mtp.sh b/benchmarks/single_node/agentic/dsv4_fp4_b300_vllm_mtp.sh index 9fdafe831..7aa4384cf 100755 --- a/benchmarks/single_node/agentic/dsv4_fp4_b300_vllm_mtp.sh +++ b/benchmarks/single_node/agentic/dsv4_fp4_b300_vllm_mtp.sh @@ -18,8 +18,8 @@ set -x # Required env vars: # MODEL, TP, CONC, KV_OFFLOADING, TOTAL_CPU_DRAM_GB, RESULT_DIR # -# TP4, TP8, and DEP8 (TP8 + DP-attention) are GPU-resident (KV_OFFLOADING=none). -# DEP4 uses KV_OFFLOADING=dram with KV_OFFLOAD_BACKEND=vllm-simple or mooncake. +# TP8 and TP4 c8 are GPU-resident. TP4 c16, DEP4, and DEP8 use DRAM offload +# with KV_OFFLOAD_BACKEND=vllm-simple or mooncake. source "$(dirname "$0")/../../benchmark_lib.sh" @@ -49,9 +49,9 @@ if [ "$DP_ATTENTION" = "true" ] && [ $((2 * CONC % TP)) -ne 0 ]; then exit 1 fi -# DEP8 (TP8 + DP-attention) is a GPU-resident, high-concurrency arm that is -# tuned separately from the smaller DEP4 arm (larger prefill token budget, -# long-prefill chunking, and a lower GPU-memory-utilization headroom). +# DEP8 (TP8 + DP-attention) is a high-concurrency SimpleCPU arm tuned separately +# from DEP4 with a larger prefill token budget and lower GPU-memory-utilization +# headroom. Both DEP arms chunk long prefills. IS_DEP8=false if [ "$DP_ATTENTION" = "true" ] && [ "$TP" -eq 8 ]; then IS_DEP8=true @@ -119,8 +119,8 @@ ROUTER_PID="" MOONCAKE_MASTER_PID="" # The generated TOTAL_CPU_DRAM_GB budget is proportional to allocated GPUs. -# On cluster:b300-nv, dram-utilization=0.80 and DEP4 resolve to roughly the -# source recipe's 280 GiB per DP rank. TP4 remains GPU-resident. +# On cluster:b300-nv, dram-utilization=0.95 gives both DEP4 and DEP8 356 GB per +# DP rank (1,424 GB and 2,849 GB total, respectively). TP arms remain GPU-resident. OFFLOAD_ARGS=() case "$KV_OFFLOAD_BACKEND" in "") @@ -238,14 +238,14 @@ if [ "$EP_SIZE" -gt 1 ]; then ) fi if [ "$DP_ATTENTION" = "true" ]; then - MODE_ARGS+=(--prefill-schedule-interval 8) + MODE_ARGS+=( + --prefill-schedule-interval 8 + --long-prefill-token-threshold 512 + ) if [ "$IS_DEP8" = "true" ]; then - # GPU-resident DEP8 gets a larger prefill token budget and chunks long - # prefills so decode latency stays bounded at high concurrency. - MODE_ARGS+=( - --max-num-batched-tokens 16384 - --long-prefill-token-threshold 4096 - ) + # DEP8 gets a larger prefill token budget; the shared long-prefill + # threshold keeps decode latency bounded under load. + MODE_ARGS+=(--max-num-batched-tokens 16384) else MODE_ARGS+=(--max-num-batched-tokens 8192) fi diff --git a/configs/nvidia-master.yaml b/configs/nvidia-master.yaml index 263dec901..f916be18d 100644 --- a/configs/nvidia-master.yaml +++ b/configs/nvidia-master.yaml @@ -1725,7 +1725,7 @@ dsv4-fp4-b300-vllm: - { tp: 8, ep: 8, dp-attn: true, conc-start: 2048, conc-end: 2048 } dsv4-fp4-b300-vllm-agentic-mtp: - image: vllm/vllm-openai:nightly-dev-x86_64-cu13.0.1-904e4ec + image: vllm/vllm-openai:nightly-dev-x86_64-cu13.0.1-426e59f model: deepseek-ai/DeepSeek-V4-Pro model-prefix: dsv4 runner: cluster:b300-nv @@ -1734,18 +1734,18 @@ dsv4-fp4-b300-vllm-agentic-mtp: multinode: false scenarios: agentic-coding: - - dram-utilization: 0.80 + - dram-utilization: 0.95 search-space: # TP8 GPU-resident + MTP (num_speculative_tokens=3) - - { tp: 8, kv-offloading: none, spec-decoding: mtp, conc-list: [1, 2, 4, 6, 8] } + - { tp: 8, kv-offloading: none, spec-decoding: mtp, conc-list: [1, 4] } # TP4 GPU-resident + MTP (num_speculative_tokens=3) - - { tp: 4, kv-offloading: none, spec-decoding: mtp, conc-list: [1, 2, 4, 6, 8, 12, 16, 20] } + - { tp: 4, kv-offloading: none, spec-decoding: mtp, conc-list: [8] } # TP4 SimpleCPU + MTP (num_speculative_tokens=3) - - { tp: 4, kv-offloading: dram, kv-offload-backend: { name: vllm-simple }, spec-decoding: mtp, conc-list: [20, 24, 28, 32, 36, 40] } + - { tp: 4, kv-offloading: dram, kv-offload-backend: { name: vllm-simple }, spec-decoding: mtp, conc-list: [16] } # DEP4 SimpleCPU + MTP (num_speculative_tokens=3) - - { tp: 4, ep: 4, dp-attn: true, kv-offloading: dram, kv-offload-backend: { name: vllm-simple }, spec-decoding: mtp, conc-list: [32, 40, 48, 56], router: { name: vllm-router, version: "0.1.14" } } + - { tp: 4, ep: 4, dp-attn: true, kv-offloading: dram, kv-offload-backend: { name: vllm-simple }, spec-decoding: mtp, conc-list: [48, 64], router: { name: vllm-router, version: "0.1.14" } } # DEP8 SimpleCPU + MTP (num_speculative_tokens=3) - - { tp: 8, ep: 8, dp-attn: true, kv-offloading: none, spec-decoding: mtp, conc-list: [64, 96, 112, 128, 144, 160, 176, 192, 224], router: { name: vllm-router, version: "0.1.14" } } + - { tp: 8, ep: 8, dp-attn: true, kv-offloading: dram, kv-offload-backend: { name: vllm-simple }, spec-decoding: mtp, conc-list: [128, 256, 384, 512, 576], router: { name: vllm-router, version: "0.1.14" } } dsv4-fp4-b300-trt: image: ghcr.io#semianalysisai/trtllm-deepseek-v4:feat-deepseek_v4-c185066 diff --git a/perf-changelog.yaml b/perf-changelog.yaml index 64f87b447..b89b34e5f 100644 --- a/perf-changelog.yaml +++ b/perf-changelog.yaml @@ -5770,3 +5770,10 @@ - "models.yaml: add mtp_flags (--speculative-algorithm EAGLE --speculative-eagle-topk 1) to DeepSeek-V4-Pro-AgentX; num-steps/num-draft-tokens are derived from DECODE_MTP_SIZE in server_sglang.sh. DeepSeek-V4's MTP head is a native NextN/EAGLE draft shipped with the model, so no --speculative-draft-model-path is needed." - "Two search-space arms: TP8/EP1 no-DP at conc-list [2,4,8,16,32], and TP8/EP8/DPA (ep=8, dp-attn=true) at conc-list [64,96,128], both on image lmsysorg/sglang-rocm:v0.5.15.post1-rocm720-mi35x-20260719." pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2309 + +- config-keys: + - dsv4-fp4-b300-vllm-agentic-mtp + description: + - "Update the B300 DSV4 AgentX MTP image and SimpleCPU sweep through c576." + pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2536 + From 66c9caa0d7f8364add0a18054b9e8473c8459fbe Mon Sep 17 00:00:00 2001 From: Cam Quilici Date: Tue, 11 Aug 2026 13:48:17 -0500 Subject: [PATCH 14/18] docs(amd): record MI355X MoRI fabric contract --- .../srt-slurm-recipes/AMD_BRINGUP.md | 25 ++++++++++++++++--- 1 file changed, 22 insertions(+), 3 deletions(-) diff --git a/benchmarks/multi_node/srt-slurm-recipes/AMD_BRINGUP.md b/benchmarks/multi_node/srt-slurm-recipes/AMD_BRINGUP.md index f8548d368..ba6f023df 100644 --- a/benchmarks/multi_node/srt-slurm-recipes/AMD_BRINGUP.md +++ b/benchmarks/multi_node/srt-slurm-recipes/AMD_BRINGUP.md @@ -5,11 +5,15 @@ This document tracks the work-in-progress integration of InferenceX AMD Slurm clusters. The project is a functional orchestration bring-up, not a performance-tuning exercise. -Current development pin for both AMD launchers: +Current development pins: - repository: `SemiAnalysisAI/srt-slurm` -- branch: `agent/amd-multinode-runtime` -- commit: `315e4b06a7e0806194a646ea21832e750e896a46` +- MI300X commit: `dd0109d4043141072ad37c043f1100332008b77f` +- MI355X commit: `5c0d5a718194198447e8ceea5c37048c2fb8eaf5` + +The MI355X pin includes the SGLang-router readiness gate: srt-slurm waits for +every advertised prefill and decode HTTP health endpoint before starting the +model gateway. The MI300X launcher uses srt-slurm's supported `--no-preflight` submission mode because the immutable squashfs files live on compute-node-local RAID rather @@ -58,6 +62,21 @@ The launcher will continue to exclude compute nodes already documented as unsuitable. It must not resume down nodes, cancel or preempt existing jobs, or alter unrelated shared software. +## MI355X MoRI fabric contract + +The MI355X DeepSeek V4 disaggregated path uses MoRI with DSCP traffic class +104. Each Ionic data NIC must therefore classify DSCP 26 as priority 3, enable +lossless PFC for priority 3, and keep the port in PFC pause mode with RX and TX +pause enabled. The cluster's management traffic additionally maps DSCP 48 to +strict priority 6. The expected scheduler split is 10% priority 0, 90% +priority 3, and a 10 Gbps rate limit for strict priority 6. + +Before admitting a node to a MoRI validation run, verify all eight NICs with +`nicctl show qos` and `nicctl show port`, and verify all eight RDMA links report +`ACTIVE` / `LINK_UP`. A node that does not match this contract must be drained +before repair and resumed only after every NIC and RDMA link passes the same +checks. Do not alter QoS while a Slurm allocation is using the node. + ## Acceptance criteria ### Aggregate From 2ec7a69d84850a53c0faa86d3e952309ec7c490c Mon Sep 17 00:00:00 2001 From: Cam Quilici Date: Tue, 11 Aug 2026 13:50:01 -0500 Subject: [PATCH 15/18] docs(amd): include MoRI fabric contract in changelog --- perf-changelog.yaml | 1 + 1 file changed, 1 insertion(+) diff --git a/perf-changelog.yaml b/perf-changelog.yaml index 9eb46ae98..c2a3bae59 100644 --- a/perf-changelog.yaml +++ b/perf-changelog.yaml @@ -5800,6 +5800,7 @@ - "Move the known-green DeepSeek V4 ROCm kernel-routing contract into each self-contained srt-slurm recipe; this avoids the unsupported AITer fused-MoE heuristic dispatch path while keeping legacy harness-only environment indirection removed." - "Declare the exact SGLang and AMD MoRI runtime fingerprint reported by the pinned ROCm image so srt-slurm identity verification is reproducible." - "Rebase the AMD runtime on the active native-router branch and wait for every advertised SGLang HTTP worker to become healthy before launching Model Gateway, preventing slow ROCm startup from permanently dropping a static prefill worker after the router's one-shot registration timeout." + - "Document the MI355X Ionic DSCP/PFC and RDMA health contract required by MoRI, including the drain-before-repair and verify-before-resume procedure." pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2555 - config-keys: From ee00360bcc170d3eb0ae4f2f306b2b93236f99f2 Mon Sep 17 00:00:00 2001 From: Cam Quilici Date: Tue, 11 Aug 2026 13:52:28 -0500 Subject: [PATCH 16/18] docs(amd): record exact MI355X srt runtime pin --- perf-changelog.yaml | 1 + 1 file changed, 1 insertion(+) diff --git a/perf-changelog.yaml b/perf-changelog.yaml index c2a3bae59..8c79e3c1f 100644 --- a/perf-changelog.yaml +++ b/perf-changelog.yaml @@ -5800,6 +5800,7 @@ - "Move the known-green DeepSeek V4 ROCm kernel-routing contract into each self-contained srt-slurm recipe; this avoids the unsupported AITer fused-MoE heuristic dispatch path while keeping legacy harness-only environment indirection removed." - "Declare the exact SGLang and AMD MoRI runtime fingerprint reported by the pinned ROCm image so srt-slurm identity verification is reproducible." - "Rebase the AMD runtime on the active native-router branch and wait for every advertised SGLang HTTP worker to become healthy before launching Model Gateway, preventing slow ROCm startup from permanently dropping a static prefill worker after the router's one-shot registration timeout." + - "Pin the MI355X launcher to SemiAnalysisAI/srt-slurm commit 5c0d5a718194198447e8ceea5c37048c2fb8eaf5, including the native-router base and SGLang worker-readiness gate." - "Document the MI355X Ionic DSCP/PFC and RDMA health contract required by MoRI, including the drain-before-repair and verify-before-resume procedure." pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2555 From a5ec626d4dbb3991c9903d2289c841c06351f987 Mon Sep 17 00:00:00 2001 From: giovanniguastiamd Date: Tue, 11 Aug 2026 21:30:10 +0200 Subject: [PATCH 17/18] [AMD] [AGENTX] GLM-5.2 FP4 MI355X SGLang Agentic MTP (#2488) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * [AMD] [WIP] [AGENTX] Adding GLM5.2 MI355X Support - Agentx Signed-off-by: Giovanni Guasti * [AMD] [WIP] [AGENTX] Adding GLM5.2 MI355X Support - Agentx Signed-off-by: Giovanni Guasti * [AMD] [WIP] [AGENTX] Adding GLM5.2 MI355X Support - Agentx Signed-off-by: Giovanni Guasti * chore: update pr-link for glm5.2-fp4-mi355x-sglang-agentic-mtp 中文：更新 glm5.2-fp4-mi355x-sglang-agentic-mtp 配置的 pr-link 至 PR #2488 * [AMD] [WIP] [AGENTX] GLM 5.2 - Update glm5.2_fp4_mi355x_sglang_mtp.sh * [AMD] [WIP] [AGENTX] GLM 5.2 - Update perf-changelog.yaml * [AMD] [WIP] [AGENTX] GLM 5.2 - Update glm5.2_fp4_mi355x_sglang_mtp.sh * fix(glm5.2-fp4-mi355x-sglang-mtp): remove trailing backslashes from SGLANG_CMD array Backslashes inside bash arrays are treated as literal characters, causing SGLang to receive arguments like 'EAGLE \' instead of 'EAGLE', which results in 'unrecognized arguments' errors. 中文：修复 SGLANG_CMD 数组中多余的行继续符（反斜杠），避免 SGLang 将其作为字面参数解析导致启动失败 * fix(glm5.2-fp4-mi355x-sglang-mtp): increase watchdog-timeout to 3600s MTP draft-token forward passes under long-context agentic load block the SGLang scheduler long enough to trip the 1800s watchdog, killing the server mid-warmup. Matches the value used by multi-node AMD disagg recipes. 中文：将 watchdog-timeout 从 1800s 增加至 3600s，避免 MTP 投机解码在长上下文推理负载下触发 SGLang 调度器看门狗超时导致服务崩溃 * [AMD] [WIP] [AGENTX] GLM 5.2 - Update glm5.2_fp4_mi355x_sglang_mtp.sh * fix(glm5.2-fp4-mi355x-sglang-agentic-mtp): set mem-fraction-static 0.80 for conc<=16 / 修正低并发场景 mem-fraction-static 参数为 0.80 0.85 OOMs the device mid-replay on 131072-token prefill chunks (observed: "Tried to allocate 6.86 GiB ... 5.15 GiB is free", run 29751563205). The comment already documented 0.80 as the correct value; the assignment was wrong. Co-Authored-By: Claude Sonnet 4.6 * [AMD][AgentX] glm5.2 fp4 mi355x sglang mtp: set SGLANG_SIMULATE_ACC_LEN=2.99 for non-eval runs Co-Authored-By: Claude Opus 4.6 * [AMD] [AGENTX] GLM 5.2 - Update glm5.2_fp4_mi355x_sglang_mtp.sh * [AMD] [AGENTX] GLM 5.2 - Update amd-master.yaml * [AMD] [AGENTX] GLM 5.2 Update Search Space --------- Signed-off-by: Giovanni Guasti Co-authored-by: ajith-sirra-amd <122240613+ajith-sirra-amd@users.noreply.github.com> Co-authored-by: Claude Sonnet 4.6 Co-authored-by: seungrokj <144636725+seungrokj@users.noreply.github.com> Co-authored-by: seungrokj --- .../agentic/glm5.2_fp4_mi355x_sglang_mtp.sh | 272 ++++++++++++++++++ configs/amd-master.yaml | 19 ++ perf-changelog.yaml | 7 + 3 files changed, 298 insertions(+) create mode 100644 benchmarks/single_node/agentic/glm5.2_fp4_mi355x_sglang_mtp.sh diff --git a/benchmarks/single_node/agentic/glm5.2_fp4_mi355x_sglang_mtp.sh b/benchmarks/single_node/agentic/glm5.2_fp4_mi355x_sglang_mtp.sh new file mode 100644 index 000000000..d7ed7a359 --- /dev/null +++ b/benchmarks/single_node/agentic/glm5.2_fp4_mi355x_sglang_mtp.sh @@ -0,0 +1,272 @@ +#!/usr/bin/env bash +set -eo pipefail +set -x + +source "$(dirname "$0")/../../benchmark_lib.sh" + + export EVAL_FRAMEWORK="lm-eval" + +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 + + +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 (host DRAM), optionally extended with Mooncake L3. +# KV_OFFLOADING=dram requires KV_OFFLOAD_BACKEND=hicache or mooncake. +# +# Per-arm L2 ratio (sizing rationale below) applies to both backends unless +# overridden via HICACHE_RATIO. 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 for the +# mooncake arm. 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. The hicache-only arm has no +# L3 to fall back on, so these ratios are unvalidated there - override with +# HICACHE_RATIO if the host OOMs or hit-rate is poor. +CACHE_ARGS=() +if agentic_kv_offload_enabled; 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}" + case "$KV_OFFLOAD_BACKEND" in + hicache) + echo "HiCache (GPU+host DRAM only): ratio=$HICACHE_RATIO, 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" + ) + ;; + mooncake) + 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 + ) + ;; + *) + echo "Error: unsupported KV_OFFLOAD_BACKEND '$KV_OFFLOAD_BACKEND' (expected: hicache or mooncake)" >&2 + exit 1 + ;; + esac +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=$((1 * CONC)) +[ "$MAX_RUNNING_REQUESTS" -gt 256 ] && MAX_RUNNING_REQUESTS=256 +CUDA_GRAPH_MAX_BS=$MAX_RUNNING_REQUESTS + +if [ "${EVAL_ONLY:-false}" != "true" ]; then + export SGLANG_SIMULATE_ACC_LEN=2.99 + export SGLANG_SIMULATE_ACC_METHOD=match-expected + export SGLANG_SIMULATE_ACC_TOKEN_MODE=real-draft-token +fi + +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[@]}" + --kv-cache-dtype fp8_e4m3 + --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" + --speculative-algorithm EAGLE + --speculative-num-steps 3 + --speculative-eagle-topk 1 + --speculative-num-draft-tokens 4 + "${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 + 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/configs/amd-master.yaml b/configs/amd-master.yaml index d26796a52..9bc224b16 100644 --- a/configs/amd-master.yaml +++ b/configs/amd-master.yaml @@ -1615,3 +1615,22 @@ dsv4-fp8-mi325x-vllm-mtp: # is 55.8%. conc128 passed cleanly on the memory-tightest SKU (MI300X); # cap 8k1k MTP at 128 (normal holds 256, 1k1k holds 512). - { tp: 8, conc-start: 4, conc-end: 128, spec-decoding: mtp } + +# GLM-5.2 FP4 agentic-coding benchmark on MI355X via SGLang with MTP speculative +# decoding. TP=4 EP=4 with KV offloading to DRAM (hicache backend) to support +# long agentic context windows. Concurrency sweep [1, 2, 4, 8, 10]. +glm5.2-fp4-mi355x-sglang-agentic-mtp: + image: lmsysorg/sglang-rocm:v0.5.16-rocm720-mi35x-20260728 + 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.8 + search-space: + - { tp: 4, ep: 4, kv-offloading: dram, kv-offload-backend: { name: hicache }, conc-list: [1, 2, 4, 8, 10, 12, 16], spec-decoding: mtp } + - { tp: 8, ep: 8, kv-offloading: dram, kv-offload-backend: { name: hicache }, conc-list: [1, 2, 4], spec-decoding: mtp } + diff --git a/perf-changelog.yaml b/perf-changelog.yaml index b89b34e5f..41e7238f8 100644 --- a/perf-changelog.yaml +++ b/perf-changelog.yaml @@ -5726,6 +5726,13 @@ - "Use supported header-based Dynamo session routing with the in-repo AIPerf build." pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2520 +- config-keys: + - glm5.2-fp4-mi355x-sglang-agentic-mtp + description: + - "Add GLM-5.2 FP4 SGLANG Single Node Agentic Support" + - "Image: lmsysorg/sglang-rocm:v0.5.16-rocm720-mi35x-20260728" + pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2488 + - config-keys: - dsv4-fp4-gb300-dynamo-sglang-agentic-agg - dsv4-fp4-gb300-dynamo-sglang-agentic-disagg From 5cb444d1c57761a34b26a118b1dded3f900d79c6 Mon Sep 17 00:00:00 2001 From: Cam Quilici Date: Tue, 11 Aug 2026 14:51:22 -0500 Subject: [PATCH 18/18] Pin MI355X runtime frontend fix --- benchmarks/multi_node/srt-slurm-recipes/AMD_BRINGUP.md | 2 +- perf-changelog.yaml | 2 +- runners/launch_mi355x-amds-srt.sh | 2 +- 3 files changed, 3 insertions(+), 3 deletions(-) diff --git a/benchmarks/multi_node/srt-slurm-recipes/AMD_BRINGUP.md b/benchmarks/multi_node/srt-slurm-recipes/AMD_BRINGUP.md index ba6f023df..cf4e1a976 100644 --- a/benchmarks/multi_node/srt-slurm-recipes/AMD_BRINGUP.md +++ b/benchmarks/multi_node/srt-slurm-recipes/AMD_BRINGUP.md @@ -9,7 +9,7 @@ Current development pins: - repository: `SemiAnalysisAI/srt-slurm` - MI300X commit: `dd0109d4043141072ad37c043f1100332008b77f` -- MI355X commit: `5c0d5a718194198447e8ceea5c37048c2fb8eaf5` +- MI355X commit: `c609754b5622f96d5c12a93149e245308d4f1e9b` The MI355X pin includes the SGLang-router readiness gate: srt-slurm waits for every advertised prefill and decode HTTP health endpoint before starting the diff --git a/perf-changelog.yaml b/perf-changelog.yaml index 8c79e3c1f..3b1285f3c 100644 --- a/perf-changelog.yaml +++ b/perf-changelog.yaml @@ -5800,7 +5800,7 @@ - "Move the known-green DeepSeek V4 ROCm kernel-routing contract into each self-contained srt-slurm recipe; this avoids the unsupported AITer fused-MoE heuristic dispatch path while keeping legacy harness-only environment indirection removed." - "Declare the exact SGLang and AMD MoRI runtime fingerprint reported by the pinned ROCm image so srt-slurm identity verification is reproducible." - "Rebase the AMD runtime on the active native-router branch and wait for every advertised SGLang HTTP worker to become healthy before launching Model Gateway, preventing slow ROCm startup from permanently dropping a static prefill worker after the router's one-shot registration timeout." - - "Pin the MI355X launcher to SemiAnalysisAI/srt-slurm commit 5c0d5a718194198447e8ceea5c37048c2fb8eaf5, including the native-router base and SGLang worker-readiness gate." + - "Pin the MI355X launcher to SemiAnalysisAI/srt-slurm commit c609754b5622f96d5c12a93149e245308d4f1e9b, including the native-router base, SGLang worker-readiness gate, and benchmark-stage frontend registry fix." - "Document the MI355X Ionic DSCP/PFC and RDMA health contract required by MoRI, including the drain-before-repair and verify-before-resume procedure." pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2555 diff --git a/runners/launch_mi355x-amds-srt.sh b/runners/launch_mi355x-amds-srt.sh index 4669cc0cf..3e2d56f4a 100755 --- a/runners/launch_mi355x-amds-srt.sh +++ b/runners/launch_mi355x-amds-srt.sh @@ -5,7 +5,7 @@ set -euo pipefail # in explicitly with CONFIG_FILE; all existing MI355X launch behavior remains # unchanged for every other row. SRT_SLURM_REPOSITORY="https://github.com/SemiAnalysisAI/srt-slurm.git" -SRT_SLURM_COMMIT="5c0d5a718194198447e8ceea5c37048c2fb8eaf5" +SRT_SLURM_COMMIT="c609754b5622f96d5c12a93149e245308d4f1e9b" SLURM_PARTITION="compute" SHARED_BASE="/it-share/gharunners2/srt-slurm" SHARED_HF_CACHE="/it-share/hf-hub-cache"