fix(pd-disagg): Fix vectorized_5d Mooncake KV Transfer

evaluation/mimo-v2.5-pro-evaluation-guide.md

@@ -120,3 +120,6 @@ Follow up customer's swe-bench accuracy verification guide.
 
 ## 10. H200 performance evaluation
 Follow up customer's shared performance data
+
+## 11. Run prefill/decode disaggregated deployment
+Please refer to the [PD disaggregated guide](mimo_pd_scripts/mimo-v2.5-pro-pd-disaggregated-guide.md).

evaluation/mimo_pd_scripts/launch_router.sh

@@ -0,0 +1,24 @@
+#!/bin/bash
+set -euo pipefail
+# launch_router.sh — SGLang PD router
+#
+# Run this on the PREFILL node after both P and D servers are ready.
+
+# Defaults are the verified host IPs.
+# Override with environment variables when running on a different network.
+PREFILL_HOST_IP="${PREFILL_HOST_IP:-172.16.1.26}"
+DECODE_HOST_IP="${DECODE_HOST_IP:-172.16.1.122}"
+LOG_DIR="${LOG_DIR:-logs/pd-disagg-router}"
+
+mkdir -p "$LOG_DIR"
+
+echo "=== Launching SGLang PD Router ==="
+echo "Prefill: http://$PREFILL_HOST_IP:30000"
+echo "Decode:  http://$DECODE_HOST_IP:30001"
+
+python3 -m sglang_router.launch_router \
+  --pd-disaggregation \
+  --prefill "http://$PREFILL_HOST_IP:30000" \
+  --decode "http://$DECODE_HOST_IP:30001" \
+  --host 0.0.0.0 --port 40000 \
+  2>&1 | tee "$LOG_DIR/router.log"

evaluation/mimo_pd_scripts/launch_tp8_noep_decode_aiter_mtp.sh

