Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
59 changes: 59 additions & 0 deletions benchmarks/multi_node/srt-slurm-recipes/configs/rebuild-deepep.sh
Original file line number Diff line number Diff line change
@@ -0,0 +1,59 @@
#!/bin/bash
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: Apache-2.0
set -eux

echo "=== Rebuilding DeepEP with kNumMaxTopK=16 for Qwen3.5 (topk=10) ==="

DEEPEP_SRC="/sgl-workspace/DeepEP"

if [ ! -d "$DEEPEP_SRC" ]; then
echo "ERROR: DeepEP source not found at $DEEPEP_SRC (mount via extra_mount)"
exit 1
fi

cd "$DEEPEP_SRC"

# Resolve NVSHMEM by locating its header, not by directory name: a bare
# `find -name nvshmem -type d` can latch onto header-less source stubs
# (e.g. tilelang's bundled tvm src/runtime/contrib/nvshmem) and the build
# then dies with "fatal error: nvshmem.h: No such file or directory".
NVSHMEM_H=$(find /usr/local /opt /usr/lib /sgl-workspace -name "nvshmem.h" -path "*/include/*" \
-not -path "*tilelang*" -not -path "*flashinfer*" -not -path "*DeepEP*" 2>/dev/null | head -1)
if [ -z "${NVSHMEM_H:-}" ]; then
echo "No NVSHMEM headers found in the image; installing nvidia-nvshmem-cu13..."
pip install --break-system-packages --quiet nvidia-nvshmem-cu13
NVSHMEM_H=$(find /usr/local/lib/python3*/dist-packages/nvidia -name "nvshmem.h" -path "*/include/*" 2>/dev/null | head -1)
fi
if [ -z "${NVSHMEM_H:-}" ]; then
echo "ERROR: NVSHMEM headers not found in image or via pip" >&2
exit 1
fi
NVSHMEM_DIR=$(dirname "$(dirname "$NVSHMEM_H")")
echo "NVSHMEM_DIR=$NVSHMEM_DIR"

# Fix missing nvshmem symlinks (container has .so.3 but not .so)
NVSHMEM_LIB="$NVSHMEM_DIR/lib"
if [ ! -f "$NVSHMEM_LIB/libnvshmem_host.so" ] && [ -f "$NVSHMEM_LIB/libnvshmem_host.so.3" ]; then
echo "Creating missing nvshmem symlinks..."
ln -sf libnvshmem_host.so.3 "$NVSHMEM_LIB/libnvshmem_host.so"
fi

# Apply kNumMaxTopK=16 patch (Qwen3.5 uses topk=10, default kNumMaxTopK=8 is insufficient)
# Note: source has both kNumMaxTopK (uppercase) and kNumMaxTopk (lowercase) as separate variables
sed -i 's/kNumMaxTopK[[:space:]]*=[[:space:]]*[0-9][0-9]*/kNumMaxTopK = 16/g' csrc/kernels/internode_ll.cu
sed -i 's/kNumMaxTopk[[:space:]]*=[[:space:]]*[0-9][0-9]*/kNumMaxTopk = 16/g' csrc/kernels/internode_ll.cu

# Verify the patch was applied
grep -q "kNumMaxTop. = 16" csrc/kernels/internode_ll.cu && echo "Patch verified: kNumMaxTopK/k=16" || {
echo "ERROR: kNumMaxTopK patch failed to apply!"; exit 1;
}

# Build with full output so we can debug failures
# set -e will auto-exit on failure
TORCH_CUDA_ARCH_LIST="10.0" \
NVSHMEM_DIR="$NVSHMEM_DIR" \
pip install -e . --no-build-isolation 2>&1

echo "=== DeepEP rebuild complete ==="
python3 -c "import deep_ep; print('deep_ep imported successfully')"
Original file line number Diff line number Diff line change
@@ -0,0 +1,155 @@
name: "qwen3.5-1p1d-dep4-dep16"

setup_script: rebuild-deepep.sh

dynamo:
hash: 46520ca59afe992fb5ef61b3197b2316f8df9b2b
install: true

frontend:
type: dynamo
enable_multiple_frontends: true
num_additional_frontends: 1
nginx_container: nginx

model:
path: "qwen3.5-fp8"
container: "lmsysorg/sglang:nightly-dev-cu13-20260709-074bb928"
precision: "fp8"

resources:
gpu_type: "gb300"
gpus_per_node: 4
prefill_nodes: 1
decode_nodes: 4
prefill_workers: 1
decode_workers: 1

backend:
type: sglang

prefill_environment:
SGLANG_JIT_DEEPGEMM_PRECOMPILE: "0"
NO_COLOR: "1"
TORCH_DISTRIBUTED_DEFAULT_TIMEOUT: "3600"
TORCH_NCCL_WATCHDOG_TIMEOUT_SEC: "3600"
TORCH_NCCL_HEARTBEAT_TIMEOUT_SEC: "3600"
PYTHONUNBUFFERED: "1"
NCCL_MNNVL_ENABLE: "1"
NCCL_CUMEM_ENABLE: "1"
MC_FORCE_MNNVL: "1"
SGLANG_DG_CACHE_DIR: "/configs/deepgemm-cache"
FLASHINFER_WORKSPACE_BASE: "/configs/flashinfer-cache"
SGLANG_DISAGGREGATION_HEARTBEAT_MAX_FAILURE: "100000"
SGLANG_DISAGGREGATION_BOOTSTRAP_TIMEOUT: "100000"
SGLANG_DISAGGREGATION_WAITING_TIMEOUT: "100000"
SGLANG_MOONCAKE_CUSTOM_MEM_POOL: "True"
SGLANG_USE_MESSAGE_QUEUE_BROADCASTER: "0"
SGLANG_DISABLE_TP_MEMORY_INBALANCE_CHECK: "1"
SGLANG_HEALTH_STARTING_OK: "1"
SGLANG_ENABLE_HEALTH_ENDPOINT_GENERATION: "0"

decode_environment:
SGLANG_JIT_DEEPGEMM_PRECOMPILE: "0"
NO_COLOR: "1"
TORCH_DISTRIBUTED_DEFAULT_TIMEOUT: "3600"
TORCH_NCCL_WATCHDOG_TIMEOUT_SEC: "3600"
TORCH_NCCL_HEARTBEAT_TIMEOUT_SEC: "3600"
PYTHONUNBUFFERED: "1"
NCCL_MNNVL_ENABLE: "1"
NCCL_CUMEM_ENABLE: "1"
MC_FORCE_MNNVL: "1"
MC_TE_METRIC: "true"
SGLANG_DG_CACHE_DIR: "/tmp/deepgemm-cache"
FLASHINFER_WORKSPACE_BASE: "/configs/flashinfer-cache"
SGLANG_DISAGGREGATION_HEARTBEAT_MAX_FAILURE: "100000"
SGLANG_DISAGGREGATION_BOOTSTRAP_TIMEOUT: "100000"
SGLANG_DISAGGREGATION_WAITING_TIMEOUT: "100000"
SGLANG_DECODE_BOOTSTRAP_TIMEOUT: "1000"
SGLANG_HACK_SEQ_BOOTSTRAP_ROOM: "1"
SGLANG_MOONCAKE_CUSTOM_MEM_POOL: "True"
SGLANG_USE_MESSAGE_QUEUE_BROADCASTER: "0"
SGLANG_DISABLE_TP_MEMORY_INBALANCE_CHECK: "1"
SGLANG_DEEPEP_NUM_MAX_DISPATCH_TOKENS_PER_RANK: "512"
SGLANG_HEALTH_CHECK_TIMEOUT: "1800"
SGLANG_HEALTH_STARTING_OK: "1"
SGLANG_ENABLE_HEALTH_ENDPOINT_GENERATION: "0"

sglang_config:
prefill:
served-model-name: "Qwen/Qwen3.5-397B-A17B-FP8"
model-path: "/model/"
quantization: "fp8"
kv-cache-dtype: "fp8_e4m3"
trust-remote-code: true

tensor-parallel-size: 4
data-parallel-size: 4
expert-parallel-size: 4
enable-dp-attention: true
enable-dp-lm-head: true

mamba-scheduler-strategy: "no_buffer"
mamba-track-interval: 2048
mamba-ssm-dtype: "bfloat16"
disaggregation-mode: "prefill"
disable-radix-cache: true
disaggregation-bootstrap-port: 31000
mem-fraction-static: 0.8
chunked-prefill-size: 65536
load-balance-method: "round_robin"
watchdog-timeout: 1000000
disable-cuda-graph: true
log-level: "info"
page-size: 64
attention-backend: "trtllm_mha"
moe-runner-backend: "flashinfer_trtllm"

decode:
served-model-name: "Qwen/Qwen3.5-397B-A17B-FP8"
model-path: "/model/"
trust-remote-code: true
quantization: "fp8"
kv-cache-dtype: "fp8_e4m3"

tensor-parallel-size: 16
data-parallel-size: 16
expert-parallel-size: 16
enable-dp-attention: true
enable-dp-lm-head: true
prefill-round-robin-balance: true

mamba-scheduler-strategy: "no_buffer"
mamba-track-interval: 128
mamba-ssm-dtype: "bfloat16"

disaggregation-mode: "decode"
disable-radix-cache: true
disaggregation-bootstrap-port: 31000

chunked-prefill-size: 4096
context-length: 4096
mem-fraction-static: 0.80
max-mamba-cache-size: 2048
max-running-requests: 2048
cuda-graph-max-bs: 128
watchdog-timeout: 1000000

page-size: 64
attention-backend: "trtllm_mha"
moe-runner-backend: "deep_gemm"
moe-a2a-backend: "deepep"
deepep-mode: "low_latency"
ep-dispatch-algorithm: "static"
eplb-algorithm: "deepseek"

decode-log-interval: 1
stream-interval: 50

benchmark:
type: "sa-bench"
isl: 1024
osl: 1024
req_rate: "inf"
random_range_ratio: 0.8
concurrencies: "512x1024x2048"
Original file line number Diff line number Diff line change
@@ -0,0 +1,123 @@
name: "qwen3.5-1p1d-tp4-tp4"