@@ -0,0 +1,127 @@
+#!/usr/bin/env bash
+set -euo pipefail
+# MiMo-v2.5-Pro SGLang 1P1D disaggregated decode node.
+
+SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
+# KV/attention settings must match prefill
+# (page 64, kv auto, vectorized_5d). CUDA graph stays on for decode.
+
+export SGLANG_USE_AITER=1
+export SGLANG_MOE_PADDING=1
+export SGLANG_ROCM_FUSED_DECODE_MLA=1
+export SGLANG_SET_CPU_AFFINITY=1
+export HSA_NO_SCRATCH_RECLAIM=1
+
+# Native InfiniBand use GID 0 is valid. GID 1 used for RoCE.
+export SGLANG_HOST_IP="${SGLANG_HOST_IP:-172.16.1.122}"
+export HOST_IP="${HOST_IP:-${SGLANG_HOST_IP}}"
+export MC_GID_INDEX="${MC_GID_INDEX:-1}"
+#export DISAGGREGATION_IB_DEVICE="${DISAGGREGATION_IB_DEVICE:-mlx5_ib0,mlx5_ib1,mlx5_ib2,mlx5_ib3,mlx5_ib4,mlx5_ib5,mlx5_ib6,mlx5_ib7}"
+export DISAGGREGATION_IB_DEVICE="${DISAGGREGATION_IB_DEVICE:-rocep121s0,rocep9s0,rocep105s0,rocep25s0,rocep249s0,rocep137s0,rocep233s0,rocep153s0}"
+export MC_TE_METRIC=1
+
+export SGLANG_DISAGGREGATION_THREAD_POOL_SIZE=12
+export SGLANG_DISAGGREGATION_BOOTSTRAP_TIMEOUT=5000
+export SGLANG_DISAGGREGATION_WAITING_TIMEOUT=5000
+export SGLANG_ENABLE_SPEC_V2=1
+export SGLANG_SPEC_NAN_DETECTION=1
+export SGLANG_SPEC_OOB_DETECTION=1
+export SGLANG_MIMO_EAGLE_HIP_NONGREEDY_VERIFY="${SGLANG_MIMO_EAGLE_HIP_NONGREEDY_VERIFY:-1}"
+export SGLANG_USE_AITER_CK_BLOCKSCALE_BPRESHUFFLE=1
+
+if [[ "${SGLANG_ABLATION_ENABLE_QUICK_REDUCE:-0}" == "1" ]]; then
+  export ROCM_QUICK_REDUCE_QUANTIZATION="${ROCM_QUICK_REDUCE_QUANTIZATION:-INT8}"
+  echo "ROCM_QUICK_REDUCE_QUANTIZATION enabled for ablation: ${ROCM_QUICK_REDUCE_QUANTIZATION}"
+elif [[ -v ROCM_QUICK_REDUCE_QUANTIZATION ]]; then
+  echo "ROCM_QUICK_REDUCE_QUANTIZATION was set to '${ROCM_QUICK_REDUCE_QUANTIZATION}', unsetting for accuracy"
+  unset ROCM_QUICK_REDUCE_QUANTIZATION
+else
+  echo "ROCM_QUICK_REDUCE_QUANTIZATION is unset"
+fi
+
+export SGLANG_MIMO_MIXED_ROUTER="${SGLANG_MIMO_MIXED_ROUTER:-0}"
+export SGLANG_FLYPA_MIMO_PREFILL="${SGLANG_FLYPA_MIMO_PREFILL:-0}"
+export TORCH_ALLOW_TF32_CUBLAS_OVERRIDE=1
+export SGLANG_AITER_KV_CACHE_LAYOUT=vectorized_5d
+export SGLANG_FLYDSL_MIMO_PREFILL="${SGLANG_FLYDSL_MIMO_PREFILL:-0}"
+export SGLANG_AITER_PA_DECODE_IMPL="${SGLANG_AITER_PA_DECODE_IMPL:-flydsl}"
+export SGLANG_FLYDSL_PA_NUM_PARTITIONS="${SGLANG_FLYDSL_PA_NUM_PARTITIONS:-16}"
+export CUDA_GRAPH_BACKEND_DECODE="${CUDA_GRAPH_BACKEND_DECODE:-full}"
+export CUDA_GRAPH_BS_DECODE="${CUDA_GRAPH_BS_DECODE:-}"
+export REASONING_PARSER="${REASONING_PARSER:-mimo}"
+export MAX_RUNNING_REQUESTS="${MAX_RUNNING_REQUESTS:-96}"
+export MEM_FRACTION_STATIC="${MEM_FRACTION_STATIC:-0.90}"
+export SWA_FULL_TOKENS_RATIO="${SWA_FULL_TOKENS_RATIO:-0.2}"
+export CHUNKED_PREFILL_SIZE="${CHUNKED_PREFILL_SIZE:-16384}"
+export KV_CACHE_DTYPE="${KV_CACHE_DTYPE:-auto}"
+export DISABLE_RADIX_CACHE="${DISABLE_RADIX_CACHE:-1}"
+
+export SGLANG_MIMO_FUSED_RMS_MOE_QUANT="${SGLANG_MIMO_FUSED_RMS_MOE_QUANT:-1}"
+export SGLANG_MIMO_FUSED_RMS_QKV_QUANT="${SGLANG_MIMO_FUSED_RMS_QKV_QUANT:-1}"
+export SGLANG_AITER_MIMO_FRESH_BF16_ASM="${SGLANG_AITER_MIMO_FRESH_BF16_ASM:-1}"
+export SGLANG_AITER_MIMO_FRESH_BF16_ASM_VARLEN="${SGLANG_AITER_MIMO_FRESH_BF16_ASM_VARLEN:-1}"
+export SGLANG_AITER_MIMO_FRESH_BF16_SWA_VARLEN="${SGLANG_AITER_MIMO_FRESH_BF16_SWA_VARLEN:-1}"
+
+unset SGLANG_SIMULATE_ACC_LEN SGLANG_SIMULATE_ACC_METHOD
+
+export MODEL="${MODEL:-/models/MiMo-V2.5-Pro}"
+export RUN_ID="${RUN_ID:-$(date -u +%Y%m%dT%H%M%SZ)}"
+export LOG_DIR="${LOG_DIR:-/root/workspace/bench_tp8_noep_mooncake}"
+export LOG_FILE="${LOG_FILE:-decode_server.log}"
+
+echo "Attention hybrid: prefill-flydsl=${SGLANG_FLYDSL_MIMO_PREFILL}, target-verify=${SGLANG_AITER_PA_DECODE_IMPL}, KV-layout=${SGLANG_AITER_KV_CACHE_LAYOUT}"
+echo "Configuration: max-running=${MAX_RUNNING_REQUESTS}, page=64, chunked-prefill=${CHUNKED_PREFILL_SIZE}, mem=${MEM_FRACTION_STATIC}, swa=${SWA_FULL_TOKENS_RATIO}, kv-cache-dtype=${KV_CACHE_DTYPE}"
+echo "PD decode: mooncake, ib-device=${DISAGGREGATION_IB_DEVICE}, gid-index=${MC_GID_INDEX}, host-ip=${SGLANG_HOST_IP}, topo=${MC_CUSTOM_TOPO_JSON}"
+echo "Server log: ${LOG_DIR}/${LOG_FILE}"
+echo "MTP: EAGLE, steps=3, top-k=1, draft-tokens=4, multi-layer=enabled"
+
+mkdir -p "${LOG_DIR}"
+
+cuda_graph_args=(--cuda-graph-backend-decode "${CUDA_GRAPH_BACKEND_DECODE}")
+if [[ -n "${CUDA_GRAPH_BS_DECODE}" ]]; then
+  read -r -a cuda_graph_bs <<< "${CUDA_GRAPH_BS_DECODE}"
+  cuda_graph_args+=(--cuda-graph-bs-decode "${cuda_graph_bs[@]}")
+fi
+
+kv_cache_args=()
+if [[ -n "${KV_CACHE_DTYPE}" && "${KV_CACHE_DTYPE}" != "auto" ]]; then
+  kv_cache_args+=(--kv-cache-dtype "${KV_CACHE_DTYPE}")
+fi
+
+radix_cache_args=()
+if [[ "${DISABLE_RADIX_CACHE}" == "1" ]]; then
+  radix_cache_args+=(--disable-radix-cache)
+  echo "Radix cache: disabled"
+else
+  echo "Radix cache: enabled"
+fi
+
+python3 -u -m sglang.launch_server \
+  --model-path "${MODEL}" \
+  --tp-size 8 \
+  --max-running-requests "${MAX_RUNNING_REQUESTS}" \
+  --host 0.0.0.0 \
+  --port 30001 \
+  --trust-remote-code \
+  --reasoning-parser "${REASONING_PARSER}" \
+  --tool-call-parser mimo \
+  --mem-fraction-static "${MEM_FRACTION_STATIC}" \
+  --swa-full-tokens-ratio "${SWA_FULL_TOKENS_RATIO}" \
+  --context-length 1048576 \
+  --chunked-prefill-size "${CHUNKED_PREFILL_SIZE}" \
+  --max-prefill-tokens 1048576 \
+  --attention-backend aiter \
+  "${kv_cache_args[@]}" \
+  --page-size 64 \
+  --disaggregation-mode decode \
+  --disaggregation-transfer-backend mooncake \
+  --disaggregation-ib-device "${DISAGGREGATION_IB_DEVICE}" \
+  --speculative-algorithm EAGLE \
+  --speculative-num-steps 3 \
+  --speculative-eagle-topk 1 \
+  --speculative-num-draft-tokens 4 \
+  --enable-multi-layer-eagle \
+  --watchdog-time 1200 \
+  "${radix_cache_args[@]}" \
+  "${cuda_graph_args[@]}" \
+  2>&1 | tee "${LOG_DIR}/${LOG_FILE}"