sbatch_directives:
mem: "0"

dynamo:
hash: 46520ca59afe992fb5ef61b3197b2316f8df9b2b
install: true

frontend:
type: dynamo
enable_multiple_frontends: true
num_additional_frontends: 1
nginx_container: nginx

model:
path: "qwen3.5-fp8"
container: "lmsysorg/sglang:nightly-dev-cu13-20260709-074bb928"
precision: "fp8"

resources:
gpu_type: "gb300"
gpus_per_node: 4
prefill_nodes: 1
decode_nodes: 1
prefill_workers: 1
decode_workers: 1

backend:
type: sglang

prefill_environment:
SGLANG_JIT_DEEPGEMM_PRECOMPILE: "0"
TORCH_DISTRIBUTED_DEFAULT_TIMEOUT: "3600"
TORCH_NCCL_WATCHDOG_TIMEOUT_SEC: "3600"
TORCH_NCCL_HEARTBEAT_TIMEOUT_SEC: "3600"
PYTHONUNBUFFERED: "1"
NCCL_MNNVL_ENABLE: "1"
NCCL_CUMEM_ENABLE: "1"
SGLANG_DG_CACHE_DIR: "/configs/deepgemm-cache"
FLASHINFER_WORKSPACE_BASE: "/configs/flashinfer-cache"
SGLANG_DISAGGREGATION_HEARTBEAT_MAX_FAILURE: "100000"
SGLANG_DISAGGREGATION_BOOTSTRAP_TIMEOUT: "100000"
SGLANG_DISAGGREGATION_WAITING_TIMEOUT: "100000"
SGLANG_USE_MESSAGE_QUEUE_BROADCASTER: "0"
SGLANG_MOONCAKE_CUSTOM_MEM_POOL: "True"
MC_FORCE_MNNVL: "1"

decode_environment:
SGLANG_JIT_DEEPGEMM_PRECOMPILE: "0"
TORCH_DISTRIBUTED_DEFAULT_TIMEOUT: "3600"
TORCH_NCCL_WATCHDOG_TIMEOUT_SEC: "3600"
TORCH_NCCL_HEARTBEAT_TIMEOUT_SEC: "3600"
PYTHONUNBUFFERED: "1"
NCCL_MNNVL_ENABLE: "1"
NCCL_CUMEM_ENABLE: "1"
SGLANG_DG_CACHE_DIR: "/configs/deepgemm-cache"
FLASHINFER_WORKSPACE_BASE: "/configs/flashinfer-cache"
SGLANG_DISAGGREGATION_HEARTBEAT_MAX_FAILURE: "100000"
SGLANG_DISAGGREGATION_BOOTSTRAP_TIMEOUT: "100000"
SGLANG_DISAGGREGATION_WAITING_TIMEOUT: "100000"
SGLANG_DECODE_BOOTSTRAP_TIMEOUT: "1000"
SGLANG_MOONCAKE_CUSTOM_MEM_POOL: "True"
SGLANG_USE_MESSAGE_QUEUE_BROADCASTER: "0"
SGLANG_HEALTH_CHECK_TIMEOUT: "3600"
SGLANG_HEALTH_STARTING_OK: "1"
SGLANG_ENABLE_HEALTH_ENDPOINT_GENERATION: "0"
MC_FORCE_MNNVL: "1"
MC_TE_METRIC: "true"
SGLANG_HACK_SEQ_BOOTSTRAP_ROOM: "1"

sglang_config:
prefill:
served-model-name: "Qwen/Qwen3.5-397B-A17B-FP8"
model-path: "/model/"
quantization: "fp8"
kv-cache-dtype: "fp8_e4m3"
trust-remote-code: true
attention-backend: "trtllm_mha"
tensor-parallel-size: 4
mamba-ssm-dtype: "bfloat16"
moe-runner-backend: "flashinfer_trtllm"
disable-radix-cache: true
max-running-requests: 1024
mem-fraction-static: 0.8
chunked-prefill-size: 16384
max-prefill-tokens: 16384
context-length: 4096
cuda-graph-max-bs: 1024
decode-log-interval: 1
stream-interval: 50
disaggregation-mode: "prefill"

decode:
served-model-name: "Qwen/Qwen3.5-397B-A17B-FP8"
model-path: "/model/"
quantization: "fp8"
kv-cache-dtype: "fp8_e4m3"
trust-remote-code: true
attention-backend: "trtllm_mha"
tensor-parallel-size: 4
mamba-ssm-dtype: "bfloat16"
moe-runner-backend: "flashinfer_trtllm"
disable-radix-cache: true
max-running-requests: 1024
mem-fraction-static: 0.8
chunked-prefill-size: 16384
max-prefill-tokens: 16384
context-length: 4096
cuda-graph-max-bs: 1024
decode-log-interval: 1
stream-interval: 50
disaggregation-mode: "decode"

benchmark:
type: "sa-bench"
isl: 1024
osl: 1024
req_rate: "inf"
num_prompts_mult: 10
num_warmup_mult: 1
random_range_ratio: 0.8
concurrencies: "1x2x4x8x16x32x64"
Loading