evaluation/mimo_pd_scripts/launch_tp8_noep_prefill_aiter_mtp.sh

@@ -0,0 +1,139 @@
+#!/usr/bin/env bash
+set -euo pipefail
+
+# MiMo-v2.5-Pro SGLang PD disaggregated prefill node.
+SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
+
+export SGLANG_USE_AITER=1
+export SGLANG_MOE_PADDING=1
+export SGLANG_SET_CPU_AFFINITY=1
+export HSA_NO_SCRATCH_RECLAIM=1
+
+# Native InfiniBand use GID 0 is valid. GID 1 used for RoCE.
+export SGLANG_HOST_IP="${SGLANG_HOST_IP:-172.16.1.26}"
+export HOST_IP="${HOST_IP:-${SGLANG_HOST_IP}}"
+export MC_GID_INDEX="${MC_GID_INDEX:-1}"
+#export DISAGGREGATION_IB_DEVICE="${DISAGGREGATION_IB_DEVICE:-mlx5_ib0,mlx5_ib1,mlx5_ib2,mlx5_ib3,mlx5_ib4,mlx5_ib5,mlx5_ib6,mlx5_ib7}"
+export DISAGGREGATION_IB_DEVICE="${DISAGGREGATION_IB_DEVICE:-rocep121s0,rocep9s0,rocep105s0,rocep25s0,rocep249s0,rocep137s0,rocep233s0,rocep153s0}"
+export MC_TE_METRIC=1
+
+export SGLANG_DISAGGREGATION_THREAD_POOL_SIZE=12
+export SGLANG_DISAGGREGATION_BOOTSTRAP_TIMEOUT=5000
+export SGLANG_DISAGGREGATION_WAITING_TIMEOUT=5000
+export SGLANG_ENABLE_SPEC_V2=1
+export SGLANG_SPEC_NAN_DETECTION=1
+export SGLANG_SPEC_OOB_DETECTION=1
+export SGLANG_MIMO_EAGLE_HIP_NONGREEDY_VERIFY="${SGLANG_MIMO_EAGLE_HIP_NONGREEDY_VERIFY:-1}"
+export SGLANG_USE_AITER_CK_BLOCKSCALE_BPRESHUFFLE=1
+
+if [[ "${SGLANG_ABLATION_ENABLE_QUICK_REDUCE:-0}" == "1" ]]; then
+  export ROCM_QUICK_REDUCE_QUANTIZATION="${ROCM_QUICK_REDUCE_QUANTIZATION:-INT8}"
+  echo "ROCM_QUICK_REDUCE_QUANTIZATION enabled for ablation: ${ROCM_QUICK_REDUCE_QUANTIZATION}"
+elif [[ -v ROCM_QUICK_REDUCE_QUANTIZATION ]]; then
+  echo "ROCM_QUICK_REDUCE_QUANTIZATION was set to '${ROCM_QUICK_REDUCE_QUANTIZATION}', unsetting for accuracy"
+  unset ROCM_QUICK_REDUCE_QUANTIZATION
+else
+  echo "ROCM_QUICK_REDUCE_QUANTIZATION is unset"
+fi
+
+export SGLANG_MIMO_MIXED_ROUTER="${SGLANG_MIMO_MIXED_ROUTER:-0}"
+export SGLANG_FLYPA_MIMO_PREFILL="${SGLANG_FLYPA_MIMO_PREFILL:-0}"
+export TORCH_ALLOW_TF32_CUBLAS_OVERRIDE=1
+# Required on this tree: native QK192/V128 + 5D pool. page-size 32 + fp8 hits
+# aiter_backend.view(-1, 32, 1, 192) and crashes.
+export SGLANG_AITER_KV_CACHE_LAYOUT=vectorized_5d
+export SGLANG_FLYDSL_MIMO_PREFILL="${SGLANG_FLYDSL_MIMO_PREFILL:-0}"
+export SGLANG_AITER_PA_DECODE_IMPL="${SGLANG_AITER_PA_DECODE_IMPL:-flydsl}"
+export SGLANG_FLYDSL_PA_NUM_PARTITIONS="${SGLANG_FLYDSL_PA_NUM_PARTITIONS:-16}"
+export CUDA_GRAPH_BACKEND_DECODE="${CUDA_GRAPH_BACKEND_DECODE:-full}"
+export CUDA_GRAPH_BS_DECODE="${CUDA_GRAPH_BS_DECODE:-}"
+export REASONING_PARSER="${REASONING_PARSER:-mimo}"
+export MAX_RUNNING_REQUESTS="${MAX_RUNNING_REQUESTS:-96}"
+export MEM_FRACTION_STATIC="${MEM_FRACTION_STATIC:-0.90}"
+# 0.01 overflowed SWA on 65k PD seqs; keep more headroom than the single-node default.
+export SWA_FULL_TOKENS_RATIO="${SWA_FULL_TOKENS_RATIO:-0.2}"
+export CHUNKED_PREFILL_SIZE="${CHUNKED_PREFILL_SIZE:-16384}"
+export KV_CACHE_DTYPE="${KV_CACHE_DTYPE:-auto}"
+export DISABLE_RADIX_CACHE="${DISABLE_RADIX_CACHE:-1}"
+export ENABLE_TWO_BATCH_OVERLAP="${ENABLE_TWO_BATCH_OVERLAP:-0}"
+export SGLANG_TBO_MIM_SEQ_LEN="${SGLANG_TBO_MIM_SEQ_LEN:-8000}"
+
+export SGLANG_MIMO_FUSED_RMS_MOE_QUANT="${SGLANG_MIMO_FUSED_RMS_MOE_QUANT:-1}"
+export SGLANG_MIMO_FUSED_RMS_QKV_QUANT="${SGLANG_MIMO_FUSED_RMS_QKV_QUANT:-1}"
+export SGLANG_AITER_MIMO_FRESH_BF16_ASM="${SGLANG_AITER_MIMO_FRESH_BF16_ASM:-1}"
+export SGLANG_AITER_MIMO_FRESH_BF16_ASM_VARLEN="${SGLANG_AITER_MIMO_FRESH_BF16_ASM_VARLEN:-1}"
+export SGLANG_AITER_MIMO_FRESH_BF16_SWA_VARLEN="${SGLANG_AITER_MIMO_FRESH_BF16_SWA_VARLEN:-1}"
+
+unset SGLANG_SIMULATE_ACC_LEN SGLANG_SIMULATE_ACC_METHOD
+
+export MODEL="${MODEL:-/models/MiMo-V2.5-Pro}"
+export RUN_ID="${RUN_ID:-$(date -u +%Y%m%dT%H%M%SZ)}"
+export LOG_DIR="${LOG_DIR:-/root/workspace/bench_tp8_noep_mooncake}"
+export LOG_FILE="${LOG_FILE:-prefill_server.log}"
+
+echo "Attention hybrid: prefill-flydsl=${SGLANG_FLYDSL_MIMO_PREFILL}, target-verify=${SGLANG_AITER_PA_DECODE_IMPL}, KV-layout=${SGLANG_AITER_KV_CACHE_LAYOUT}"
+echo "Configuration: max-running=${MAX_RUNNING_REQUESTS}, page=64, chunked-prefill=${CHUNKED_PREFILL_SIZE}, mem=${MEM_FRACTION_STATIC}, swa=${SWA_FULL_TOKENS_RATIO}, kv-cache-dtype=${KV_CACHE_DTYPE}"
+echo "PD prefill: mooncake, ib-device=${DISAGGREGATION_IB_DEVICE}, gid-index=${MC_GID_INDEX}, host-ip=${SGLANG_HOST_IP}, tbo=${ENABLE_TWO_BATCH_OVERLAP}"
+echo "Server log: ${LOG_DIR}/${LOG_FILE}"
+echo "MTP: EAGLE, steps=3, top-k=1, draft-tokens=4, multi-layer=enabled"
+
+mkdir -p "${LOG_DIR}"
+
+cuda_graph_args=(--cuda-graph-backend-decode "${CUDA_GRAPH_BACKEND_DECODE}")
+if [[ -n "${CUDA_GRAPH_BS_DECODE}" ]]; then
+  read -r -a cuda_graph_bs <<< "${CUDA_GRAPH_BS_DECODE}"
+  cuda_graph_args+=(--cuda-graph-bs-decode "${cuda_graph_bs[@]}")
+fi
+
+kv_cache_args=()
+if [[ -n "${KV_CACHE_DTYPE}" && "${KV_CACHE_DTYPE}" != "auto" ]]; then
+  kv_cache_args+=(--kv-cache-dtype "${KV_CACHE_DTYPE}")
+fi
+
+radix_cache_args=()
+if [[ "${DISABLE_RADIX_CACHE}" == "1" ]]; then
+  radix_cache_args+=(--disable-radix-cache)
+  echo "Radix cache: disabled"
+else
+  echo "Radix cache: enabled"
+fi
+
+tbo_args=()
+if [[ "${ENABLE_TWO_BATCH_OVERLAP}" == "1" ]]; then
+  tbo_args+=(--enable-two-batch-overlap)
+  echo "Two-batch overlap: enabled (min-isl=${SGLANG_TBO_MIM_SEQ_LEN})"
+else
+  echo "Two-batch overlap: disabled"
+fi
+
+python3 -u -m sglang.launch_server \
+  --model-path "${MODEL}" \
+  --tp-size 8 \
+  --max-running-requests "${MAX_RUNNING_REQUESTS}" \
+  --host 0.0.0.0 \
+  --port 30000 \
+  --trust-remote-code \
+  --reasoning-parser "${REASONING_PARSER}" \
+  --tool-call-parser mimo \
+  --mem-fraction-static "${MEM_FRACTION_STATIC}" \
+  --swa-full-tokens-ratio "${SWA_FULL_TOKENS_RATIO}" \
+  --context-length 1048576 \
+  --chunked-prefill-size "${CHUNKED_PREFILL_SIZE}" \
+  --max-prefill-tokens 1048576 \
+  --attention-backend aiter \
+  "${kv_cache_args[@]}" \
+  --page-size 64 \
+  --disaggregation-mode prefill \
+  --disaggregation-transfer-backend mooncake \
+  --disaggregation-ib-device "${DISAGGREGATION_IB_DEVICE}" \
+  --speculative-algorithm EAGLE \
+  --speculative-num-steps 3 \
+  --speculative-eagle-topk 1 \
+  --speculative-num-draft-tokens 4 \
+  --enable-multi-layer-eagle \
+  --watchdog-time 1200 \
+  --disable-cuda-graph \
+  "${tbo_args[@]}" \
+  "${radix_cache_args[@]}" \
+  "${cuda_graph_args[@]}" \
+  2>&1 | tee "${LOG_DIR}/${LOG_FILE}"

evaluation/mimo_pd_scripts/mimo-v2.5-pro-pd-disaggregated-guide.md

@@ -0,0 +1,98 @@
+## 1. Prepare docker image for gfx950 with AINIC
+### Prefill container on node 0
+```bash
+docker run -it --name mimo-sgl-opt-gfx950-prefill --ipc=host --network=host --privileged --security-opt seccomp=unconfined --cap-add=CAP_SYS_ADMIN --cap-add=IPC_LOCK --cap-add=SYS_PTRACE --device=/dev/kfd --device=/dev/dri --device=/dev/mem --device=/dev/infiniband -v $HOME:/root/workspace  -v /data/models:/models  -v /usr/lib/x86_64-linux-gnu/libibverbs.so.1.14.39.0:/lib/x86_64-linux-gnu/libibverbs.so.1 -v /usr/lib/x86_64-linux-gnu/libibverbs/libionic-rdmav34.so:/usr/lib/x86_64-linux-gnu/libibverbs/libionic-rdmav34.so -v /usr/lib/x86_64-linux-gnu/libionic.so:/usr/lib/x86_64-linux-gnu/libionic.so -v /usr/lib/x86_64-linux-gnu/libibverbs/libbnxt_re-rdmav34.so:/usr/lib/x86_64-linux-gnu/libibverbs/libbnxt_re-rdmav34.so -v /etc/libibverbs.d:/etc/libibverbs.d:ro rocm/sgl-dev:v0.5.11-rocm720-mi35x-20260510
+```
+
+### Decode container on node 1
+```bash
+docker run -it --name mimo-sgl-opt-gfx950-decode --ipc=host --network=host --privileged --security-opt seccomp=unconfined --cap-add=CAP_SYS_ADMIN --cap-add=IPC_LOCK --cap-add=SYS_PTRACE --device=/dev/kfd --device=/dev/dri --device=/dev/mem --device=/dev/infiniband -v $HOME:/root/workspace  -v /data/models:/models  -v /usr/lib/x86_64-linux-gnu/libibverbs.so.1.14.39.0:/lib/x86_64-linux-gnu/libibverbs.so.1 -v /usr/lib/x86_64-linux-gnu/libibverbs/libionic-rdmav34.so:/usr/lib/x86_64-linux-gnu/libibverbs/libionic-rdmav34.so -v /usr/lib/x86_64-linux-gnu/libionic.so:/usr/lib/x86_64-linux-gnu/libionic.so -v /usr/lib/x86_64-linux-gnu/libibverbs/libbnxt_re-rdmav34.so:/usr/lib/x86_64-linux-gnu/libibverbs/libbnxt_re-rdmav34.so -v /etc/libibverbs.d:/etc/libibverbs.d:ro rocm/sgl-dev:v0.5.11-rocm720-mi35x-20260510
+```
+## 2. Environment setup for both prefill container and decode container
+
+### Clean up pre-installed environment
+```bash
+pip uninstall sglang sglang-kernel sgl-kernel amd-aiter flydsl mimo-flydsl-kernels -y
+rm -rf /sgl-workspace/sglang /sgl-workspace/aiter
+```
+
+### Install sglang
+```bash
+cd /root/workspace
+git clone https://github.com/sammysun0711/sglang -b mimo-opt
+cd sglang && pip install --upgrade pip && cd sgl-kernel && python3 setup_rocm.py install
+cd ..  && rm -rf python/pyproject.toml && mv python/pyproject_other.toml python/pyproject.toml && pip install -e "python[all_hip]"
+cd ..
+```
+### Install AITER & pyhip dependency
+```bash
+git clone https://github.com/sammysun0711/aiter -b mimo-opt
+cd aiter
+git submodule update --init 3rdparty/composable_kernel
+git -C 3rdparty/composable_kernel apply < patches/composable_kernel/mimo_page64_qk192_v128_batch_prefill.patch
+pip install -e .
+cd ..
+```
+
+### Install FlyDSL
+FlyDSL installation should be after AITER installation, since it will update FlyDSL 0.2.4 runtime and FlyDSL FA & PA decode kernels.
+```bash
+git clone https://github.com/sammysun0711/FlyDSL -b mimo-opt
+cd FlyDSL/wheels && python3 -m pip install --no-deps --force-reinstall flydsl-0.2.4-*.whl mimo_flydsl_kernels-0.1.3-*.whl
+cd ../..
+```
+
+### Check RDMA status
+```bash
+rdma link show
+ibv_devinfo -v | head -30
+```
+
+### Setup passwordless ssh connection
+```bash
+git clone https://github.com/sammysun0711/llm-distributed-inference.git
+cd llm-distributed-inference/sglang/
+# On node 0, create remote ssh access in prefill container to node 1 decode container and vice versa
+```bash
+./scripts/setup_docker_passwdless_ssh.sh mi355-gpu-16 
+```
+checkout that node 0 and node 1 can access via ssh without password.
+```shell
+ssh mi355-gpu-16
+```
+
+## 3. Launch PD disagg prefill server on node 0
+Please note: update `SGLANG_HOST_IP`, `MC_GID_INDEX` and `DISAGGREGATION_IB_DEVICE` based on network settings.
+
+TBO on prefill node is disable by default, enable via `ENABLE_TWO_BATCH_OVERLAP=1`
+```bash
+cd /root/workspace/sglang/evaluation/mimo_pd_scripts
+./launch_tp8_noep_prefill_aiter_mtp.sh
+```
+
+## 4. Launch PD disagg decode server on node 1
+Please note: update `SGLANG_HOST_IP`, `MC_GID_INDEX` and `DISAGGREGATION_IB_DEVICE` based on network settings
+```bash
+cd /root/workspace/sglang/evaluation/mimo_pd_scripts
+./launch_tp8_noep_decode_aiter_mtp.sh
+```
+
+## 5. Launch PD disagg router on node 0
+Please note: update `PREFILL_HOST_IP` and `DECODE_HOST_IP` based on network settings
+```bash
+cd /root/workspace/sglang/evaluation/mimo_pd_scripts
+./launch_router.sh
+```
+
+## 6. Run PD disagg benchmark
+### Run prefill benchmark on node 0
+```bash
+cd /root/workspace/sglang/evaluation/mimo_pd_scripts
+./run_benchmark_mimo_pro_prefill.sh
+```
+
+### Run decode benchmark on node 0
+```bash
+cd /root/workspace/sglang/evaluation/mimo_pd_scripts
+./run_benchmark_mimo_pro_decode.sh
+```