forked from Karylab-cklius/vllm
Compare commits
| Author | SHA1 | Date | |
|---|---|---|---|
|
|
c2de222aa3 | ||
|
|
88d64448f4 | ||
|
|
aef79c52cc | ||
|
|
aeb03ccd3f | ||
|
|
0d556d62ad | ||
|
|
1337c7bea5 | ||
|
|
666d1381db |
@@ -1,23 +0,0 @@
|
||||
name: vllm_rocm_ci
|
||||
job_dirs:
|
||||
- ".buildkite/hardware_tests"
|
||||
run_all_patterns:
|
||||
- "docker/Dockerfile.rocm"
|
||||
- "docker/Dockerfile.rocm_base"
|
||||
- "docker/ci-rocm.hcl"
|
||||
- "docker/docker-bake-rocm.hcl"
|
||||
- ".buildkite/hardware_tests/amd.yaml"
|
||||
- ".buildkite/scripts/ci-bake-rocm.sh"
|
||||
- ".buildkite/scripts/hardware_ci/run-amd-test.py"
|
||||
- ".buildkite/scripts/hardware_ci/run-amd-test.sh"
|
||||
- "CMakeLists.txt"
|
||||
- "requirements/common.txt"
|
||||
- "requirements/rocm.txt"
|
||||
- "requirements/build/rocm.txt"
|
||||
- "requirements/test/rocm.txt"
|
||||
- "setup.py"
|
||||
- "csrc/"
|
||||
- "cmake/"
|
||||
run_all_exclude_patterns:
|
||||
- "csrc/cpu/"
|
||||
- "cmake/cpu_extension.cmake"
|
||||
@@ -1,73 +1,22 @@
|
||||
group: Hardware - AMD Build
|
||||
group: Hardware - AMD Build
|
||||
steps:
|
||||
# Ensure ci_base is up-to-date before building the test image.
|
||||
# Compares a content hash of ci_base-affecting files against the remote
|
||||
# image label. If hashes match the build is skipped (< 30 s); if they
|
||||
# differ ci_base is rebuilt and pushed automatically.
|
||||
- label: "AMD: :docker: ensure ci_base"
|
||||
key: ensure-ci-base-amd
|
||||
- label: "AMD: :docker: build image"
|
||||
key: image-build-amd
|
||||
depends_on: []
|
||||
device: amd_cpu
|
||||
no_plugin: true
|
||||
commands:
|
||||
- bash .buildkite/scripts/ci-bake-rocm.sh ci-base-rocm-ci-with-deps
|
||||
- >
|
||||
docker build
|
||||
--build-arg max_jobs=16
|
||||
--build-arg REMOTE_VLLM=1
|
||||
--build-arg ARG_PYTORCH_ROCM_ARCH='gfx90a;gfx942;gfx950'
|
||||
--build-arg VLLM_BRANCH=$BUILDKITE_COMMIT
|
||||
--tag "rocm/vllm-ci:${BUILDKITE_COMMIT}"
|
||||
-f docker/Dockerfile.rocm
|
||||
--target test
|
||||
--no-cache
|
||||
--progress plain .
|
||||
- docker push "rocm/vllm-ci:${BUILDKITE_COMMIT}"
|
||||
env:
|
||||
DOCKER_BUILDKIT: "1"
|
||||
VLLM_BAKE_FILE: "docker/docker-bake-rocm.hcl"
|
||||
PYTORCH_ROCM_ARCH: "gfx90a;gfx942;gfx950"
|
||||
REMOTE_VLLM: "1"
|
||||
VLLM_BRANCH: "$BUILDKITE_COMMIT"
|
||||
retry:
|
||||
automatic:
|
||||
- exit_status: -1 # Agent was lost
|
||||
limit: 1
|
||||
- exit_status: -10 # Agent was lost
|
||||
limit: 1
|
||||
|
||||
- label: "AMD: :docker: build test image and artifacts"
|
||||
key: image-build-amd
|
||||
depends_on:
|
||||
- ensure-ci-base-amd
|
||||
device: amd_cpu
|
||||
no_plugin: true
|
||||
commands:
|
||||
- |
|
||||
if [[ "${ROCM_CI_ARTIFACT_ONLY:-0}" == "1" ]]; then
|
||||
echo "ROCM_CI_ARTIFACT_ONLY=1; building ROCm wheel artifact only"
|
||||
IMAGE_TAG="" bash .buildkite/scripts/ci-bake-rocm.sh test-rocm-ci-with-artifacts
|
||||
else
|
||||
bash .buildkite/scripts/ci-bake-rocm.sh test-rocm-ci-with-wheel
|
||||
fi
|
||||
- |
|
||||
docker run --rm --network=none --entrypoint /bin/bash "rocm/vllm-ci:${BUILDKITE_COMMIT}" -ec '
|
||||
if [ ! -d /vllm-workspace ]; then echo Missing directory: /vllm-workspace >&2; exit 1; fi
|
||||
if [ ! -d /vllm-workspace/tests ]; then echo Missing directory: /vllm-workspace/tests >&2; exit 1; fi
|
||||
if [ ! -d /vllm-workspace/src/vllm ]; then echo Missing directory: /vllm-workspace/src/vllm >&2; exit 1; fi
|
||||
if [ ! -x /vllm-workspace/src/vllm/vllm-rs ]; then echo Missing executable: /vllm-workspace/src/vllm/vllm-rs >&2; exit 1; fi
|
||||
command -v python3
|
||||
command -v uv
|
||||
command -v pytest
|
||||
if ! command -v amd-smi >/dev/null 2>&1 && ! command -v rocminfo >/dev/null 2>&1; then
|
||||
echo No ROCm CLI found in image >&2
|
||||
exit 1
|
||||
fi
|
||||
python3 - <<PY
|
||||
import torch, vllm
|
||||
print(torch.__version__)
|
||||
print(vllm.__version__)
|
||||
PY
|
||||
echo AMD image smoke OK
|
||||
'
|
||||
env:
|
||||
DOCKER_BUILDKIT: "1"
|
||||
VLLM_BAKE_FILE: "docker/docker-bake-rocm.hcl"
|
||||
PYTORCH_ROCM_ARCH: "gfx90a;gfx942;gfx950"
|
||||
IMAGE_TAG: "rocm/vllm-ci:$BUILDKITE_COMMIT"
|
||||
REMOTE_VLLM: "1"
|
||||
VLLM_BRANCH: "$BUILDKITE_COMMIT"
|
||||
retry:
|
||||
automatic:
|
||||
- exit_status: -1 # Agent was lost
|
||||
limit: 1
|
||||
- exit_status: -10 # Agent was lost
|
||||
limit: 1
|
||||
|
||||
@@ -16,7 +16,6 @@ steps:
|
||||
- tests/kernels/test_onednn.py
|
||||
- tests/kernels/test_awq_int4_to_int8.py
|
||||
- tests/kernels/quantization/test_cpu_fp8_scaled_mm.py
|
||||
- tests/kernels/mamba/cpu/test_cpu_gdn_ops.py
|
||||
commands:
|
||||
- |
|
||||
bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 30m "
|
||||
@@ -25,22 +24,20 @@ steps:
|
||||
pytest -x -v -s tests/kernels/moe/test_cpu_quant_fused_moe.py
|
||||
pytest -x -v -s tests/kernels/test_onednn.py
|
||||
pytest -x -v -s tests/kernels/test_awq_int4_to_int8.py
|
||||
pytest -x -v -s tests/kernels/quantization/test_cpu_fp8_scaled_mm.py
|
||||
pytest -x -v -s tests/kernels/mamba/cpu/test_cpu_gdn_ops.py"
|
||||
pytest -x -v -s tests/kernels/quantization/test_cpu_fp8_scaled_mm.py"
|
||||
|
||||
# Note: SDE can't be downloaded from CI host because of AWS WAF
|
||||
# - label: CPU-Compatibility Tests
|
||||
# depends_on: []
|
||||
# device: intel_cpu
|
||||
# no_plugin: true
|
||||
# source_file_dependencies:
|
||||
# - cmake/cpu_extension.cmake
|
||||
# - setup.py
|
||||
# - vllm/platforms/cpu.py
|
||||
# commands:
|
||||
# - |
|
||||
# bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 20m "
|
||||
# bash .buildkite/scripts/hardware_ci/run-cpu-compatibility-test.sh"
|
||||
- label: CPU-Compatibility Tests
|
||||
depends_on: []
|
||||
device: intel_cpu
|
||||
no_plugin: true
|
||||
source_file_dependencies:
|
||||
- cmake/cpu_extension.cmake
|
||||
- setup.py
|
||||
- vllm/platforms/cpu.py
|
||||
commands:
|
||||
- |
|
||||
bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 20m "
|
||||
bash .buildkite/scripts/hardware_ci/run-cpu-compatibility-test.sh"
|
||||
|
||||
- label: CPU-Language Generation and Pooling Model Tests
|
||||
depends_on: []
|
||||
@@ -57,41 +54,23 @@ steps:
|
||||
pytest -x -v -s tests/models/language/generation -m cpu_model
|
||||
pytest -x -v -s tests/models/language/pooling -m cpu_model"
|
||||
|
||||
- label: CPU-ModelRunnerV2 Tests
|
||||
depends_on: []
|
||||
device: intel_cpu
|
||||
no_plugin: true
|
||||
soft_fail: true
|
||||
source_file_dependencies:
|
||||
- vllm/v1/worker/cpu/
|
||||
- vllm/v1/worker/gpu/
|
||||
- vllm/v1/sample/ops/topk_topp_triton.py
|
||||
- vllm/v1/sample/ops/topk_topp_sampler.py
|
||||
- tests/v1/sample/test_topk_topp_sampler.py
|
||||
commands:
|
||||
- |
|
||||
bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 45m "
|
||||
uv pip install git+https://github.com/triton-lang/triton-cpu.git@270e696d
|
||||
VLLM_USE_V2_MODEL_RUNNER=1 pytest -x -v -s tests/models/language/generation/test_granite.py -m cpu_model
|
||||
# TODO: move to CPU-Kernel Tests once triton-cpu has a pre-built wheel
|
||||
pytest -x -v -s tests/v1/sample/test_topk_topp_sampler.py::TestTritonTopkTopp"
|
||||
|
||||
- label: CPU-Quantization Model Tests
|
||||
depends_on: []
|
||||
device: intel_cpu
|
||||
no_plugin: true
|
||||
source_file_dependencies:
|
||||
- csrc/cpu/
|
||||
- vllm/model_executor/layers/quantization/cpu_wna16.py
|
||||
- vllm/model_executor/layers/quantization/auto_gptq.py
|
||||
- vllm/model_executor/layers/quantization/compressed_tensors/schemes/compressed_tensors_w8a8_int8.py
|
||||
- vllm/model_executor/kernels/linear/mixed_precision/cpu.py
|
||||
- vllm/model_executor/kernels/linear/scaled_mm/cpu.py
|
||||
- vllm/model_executor/layers/quantization/kernels/scaled_mm/cpu.py
|
||||
- vllm/model_executor/layers/quantization/kernels/mixed_precision/cpu.py
|
||||
- vllm/model_executor/layers/fused_moe/experts/cpu_moe.py
|
||||
- tests/quantization/test_compressed_tensors.py
|
||||
- tests/quantization/test_cpu_wna16.py
|
||||
commands:
|
||||
- |
|
||||
bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 45m "
|
||||
bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 30m "
|
||||
pytest -x -v -s tests/quantization/test_compressed_tensors.py::test_compressed_tensors_w8a8_logprobs
|
||||
pytest -x -v -s tests/quantization/test_cpu_wna16.py"
|
||||
|
||||
|
||||
@@ -13,60 +13,6 @@ steps:
|
||||
- exit_status: -10 # Agent was lost
|
||||
limit: 2
|
||||
|
||||
- label: ":docker: :smoking: Non-root smoke tests"
|
||||
key: image-build-smoke-test
|
||||
depends_on:
|
||||
- image-build
|
||||
commands:
|
||||
# Smoke 1: the default (root) image must still be importable
|
||||
# under a non-root UID via `--user 2000:0`. Validates the `vllm` passwd
|
||||
# entry + group-0-writable /home/vllm + uv path cleanup from #31959.
|
||||
# Uses `import vllm` rather than `vllm serve --help` because the latter
|
||||
# instantiates `VllmConfig` which requires a GPU attached to the
|
||||
# container.
|
||||
- docker run --rm --user 2000:0 --entrypoint python3 "$IMAGE_TAG" -c "import vllm; print(vllm.__version__)"
|
||||
# Smoke 2: assert the non-root enabling invariants are baked
|
||||
# into the image. Runs as UID 2000:0 via a shell so we can verify
|
||||
# filesystem perms + passwd/group file state + wrapper presence without
|
||||
# triggering vLLM's GPU-requiring config-init path. The opt-in
|
||||
# `vllm-openai-nonroot` target adds only `USER vllm`, `WORKDIR
|
||||
# /home/vllm`, and an `ENTRYPOINT` override on top of these invariants;
|
||||
# its build correctness is reviewed at the Dockerfile level. Wrapper
|
||||
# logic is covered separately by the pre-commit hook
|
||||
# `test-nonroot-entrypoint` (see .pre-commit-config.yaml).
|
||||
- |
|
||||
docker run --rm --user 2000:0 --entrypoint /bin/sh "$IMAGE_TAG" -ec '
|
||||
if ! getent passwd 2000 | grep -q ^vllm:; then
|
||||
echo FAIL: UID 2000 != vllm
|
||||
exit 1
|
||||
fi
|
||||
if ! id -gn 2>/dev/null | grep -qx root; then
|
||||
echo FAIL: GID 0 not root group
|
||||
exit 1
|
||||
fi
|
||||
touch /home/vllm/.smoke && rm /home/vllm/.smoke
|
||||
touch /opt/uv/cache/.smoke && rm /opt/uv/cache/.smoke
|
||||
if ! test -x /usr/local/bin/vllm-nonroot-entrypoint.sh; then
|
||||
echo FAIL: wrapper missing
|
||||
exit 1
|
||||
fi
|
||||
if ! test -w /etc/passwd; then
|
||||
echo FAIL: /etc/passwd not group-writable
|
||||
exit 1
|
||||
fi
|
||||
if ! test -w /etc/group; then
|
||||
echo FAIL: /etc/group not group-writable
|
||||
exit 1
|
||||
fi
|
||||
echo non-root invariants OK
|
||||
'
|
||||
retry:
|
||||
automatic:
|
||||
- exit_status: -1 # Agent was lost
|
||||
limit: 2
|
||||
- exit_status: -10 # Agent was lost
|
||||
limit: 2
|
||||
|
||||
- label: ":docker: Build CPU image"
|
||||
key: image-build-cpu
|
||||
depends_on: []
|
||||
@@ -110,21 +56,3 @@ steps:
|
||||
limit: 2
|
||||
- exit_status: -10 # Agent was lost
|
||||
limit: 2
|
||||
|
||||
- label: ":docker: Build arm64 image"
|
||||
key: arm64-image-build
|
||||
depends_on: []
|
||||
source_file_dependencies:
|
||||
- ".buildkite/image_build/image_build.yaml"
|
||||
- ".buildkite/image_build/image_build_arm64.sh"
|
||||
- "docker/Dockerfile"
|
||||
commands:
|
||||
- .buildkite/image_build/image_build_arm64.sh $REGISTRY $REPO $BUILDKITE_COMMIT
|
||||
env:
|
||||
DOCKER_BUILDKIT: "1"
|
||||
retry:
|
||||
automatic:
|
||||
- exit_status: -1 # Agent was lost
|
||||
limit: 2
|
||||
- exit_status: -10 # Agent was lost
|
||||
limit: 2
|
||||
|
||||
@@ -1,37 +0,0 @@
|
||||
#!/bin/bash
|
||||
set -e
|
||||
|
||||
if [[ $# -lt 3 ]]; then
|
||||
echo "Usage: $0 <registry> <repo> <commit>"
|
||||
exit 1
|
||||
fi
|
||||
|
||||
REGISTRY=$1
|
||||
REPO=$2
|
||||
BUILDKITE_COMMIT=$3
|
||||
|
||||
# authenticate with AWS ECR
|
||||
aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY" || true
|
||||
|
||||
# skip build if image already exists
|
||||
if [[ -z $(docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-arm64) ]]; then
|
||||
echo "Image not found, proceeding with build..."
|
||||
else
|
||||
echo "Image found"
|
||||
exit 0
|
||||
fi
|
||||
|
||||
# build (Grace/GH200 is the arm64 GPU target; sm_90)
|
||||
docker build --file docker/Dockerfile \
|
||||
--platform linux/arm64 \
|
||||
--build-arg max_jobs=16 \
|
||||
--build-arg nvcc_threads=4 \
|
||||
--build-arg torch_cuda_arch_list="9.0" \
|
||||
--build-arg USE_SCCACHE=1 \
|
||||
--build-arg buildkite_commit="$BUILDKITE_COMMIT" \
|
||||
--tag "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-arm64 \
|
||||
--target test \
|
||||
--progress plain .
|
||||
|
||||
# push
|
||||
docker push "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-arm64
|
||||
@@ -11,7 +11,7 @@ REPO=$2
|
||||
BUILDKITE_COMMIT=$3
|
||||
|
||||
# authenticate with AWS ECR
|
||||
aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY" || true
|
||||
aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY"
|
||||
|
||||
# skip build if image already exists
|
||||
if [[ -z $(docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-hpu) ]]; then
|
||||
|
||||
@@ -11,8 +11,8 @@ REPO=$2
|
||||
BUILDKITE_COMMIT=$3
|
||||
|
||||
# authenticate with AWS ECR
|
||||
aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY" || true
|
||||
aws ecr get-login-password --region us-east-1 | docker login --username AWS --password-stdin 936637512419.dkr.ecr.us-east-1.amazonaws.com || true
|
||||
aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY"
|
||||
aws ecr get-login-password --region us-east-1 | docker login --username AWS --password-stdin 936637512419.dkr.ecr.us-east-1.amazonaws.com
|
||||
|
||||
# skip build if image already exists
|
||||
if ! docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-xpu &> /dev/null; then
|
||||
|
||||
@@ -1,22 +0,0 @@
|
||||
group: Basic Correctness
|
||||
depends_on:
|
||||
- image-build-xpu
|
||||
steps:
|
||||
- label: XPU Sleep Mode
|
||||
timeout_in_minutes: 30
|
||||
device: intel_gpu
|
||||
no_plugin: true
|
||||
working_dir: "."
|
||||
env:
|
||||
REGISTRY: "public.ecr.aws/q9t5s3a7"
|
||||
REPO: "vllm-ci-test-repo"
|
||||
VLLM_TEST_DEVICE: "xpu"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/basic_correctness/test_cumem.py
|
||||
commands:
|
||||
- >-
|
||||
bash .buildkite/scripts/hardware_ci/run-intel-test.sh
|
||||
'cd tests &&
|
||||
export VLLM_WORKER_MULTIPROC_METHOD=spawn &&
|
||||
pytest -v -s basic_correctness/test_mem.py::test_end_to_end'
|
||||
@@ -1,23 +0,0 @@
|
||||
group: Expert Parallelism
|
||||
depends_on:
|
||||
- image-build-xpu
|
||||
steps:
|
||||
- label: EPLB Algorithm
|
||||
key: eplb-algorithm
|
||||
timeout_in_minutes: 45
|
||||
device: intel_gpu
|
||||
no_plugin: true
|
||||
working_dir: "."
|
||||
env:
|
||||
REGISTRY: "public.ecr.aws/q9t5s3a7"
|
||||
REPO: "vllm-ci-test-repo"
|
||||
VLLM_TEST_DEVICE: "xpu"
|
||||
source_file_dependencies:
|
||||
- vllm/distributed/eplb
|
||||
- tests/distributed/test_eplb_algo.py
|
||||
- tests/distributed/test_eplb_utils.py
|
||||
commands:
|
||||
- >-
|
||||
bash .buildkite/scripts/hardware_ci/run-intel-test.sh
|
||||
'cd tests &&
|
||||
pytest -v -s distributed/test_eplb_algo.py'
|
||||
@@ -38,17 +38,7 @@ steps:
|
||||
REPO: "vllm-ci-test-repo"
|
||||
VLLM_TEST_DEVICE: "xpu"
|
||||
source_file_dependencies:
|
||||
- vllm/config/
|
||||
- vllm/distributed/
|
||||
- vllm/engine/
|
||||
- vllm/inputs/
|
||||
- vllm/logger.py
|
||||
- vllm/model_executor/
|
||||
- vllm/platforms/
|
||||
- vllm/sampling_params.py
|
||||
- vllm/transformers_utils/
|
||||
- vllm/utils/
|
||||
- vllm/v1/
|
||||
- vllm/
|
||||
- tests/v1/sample
|
||||
- tests/v1/logits_processors
|
||||
- tests/v1/test_oracle.py
|
||||
@@ -62,126 +52,4 @@ steps:
|
||||
pytest -v -s v1/logits_processors --ignore=v1/logits_processors/test_custom_online.py --ignore=v1/logits_processors/test_custom_offline.py &&
|
||||
pytest -v -s v1/test_oracle.py &&
|
||||
pytest -v -s v1/test_request.py &&
|
||||
pytest -v -s v1/test_outputs.py &&
|
||||
pytest -v -s v1/sample/test_topk_topp_sampler.py'
|
||||
|
||||
- label: XPU CPU Offload
|
||||
timeout_in_minutes: 60
|
||||
device: intel_gpu
|
||||
no_plugin: true
|
||||
working_dir: "."
|
||||
env:
|
||||
REGISTRY: "public.ecr.aws/q9t5s3a7"
|
||||
REPO: "vllm-ci-test-repo"
|
||||
VLLM_TEST_DEVICE: "xpu"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- vllm/v1/kv_offload/
|
||||
- vllm/v1/kv_connector/
|
||||
- tests/v1/kv_offload/
|
||||
- tests/v1/kv_connector/unit/test_offloading_connector.py
|
||||
commands:
|
||||
- >-
|
||||
bash .buildkite/scripts/hardware_ci/run-intel-test.sh
|
||||
'export VLLM_WORKER_MULTIPROC_METHOD=spawn &&
|
||||
cd tests &&
|
||||
pytest -v -s v1/kv_offload &&
|
||||
pytest -v -s v1/kv_connector/unit/test_offloading_connector.py'
|
||||
|
||||
- label: Regression
|
||||
key: regression
|
||||
timeout_in_minutes: 30
|
||||
device: intel_gpu
|
||||
no_plugin: true
|
||||
working_dir: "."
|
||||
env:
|
||||
REGISTRY: "public.ecr.aws/q9t5s3a7"
|
||||
REPO: "vllm-ci-test-repo"
|
||||
VLLM_TEST_DEVICE: "xpu"
|
||||
source_file_dependencies:
|
||||
- vllm/config/
|
||||
- vllm/distributed/
|
||||
- vllm/engine/
|
||||
- vllm/inputs/
|
||||
- vllm/model_executor/
|
||||
- vllm/multimodal/
|
||||
- vllm/platforms/
|
||||
- vllm/sampling_params.py
|
||||
- vllm/transformers_utils/
|
||||
- vllm/utils/
|
||||
- vllm/v1/
|
||||
- tests/test_regression
|
||||
commands:
|
||||
- >-
|
||||
bash .buildkite/scripts/hardware_ci/run-intel-test.sh
|
||||
'pip install modelscope &&
|
||||
cd tests &&
|
||||
pytest -v -s test_regression.py'
|
||||
|
||||
- label: Metrics, Tracing (2 GPUs)
|
||||
key: metrics-tracing-2-gpus
|
||||
timeout_in_minutes: 30
|
||||
num_devices: 2
|
||||
device: intel_gpu
|
||||
no_plugin: true
|
||||
working_dir: "."
|
||||
env:
|
||||
REGISTRY: "public.ecr.aws/q9t5s3a7"
|
||||
REPO: "vllm-ci-test-repo"
|
||||
VLLM_TEST_DEVICE: "xpu"
|
||||
source_file_dependencies:
|
||||
- vllm/config/
|
||||
- vllm/distributed/
|
||||
- vllm/engine/
|
||||
- vllm/inputs/
|
||||
- vllm/model_executor/
|
||||
- vllm/multimodal/
|
||||
- vllm/platforms/
|
||||
- vllm/sampling_params.py
|
||||
- vllm/tracing/
|
||||
- vllm/transformers_utils/
|
||||
- vllm/utils/
|
||||
- vllm/v1/
|
||||
- tests/v1/tracing
|
||||
commands:
|
||||
- >-
|
||||
bash .buildkite/scripts/hardware_ci/run-intel-test.sh
|
||||
'pip install opentelemetry-sdk\>=1.26.0 opentelemetry-api\>=1.26.0 opentelemetry-exporter-otlp\>=1.26.0 opentelemetry-semantic-conventions-ai\>=0.4.1 &&
|
||||
cd tests &&
|
||||
pytest -v -s v1/tracing'
|
||||
|
||||
- label: Async Engine, Inputs, Utils, Worker
|
||||
key: async-engine-inputs-utils-worker
|
||||
timeout_in_minutes: 30
|
||||
device: intel_gpu
|
||||
no_plugin: true
|
||||
working_dir: "."
|
||||
env:
|
||||
REGISTRY: "public.ecr.aws/q9t5s3a7"
|
||||
REPO: "vllm-ci-test-repo"
|
||||
VLLM_TEST_DEVICE: "xpu"
|
||||
source_file_dependencies:
|
||||
- vllm/assets/
|
||||
- vllm/config/
|
||||
- vllm/distributed/
|
||||
- vllm/engine/
|
||||
- vllm/inputs/
|
||||
- vllm/model_executor/
|
||||
- vllm/multimodal/
|
||||
- vllm/platforms/
|
||||
- vllm/sampling_params.py
|
||||
- vllm/tokenizers/
|
||||
- vllm/transformers_utils/
|
||||
- vllm/utils/
|
||||
- vllm/v1/
|
||||
- tests/detokenizer
|
||||
- tests/multimodal
|
||||
- tests/utils_
|
||||
commands:
|
||||
- >-
|
||||
bash .buildkite/scripts/hardware_ci/run-intel-test.sh
|
||||
'cd tests &&
|
||||
pip install av &&
|
||||
pytest -v -s detokenizer &&
|
||||
pytest -v -s -m "not cpu_test" ./multimodal &&
|
||||
pytest -v -s utils_ --ignore=utils_/test_mem_utils.py'
|
||||
pytest -v -s v1/test_outputs.py'
|
||||
|
||||
@@ -1,111 +0,0 @@
|
||||
group: Models - Multimodal
|
||||
depends_on:
|
||||
- image-build-xpu
|
||||
steps:
|
||||
- label: "Multi-Modal Models (Standard) 1: qwen2"
|
||||
key: multi-modal-models-standard-1-qwen2
|
||||
timeout_in_minutes: 45
|
||||
device: intel_gpu
|
||||
no_plugin: true
|
||||
working_dir: "."
|
||||
env:
|
||||
REGISTRY: "public.ecr.aws/q9t5s3a7"
|
||||
REPO: "vllm-ci-test-repo"
|
||||
VLLM_TEST_DEVICE: "xpu"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/models/multimodal
|
||||
commands:
|
||||
- >-
|
||||
bash .buildkite/scripts/hardware_ci/run-intel-test.sh
|
||||
'pip install av git+https://github.com/TIGER-AI-Lab/Mantis.git &&
|
||||
cd tests &&
|
||||
pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "qwen2" &&
|
||||
pytest -v -s models/multimodal/generation/test_ultravox.py -m core_model'
|
||||
|
||||
- label: "Multi-Modal Models (Standard) 2: qwen3 + gemma"
|
||||
key: multi-modal-models-standard-2-qwen3-gemma
|
||||
timeout_in_minutes: 45
|
||||
device: intel_gpu
|
||||
no_plugin: true
|
||||
working_dir: "."
|
||||
env:
|
||||
REGISTRY: "public.ecr.aws/q9t5s3a7"
|
||||
REPO: "vllm-ci-test-repo"
|
||||
VLLM_TEST_DEVICE: "xpu"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/models/multimodal
|
||||
commands:
|
||||
- >-
|
||||
bash .buildkite/scripts/hardware_ci/run-intel-test.sh
|
||||
'pip install git+https://github.com/TIGER-AI-Lab/Mantis.git &&
|
||||
cd tests &&
|
||||
pytest -v -s models/multimodal/generation/test_qwen2_5_vl.py -m core_model'
|
||||
|
||||
- label: "Multi-Modal Models (Standard) 3: llava + qwen2_vl"
|
||||
key: multi-modal-models-standard-3-llava-qwen2-vl
|
||||
timeout_in_minutes: 45
|
||||
device: intel_gpu
|
||||
no_plugin: true
|
||||
working_dir: "."
|
||||
env:
|
||||
REGISTRY: "public.ecr.aws/q9t5s3a7"
|
||||
REPO: "vllm-ci-test-repo"
|
||||
VLLM_TEST_DEVICE: "xpu"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/models/multimodal
|
||||
commands:
|
||||
- >-
|
||||
bash .buildkite/scripts/hardware_ci/run-intel-test.sh
|
||||
'pip install git+https://github.com/TIGER-AI-Lab/Mantis.git &&
|
||||
cd tests &&
|
||||
pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "not qwen2 and not qwen3 and not gemma" &&
|
||||
pytest -v -s models/multimodal/generation/test_qwen2_vl.py -m core_model'
|
||||
|
||||
- label: "Multi-Modal Models (Standard) 4: other + whisper"
|
||||
key: multi-modal-models-standard-4-other-whisper
|
||||
timeout_in_minutes: 45
|
||||
device: intel_gpu
|
||||
no_plugin: true
|
||||
working_dir: "."
|
||||
env:
|
||||
REGISTRY: "public.ecr.aws/q9t5s3a7"
|
||||
REPO: "vllm-ci-test-repo"
|
||||
VLLM_TEST_DEVICE: "xpu"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/models/multimodal
|
||||
commands:
|
||||
- >-
|
||||
bash .buildkite/scripts/hardware_ci/run-intel-test.sh
|
||||
'pip install av git+https://github.com/TIGER-AI-Lab/Mantis.git &&
|
||||
cd tests &&
|
||||
pytest -v -s models/multimodal -m core_model --ignore models/multimodal/generation/test_common.py --ignore models/multimodal/generation/test_ultravox.py --ignore models/multimodal/generation/test_qwen2_5_vl.py --ignore models/multimodal/generation/test_qwen2_vl.py --ignore models/multimodal/generation/test_whisper.py --ignore models/multimodal/generation/test_memory_leak.py --ignore models/multimodal/processing'
|
||||
|
||||
- label: Multi-Modal Processor # 44min
|
||||
key: multi-modal-processor
|
||||
timeout_in_minutes: 45
|
||||
device: intel_gpu
|
||||
no_plugin: true
|
||||
working_dir: "."
|
||||
env:
|
||||
REGISTRY: "public.ecr.aws/q9t5s3a7"
|
||||
REPO: "vllm-ci-test-repo"
|
||||
VLLM_TEST_DEVICE: "xpu"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/models/multimodal
|
||||
- tests/models/registry.py
|
||||
commands:
|
||||
- >-
|
||||
bash .buildkite/scripts/hardware_ci/run-intel-test.sh
|
||||
'pip install av matplotlib ftfy git+https://github.com/TIGER-AI-Lab/Mantis.git &&
|
||||
pip install open-clip-torch --no-deps &&
|
||||
cd tests &&
|
||||
pytest -v -s models/multimodal/processing/test_tensor_schema.py
|
||||
--deselect "tests/models/multimodal/processing/test_tensor_schema.py::test_model_tensor_schema[mistralai/Mistral-Large-3-675B-Instruct-2512-NVFP4]"
|
||||
--deselect "tests/models/multimodal/processing/test_tensor_schema.py::test_model_tensor_schema[Qwen/Qwen2.5-Omni-7B-AWQ]"
|
||||
--num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB'
|
||||
parallelism: 4
|
||||
@@ -39,11 +39,7 @@ steps:
|
||||
python3 examples/basic/offline_inference/generate.py --model nvidia/Llama-3.1-8B-Instruct-FP8 --block-size 64 --enforce-eager --quantization modelopt --kv-cache-dtype fp8 --attention-backend TRITON_ATTN --max-model-len 4096 &&
|
||||
python3 examples/basic/offline_inference/generate.py --model superjob/Qwen3-4B-Instruct-2507-GPTQ-Int4 --block-size 64 --enforce-eager --max-model-len 8192 &&
|
||||
python3 examples/basic/offline_inference/generate.py --model ibm-research/PowerMoE-3b --block-size 64 --enforce-eager -tp 2 &&
|
||||
python3 examples/basic/offline_inference/generate.py --model ibm-research/PowerMoE-3b --block-size 64 --enforce-eager -tp 2 --enable-expert-parallel &&
|
||||
python3 examples/basic/offline_inference/generate.py --model superjob/Qwen3-4B-Instruct-2507-GPTQ-Int4 --max-model-len 8192 &&
|
||||
VLLM_XPU_FUSED_MOE_USE_REF=1 python3 examples/basic/offline_inference/generate.py --model Qwen/Qwen3-30B-A3B-Instruct-2507-FP8 --enforce-eager -tp 2 --max-model-len 8192 &&
|
||||
python3 examples/basic/offline_inference/generate.py --model INCModel/Qwen3-30B-A3B-Instruct-2507-MXFP4-LLMC --enforce-eager -tp 2 --max-model-len 8192
|
||||
'
|
||||
python3 examples/basic/offline_inference/generate.py --model ibm-research/PowerMoE-3b --block-size 64 --enforce-eager -tp 2 --enable-expert-parallel'
|
||||
- label: "XPU V1 test"
|
||||
depends_on:
|
||||
- image-build-xpu
|
||||
@@ -68,22 +64,3 @@ steps:
|
||||
pytest -v -s v1/test_serial_utils.py &&
|
||||
pytest -v -s v1/spec_decode --ignore=v1/spec_decode/test_max_len.py --ignore=v1/spec_decode/test_speculators_eagle3.py --ignore=v1/spec_decode/test_acceptance_length.py &&
|
||||
pytest -v -s v1/kv_connector/unit --ignore=v1/kv_connector/unit/test_multi_connector.py --ignore=v1/kv_connector/unit/test_example_connector.py --ignore=v1/kv_connector/unit/test_lmcache_integration.py --ignore=v1/kv_connector/unit/test_hf3fs_client.py --ignore=v1/kv_connector/unit/test_hf3fs_connector.py --ignore=v1/kv_connector/unit/test_hf3fs_metadata_server.py --ignore=v1/kv_connector/unit/test_offloading_connector.py'
|
||||
- label: "XPU server test"
|
||||
depends_on:
|
||||
- image-build-xpu
|
||||
timeout_in_minutes: 30
|
||||
device: intel_gpu
|
||||
no_plugin: true
|
||||
env:
|
||||
REGISTRY: "public.ecr.aws/q9t5s3a7"
|
||||
REPO: "vllm-ci-test-repo"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- .buildkite/intel_jobs/test-intel.yaml
|
||||
commands:
|
||||
- >-
|
||||
bash .buildkite/scripts/hardware_ci/run-intel-test.sh
|
||||
'pip install av &&
|
||||
cd tests &&
|
||||
pytest -v -s entrypoints/multimodal/openai/chat_completion/test_audio_in_video.py &&
|
||||
pytest -v -s benchmarks/test_serve_cli.py'
|
||||
|
||||
@@ -6,7 +6,9 @@ tasks:
|
||||
value: 0.7142
|
||||
- name: "exact_match,flexible-extract"
|
||||
value: 0.4579
|
||||
moe_backend: "flashinfer_cutlass"
|
||||
env_vars:
|
||||
VLLM_USE_FLASHINFER_MOE_FP8: "1"
|
||||
VLLM_FLASHINFER_MOE_BACKEND: "throughput"
|
||||
limit: 1319
|
||||
num_fewshot: 5
|
||||
max_model_len: 262144
|
||||
|
||||
@@ -68,10 +68,6 @@ def launch_lm_eval(eval_config, tp_size):
|
||||
if current_platform.is_rocm() and "Nemotron-3" in eval_config["model_name"]:
|
||||
model_args += "attention_backend=TRITON_ATTN"
|
||||
|
||||
moe_backend = eval_config.get("moe_backend", None)
|
||||
if moe_backend is not None:
|
||||
model_args += f"moe_backend={moe_backend},"
|
||||
|
||||
env_vars = eval_config.get("env_vars", None)
|
||||
with scoped_env_vars(env_vars):
|
||||
results = lm_eval.simple_evaluate(
|
||||
|
||||
@@ -2,6 +2,7 @@
|
||||
{
|
||||
"test_name": "latency_llama8B_tp1",
|
||||
"environment_variables": {
|
||||
"VLLM_RPC_TIMEOUT": 100000,
|
||||
"VLLM_ALLOW_LONG_MAX_MODEL_LEN": 1,
|
||||
"VLLM_ENGINE_ITERATION_TIMEOUT_S": 120,
|
||||
"VLLM_CPU_KVCACHE_SPACE": 40
|
||||
|
||||
@@ -2,6 +2,7 @@
|
||||
{
|
||||
"test_name": "latency_llama8B_tp2",
|
||||
"environment_variables": {
|
||||
"VLLM_RPC_TIMEOUT": 100000,
|
||||
"VLLM_ALLOW_LONG_MAX_MODEL_LEN": 1,
|
||||
"VLLM_ENGINE_ITERATION_TIMEOUT_S": 120,
|
||||
"VLLM_CPU_SGL_KERNEL": 1,
|
||||
|
||||
@@ -13,6 +13,7 @@
|
||||
200
|
||||
],
|
||||
"server_environment_variables": {
|
||||
"VLLM_RPC_TIMEOUT": 100000,
|
||||
"VLLM_ALLOW_LONG_MAX_MODEL_LEN": 1,
|
||||
"VLLM_ENGINE_ITERATION_TIMEOUT_S": 120,
|
||||
"VLLM_CPU_SGL_KERNEL": 1,
|
||||
|
||||
@@ -5,6 +5,7 @@
|
||||
],
|
||||
"max_concurrency_list": [12, 16, 24, 32, 64, 128, 200],
|
||||
"server_environment_variables": {
|
||||
"VLLM_RPC_TIMEOUT": 100000,
|
||||
"VLLM_ENGINE_ITERATION_TIMEOUT_S": 120
|
||||
},
|
||||
"server_parameters": {
|
||||
|
||||
@@ -9,6 +9,7 @@
|
||||
128
|
||||
],
|
||||
"server_environment_variables": {
|
||||
"VLLM_RPC_TIMEOUT": 100000,
|
||||
"VLLM_ALLOW_LONG_MAX_MODEL_LEN": 1,
|
||||
"VLLM_ENGINE_ITERATION_TIMEOUT_S": 120,
|
||||
"VLLM_CPU_SGL_KERNEL": 1,
|
||||
|
||||
@@ -5,6 +5,7 @@
|
||||
],
|
||||
"max_concurrency_list": [12, 16, 24, 32, 64, 128, 200],
|
||||
"server_environment_variables": {
|
||||
"VLLM_RPC_TIMEOUT": 100000,
|
||||
"VLLM_ALLOW_LONG_MAX_MODEL_LEN": 1,
|
||||
"VLLM_ENGINE_ITERATION_TIMEOUT_S": 120,
|
||||
"VLLM_CPU_SGL_KERNEL": 1,
|
||||
|
||||
@@ -5,6 +5,7 @@
|
||||
],
|
||||
"max_concurrency_list": [12, 16, 24, 32, 64, 128, 200],
|
||||
"server_environment_variables": {
|
||||
"VLLM_RPC_TIMEOUT": 100000,
|
||||
"VLLM_ALLOW_LONG_MAX_MODEL_LEN": 1,
|
||||
"VLLM_ENGINE_ITERATION_TIMEOUT_S": 120,
|
||||
"VLLM_CPU_SGL_KERNEL": 1,
|
||||
|
||||
@@ -1,112 +1,77 @@
|
||||
{
|
||||
"defaults": {
|
||||
"qps_list": [
|
||||
"inf"
|
||||
],
|
||||
"max_concurrency_list": [12, 16, 24, 32, 64, 128, 200],
|
||||
"server_parameters": {
|
||||
"model": "meta-llama/Llama-3.1-8B-Instruct",
|
||||
"tensor_parallel_size": 1,
|
||||
"dtype": "bfloat16"
|
||||
[
|
||||
{
|
||||
"test_name": "serving_llama8B_tp1_sharegpt",
|
||||
"qps_list": [1, 4, 16, "inf"],
|
||||
"server_parameters": {
|
||||
"model": "meta-llama/Meta-Llama-3.1-8B-Instruct",
|
||||
"tensor_parallel_size": 1,
|
||||
"disable_log_stats": "",
|
||||
"load_format": "dummy"
|
||||
},
|
||||
"client_parameters": {
|
||||
"model": "meta-llama/Meta-Llama-3.1-8B-Instruct",
|
||||
"backend": "vllm",
|
||||
"dataset_name": "sharegpt",
|
||||
"dataset_path": "./ShareGPT_V3_unfiltered_cleaned_split.json",
|
||||
"temperature": 0,
|
||||
"num_prompts": 200
|
||||
}
|
||||
},
|
||||
"client_parameters": {
|
||||
"model": "meta-llama/Llama-3.1-8B-Instruct",
|
||||
"backend": "vllm",
|
||||
"ignore-eos": "",
|
||||
"temperature": 0,
|
||||
"num_prompts": 200
|
||||
{
|
||||
"test_name": "serving_llama70B_tp4_sharegpt",
|
||||
"qps_list": [1, 4, 16, "inf"],
|
||||
"server_parameters": {
|
||||
"model": "meta-llama/Meta-Llama-3.1-70B-Instruct",
|
||||
"tensor_parallel_size": 4,
|
||||
"disable_log_stats": "",
|
||||
"load_format": "dummy"
|
||||
},
|
||||
"client_parameters": {
|
||||
"model": "meta-llama/Meta-Llama-3.1-70B-Instruct",
|
||||
"backend": "vllm",
|
||||
"dataset_name": "sharegpt",
|
||||
"dataset_path": "./ShareGPT_V3_unfiltered_cleaned_split.json",
|
||||
"temperature": 0,
|
||||
"num_prompts": 200
|
||||
}
|
||||
},
|
||||
{
|
||||
"test_name": "serving_mixtral8x7B_tp2_sharegpt",
|
||||
"qps_list": [1, 4, 16, "inf"],
|
||||
"server_parameters": {
|
||||
"model": "mistralai/Mixtral-8x7B-Instruct-v0.1",
|
||||
"tensor_parallel_size": 2,
|
||||
"disable_log_stats": "",
|
||||
"load_format": "dummy"
|
||||
},
|
||||
"client_parameters": {
|
||||
"model": "mistralai/Mixtral-8x7B-Instruct-v0.1",
|
||||
"backend": "vllm",
|
||||
"dataset_name": "sharegpt",
|
||||
"dataset_path": "./ShareGPT_V3_unfiltered_cleaned_split.json",
|
||||
"temperature": 0,
|
||||
"num_prompts": 200
|
||||
}
|
||||
},
|
||||
{
|
||||
"test_name": "serving_llama70B_tp4_sharegpt_specdecode",
|
||||
"qps_list": [2],
|
||||
"server_parameters": {
|
||||
"model": "meta-llama/Meta-Llama-3.1-70B-Instruct",
|
||||
"tensor_parallel_size": 4,
|
||||
"speculative_config": {
|
||||
"model": "turboderp/Qwama-0.5B-Instruct",
|
||||
"num_speculative_tokens": 4,
|
||||
"draft_tensor_parallel_size": 1
|
||||
}
|
||||
},
|
||||
"client_parameters": {
|
||||
"model": "meta-llama/Meta-Llama-3.1-70B-Instruct",
|
||||
"backend": "vllm",
|
||||
"dataset_name": "sharegpt",
|
||||
"dataset_path": "./ShareGPT_V3_unfiltered_cleaned_split.json",
|
||||
"temperature": 0,
|
||||
"num_prompts": 200
|
||||
}
|
||||
}
|
||||
},
|
||||
"tests": [
|
||||
{
|
||||
"test_name": "serving_llama8B_tp1_sharegpt",
|
||||
"server_parameters": {
|
||||
"tensor_parallel_size": 1
|
||||
},
|
||||
"client_parameters": {
|
||||
"dataset_name": "sharegpt",
|
||||
"dataset_path": "./ShareGPT_V3_unfiltered_cleaned_split.json"
|
||||
}
|
||||
},
|
||||
{
|
||||
"dataset_name": "sharegpt",
|
||||
"dataset_path": "./ShareGPT_V3_unfiltered_cleaned_split.json"
|
||||
}
|
||||
},
|
||||
{
|
||||
"test_name": "serving_llama8B_tp1_random_128_128",
|
||||
"server_parameters": {
|
||||
"tensor_parallel_size": 1
|
||||
},
|
||||
"client_parameters": {
|
||||
"dataset_name": "random",
|
||||
"random-input-len": 128,
|
||||
"random-output-len": 128
|
||||
}
|
||||
},
|
||||
{
|
||||
"test_name": "serving_llama8B_tp1_random_128_2048",
|
||||
"server_parameters": {
|
||||
"tensor_parallel_size": 1
|
||||
},
|
||||
"client_parameters": {
|
||||
"dataset_name": "random",
|
||||
"random-input-len": 128,
|
||||
"random-output-len": 2048
|
||||
}
|
||||
},
|
||||
{
|
||||
"test_name": "serving_llama8B_tp1_random_2048_128",
|
||||
"server_parameters": {
|
||||
"tensor_parallel_size": 1
|
||||
},
|
||||
"client_parameters": {
|
||||
"dataset_name": "random",
|
||||
"random-input-len": 2048,
|
||||
"random-output-len": 128
|
||||
}
|
||||
},
|
||||
{
|
||||
"test_name": "serving_llama8B_tp1_random_2048_2048",
|
||||
"server_parameters": {
|
||||
"tensor_parallel_size": 1
|
||||
},
|
||||
"client_parameters": {
|
||||
"dataset_name": "random",
|
||||
"random-input-len": 2048,
|
||||
"random-output-len": 2048
|
||||
}
|
||||
},
|
||||
{
|
||||
"test_name": "serving_llama70B_tp4_random_128_128",
|
||||
"server_parameters": {
|
||||
"model": "meta-llama/Llama-3.3-70B-Instruct",
|
||||
"async_scheduling": "",
|
||||
"no_enable_prefix_caching": "",
|
||||
"max_num_batched_tokens": 8192
|
||||
},
|
||||
"client_parameters": {
|
||||
"model": "meta-llama/Llama-3.3-70B-Instruct",
|
||||
"dataset_name": "random",
|
||||
"random-input-len": 128,
|
||||
"random-output-len": 128
|
||||
}
|
||||
},
|
||||
{
|
||||
"test_name": "serving_gemma4-e4b_tp1_random_128_128",
|
||||
"server_parameters": {
|
||||
"model": "google/gemma-4-E4B-it",
|
||||
"enable_auto_tool_choice": "",
|
||||
"tool_call_parser": "gemma4",
|
||||
"chat_template": "examples/tool_chat_template_gemma4.jinja",
|
||||
"reasoning_parser": "gemma4"
|
||||
},
|
||||
"client_parameters": {
|
||||
"model": "google/gemma-4-E4B-it",
|
||||
"dataset_name": "random",
|
||||
"random-input-len": 128,
|
||||
"random-output-len": 128
|
||||
}
|
||||
}
|
||||
]
|
||||
}
|
||||
]
|
||||
|
||||
@@ -2,6 +2,7 @@
|
||||
{
|
||||
"test_name": "throughput_llama8B_tp1",
|
||||
"environment_variables": {
|
||||
"VLLM_RPC_TIMEOUT": 100000,
|
||||
"VLLM_ALLOW_LONG_MAX_MODEL_LEN": 1,
|
||||
"VLLM_ENGINE_ITERATION_TIMEOUT_S": 120,
|
||||
"VLLM_CPU_KVCACHE_SPACE": 40
|
||||
|
||||
@@ -2,6 +2,7 @@
|
||||
{
|
||||
"test_name": "throughput_llama8B_tp2",
|
||||
"environment_variables": {
|
||||
"VLLM_RPC_TIMEOUT": 100000,
|
||||
"VLLM_ALLOW_LONG_MAX_MODEL_LEN": 1,
|
||||
"VLLM_ENGINE_ITERATION_TIMEOUT_S": 120,
|
||||
"VLLM_CPU_SGL_KERNEL": 1,
|
||||
|
||||
@@ -1,25 +1,12 @@
|
||||
# CUDA architecture lists — following PyTorch RELEASE.md
|
||||
# (https://github.com/pytorch/pytorch/blob/main/RELEASE.md)
|
||||
# SM86 included for broader Ampere coverage; SM89 for marlin fp8 support
|
||||
# These requested arches are filtered by CMake's CUDA_SUPPORTED_ARCHS before
|
||||
# per-kernel arch selection. Do not add +PTX here: top-level +PTX is stripped
|
||||
# during that filtering, so kernels that need PTX must request it locally.
|
||||
env:
|
||||
# for CUDA >=13, sm_100+ targets have family specifiers (see CMakeLists.txt)
|
||||
# so targets like 10.3 and 12.1 are automatically supported with this list
|
||||
CUDA_ARCH_X86: "7.5 8.0 8.6 8.9 9.0 10.0 12.0"
|
||||
# aarch64-only targets: Orin (8.7), Thor (11.0, CUDA 13+)
|
||||
CUDA_ARCH_AARCH64: "8.0 8.7 8.9 9.0 10.0 11.0 12.0"
|
||||
|
||||
# for CUDA <13, we need to specify all needed targets
|
||||
# some targets (10.3, 12.1) are skipped to limit the wheel size (< 500MB)
|
||||
# please use CUDA 13 wheels or compile yourself on these new devices
|
||||
CUDA_ARCH_X86: "7.5 8.0 8.6 8.9 9.0 10.0 12.0+PTX"
|
||||
# aarch64 only architectures: 8.7 for Orin, 11.0 for Thor (since CUDA 13)
|
||||
CUDA_ARCH_AARCH64: "8.0 8.7 8.9 9.0 10.0 11.0 12.0+PTX"
|
||||
CUDA_ARCH_X86_CU129: "7.5 8.0 8.6 8.9 9.0 10.0 12.0"
|
||||
CUDA_ARCH_AARCH64_CU129: "8.0 8.7 8.9 9.0 10.0 12.0"
|
||||
|
||||
# pre-built mooncake wheels
|
||||
# the manylinux_2_35 wheel has compatibility issue on Ubuntu 24.04
|
||||
# so we use different wheels for the time being
|
||||
MOONCAKE_WHEEL_AARCH64_2_35: "https://vllm-wheels.s3.amazonaws.com/mooncake/mooncake_transfer_engine-0.3.10.post2-0da9dfea3-cp312-cp312-manylinux_2_35_aarch64.whl"
|
||||
MOONCAKE_WHEEL_AARCH64_2_39: "https://vllm-wheels.s3.amazonaws.com/mooncake/mooncake_transfer_engine-0.3.10.post2-0da9dfea3-cp312-cp312-manylinux_2_39_aarch64.whl"
|
||||
MOONCAKE_WHEEL_X86_64: "https://vllm-wheels.s3.amazonaws.com/mooncake/mooncake_transfer_engine-0.3.10.post2-0da9dfea3-cp312-cp312-manylinux_2_35_x86_64.whl"
|
||||
@@ -750,7 +737,7 @@ steps:
|
||||
- "bash tools/vllm-rocm/generate-rocm-wheels-root-index.sh"
|
||||
env:
|
||||
S3_BUCKET: "vllm-wheels"
|
||||
VARIANT: "rocm723"
|
||||
VARIANT: "rocm722"
|
||||
|
||||
# ROCm Job 6: Build ROCm Release Docker Image
|
||||
- label: ":docker: Build release image - x86_64 - ROCm"
|
||||
@@ -859,6 +846,7 @@ steps:
|
||||
allow_failure: true
|
||||
- step: build-cpu-release-image-arm64
|
||||
allow_failure: true
|
||||
if: build.env("NIGHTLY") != "1"
|
||||
|
||||
- label: "Publish release images to DockerHub"
|
||||
depends_on:
|
||||
|
||||
File diff suppressed because it is too large
Load Diff
@@ -13,8 +13,5 @@ INPUT_FILE="$1"
|
||||
# Strip timestamps
|
||||
sed -i 's/^\[[0-9]\{4\}-[0-9]\{2\}-[0-9]\{2\}T[0-9]\{2\}:[0-9]\{2\}:[0-9]\{2\}Z\] //' "$INPUT_FILE"
|
||||
|
||||
# Strip Buildkite inline timestamp markers (ESC _bk;t=<ms> BEL)
|
||||
sed -i 's/\x1B_bk;t=[0-9]*\x07//g' "$INPUT_FILE"
|
||||
|
||||
# Strip colorization
|
||||
sed -i -r 's/\x1B\[[0-9;]*[mK]//g' "$INPUT_FILE"
|
||||
|
||||
@@ -1,178 +1,55 @@
|
||||
#!/bin/bash
|
||||
# Fetch vLLM Buildkite CI logs (public; no login required).
|
||||
# Usage: ./ci-fetch-log.sh <buildkite_job_url> [output_file]
|
||||
# ./ci-fetch-log.sh <build_number> <job_uuid> [output_file]
|
||||
#
|
||||
# Usage:
|
||||
# ci-fetch-log.sh [--soft|--all] --pr [<PR>] failed jobs in the PR's latest
|
||||
# build (current branch if omitted)
|
||||
# ci-fetch-log.sh [--soft|--all] <build_url> failed jobs in that build
|
||||
# ci-fetch-log.sh <job_url> [output] one job; both #<job_uuid> and
|
||||
# ?sid=<id> URL forms work
|
||||
# ci-fetch-log.sh <build> <job_uuid> [output]
|
||||
# Downloads the raw log for a Buildkite job from the public, unauthenticated
|
||||
# /organizations/<org>/pipelines/<pipeline>/builds/<n>/jobs/<uuid>/download
|
||||
# endpoint, then strips ANSI/timestamps via ci-clean-log.sh.
|
||||
#
|
||||
# --soft also fetches soft-failed jobs; --all fetches every finished job.
|
||||
# Saves each log as ci-<build>-<job-name>.log (ANSI/timestamps stripped) and
|
||||
# prints "<file>\t<job name>" per job. [output] is single-job only; "-"
|
||||
# streams to stdout. Existing files are kept; CI_FETCH_LOG_FORCE=1 refetches.
|
||||
# Find <build_number> and <job_uuid> via:
|
||||
# gh pr checks <PR> --repo vllm-project/vllm
|
||||
# Each failing row's URL is .../builds/<build_number>#<job_uuid>.
|
||||
|
||||
set -euo pipefail
|
||||
|
||||
ORG="vllm"
|
||||
PIPELINE="ci"
|
||||
UA="vllm-ci-fetch-log"
|
||||
UUID_RE='[0-9a-f]{8}-[0-9a-f]{4}-[0-9a-f]{4}-[0-9a-f]{4}-[0-9a-f]{12}'
|
||||
|
||||
usage() {
|
||||
sed -n '2,15p' "$0" | sed 's/^# \{0,1\}//'
|
||||
echo "Usage: $0 <buildkite_job_url> [output_file]"
|
||||
echo " $0 <build_number> <job_uuid> [output_file]"
|
||||
exit 1
|
||||
}
|
||||
|
||||
die() {
|
||||
echo "$1" >&2
|
||||
exit 1
|
||||
}
|
||||
if [ $# -lt 1 ]; then usage; fi
|
||||
|
||||
BUILD="" JOB="" SID="" OUT=""
|
||||
SCOPE="failed"
|
||||
|
||||
while :; do
|
||||
case "${1:-}" in
|
||||
--soft) SCOPE="soft" ;;
|
||||
--all) SCOPE="all" ;;
|
||||
*) break ;;
|
||||
esac
|
||||
shift
|
||||
done
|
||||
|
||||
case "${1:-}" in
|
||||
--pr)
|
||||
PR="${2:-}"
|
||||
# gh pr checks exits non-zero when checks are failing; that is the
|
||||
# expected case here.
|
||||
URL=$(gh pr checks ${PR:+"$PR"} --repo vllm-project/vllm 2>/dev/null |
|
||||
grep -oE "https://buildkite.com/${ORG}/${PIPELINE}/builds/[0-9]+" |
|
||||
sort -t/ -k7 -n | tail -1 || true)
|
||||
[ -n "$URL" ] || die "No Buildkite build found via: gh pr checks ${PR:-<current branch>}"
|
||||
BUILD="${URL##*/}"
|
||||
;;
|
||||
https://*)
|
||||
if [[ "$1" == https://* ]]; then
|
||||
BUILD=$(echo "$1" | sed -nE 's#.*/builds/([0-9]+).*#\1#p')
|
||||
JOB=$(echo "$1" | grep -oE "#${UUID_RE}" | head -n 1 | cut -c2- || true)
|
||||
SID=$(echo "$1" | grep -oE "[?&]sid=${UUID_RE}" | head -n 1 | sed 's/.*sid=//' || true)
|
||||
OUT="${2:-}"
|
||||
[ -n "$BUILD" ] || die "Could not parse build number from: $1"
|
||||
;;
|
||||
[0-9]*)
|
||||
[ $# -ge 2 ] || usage
|
||||
JOB=$(echo "$1" | grep -oE '[0-9a-f]{8}-[0-9a-f-]+' | head -n 1)
|
||||
OUT="${2:-ci-${BUILD}-${JOB:0:8}.log}"
|
||||
else
|
||||
if [ $# -lt 2 ]; then usage; fi
|
||||
BUILD="$1"
|
||||
JOB="$2"
|
||||
OUT="${3:-}"
|
||||
;;
|
||||
*)
|
||||
OUT="${3:-ci-${BUILD}-${JOB:0:8}.log}"
|
||||
fi
|
||||
|
||||
if [ -z "$BUILD" ] || [ -z "$JOB" ]; then
|
||||
echo "Could not parse build number or job UUID from: $1" >&2
|
||||
usage
|
||||
;;
|
||||
esac
|
||||
fi
|
||||
|
||||
COOKIES=$(mktemp)
|
||||
JOBS_TSV=$(mktemp)
|
||||
trap 'rm -f "$COOKIES" "$JOBS_TSV"' EXIT
|
||||
trap 'rm -f "$COOKIES"' EXIT
|
||||
|
||||
# Buildkite issues a session cookie on first hit; later requests need it.
|
||||
curl -fsSL -c "$COOKIES" -A "$UA" \
|
||||
# Buildkite issues a session cookie on first hit; subsequent /download needs it.
|
||||
curl -fsSL -c "$COOKIES" -A "vllm-ci-fetch-log" \
|
||||
"https://buildkite.com/${ORG}/${PIPELINE}/builds/${BUILD}" -o /dev/null
|
||||
|
||||
# The build's job list (id, step uuid, state, name) is served as JSON from
|
||||
# the user-facing /data/jobs endpoint. Flatten it to TSV for easy filtering:
|
||||
# job_id step_uuid failed soft_failed finished slug name
|
||||
curl -fsSL -b "$COOKIES" -A "$UA" \
|
||||
"https://buildkite.com/${ORG}/${PIPELINE}/builds/${BUILD}/data/jobs" |
|
||||
python3 -c '
|
||||
import json, re, sys
|
||||
curl -fsSL -b "$COOKIES" -A "vllm-ci-fetch-log" \
|
||||
"https://buildkite.com/organizations/${ORG}/pipelines/${PIPELINE}/builds/${BUILD}/jobs/${JOB}/download" \
|
||||
-o "$OUT"
|
||||
|
||||
data = json.load(sys.stdin)
|
||||
if data.get("has_next_page"):
|
||||
print("warning: job list is paginated; some jobs not shown", file=sys.stderr)
|
||||
for r in data["records"]:
|
||||
if r.get("type") != "script":
|
||||
continue
|
||||
name = (r.get("name") or "").replace("\t", " ").replace("\n", " ")
|
||||
slug = re.sub(r"[^a-z0-9]+", "-", name.lower()).strip("-")[:60]
|
||||
print("\t".join([
|
||||
r["id"],
|
||||
r.get("step_uuid") or "",
|
||||
str(r.get("passed") is False),
|
||||
str(bool(r.get("soft_failed"))),
|
||||
str(bool(r.get("finished_at"))),
|
||||
slug,
|
||||
name,
|
||||
]))
|
||||
' >"$JOBS_TSV" || die "Could not list jobs for build ${BUILD}"
|
||||
bash "$(dirname "$0")/ci-clean-log.sh" "$OUT"
|
||||
|
||||
if [ -n "$SID" ] && [ -z "$JOB" ]; then
|
||||
# The ?sid= in builds/<N>/list URLs is the *step* uuid, not the job uuid.
|
||||
JOB=$(awk -F'\t' -v s="$SID" '$1 == s || $2 == s {print $1; exit}' "$JOBS_TSV")
|
||||
[ -n "$JOB" ] || die "No job matching sid=${SID} in build ${BUILD}"
|
||||
fi
|
||||
|
||||
fetch_job() { # <job_uuid> <output_file>
|
||||
curl -fsSL -b "$COOKIES" -A "$UA" \
|
||||
"https://buildkite.com/organizations/${ORG}/pipelines/${PIPELINE}/builds/${BUILD}/jobs/$1/download" \
|
||||
-o "$2"
|
||||
bash "$(dirname "$0")/ci-clean-log.sh" "$2"
|
||||
}
|
||||
|
||||
if [ -n "$JOB" ]; then
|
||||
# Single-job mode.
|
||||
NAME=$(awk -F'\t' -v j="$JOB" '$1 == j {print $7; exit}' "$JOBS_TSV")
|
||||
SLUG=$(awk -F'\t' -v j="$JOB" '$1 == j {print $6; exit}' "$JOBS_TSV")
|
||||
[ -n "$OUT" ] || OUT="ci-${BUILD}-${SLUG:-${JOB:0:13}}.log"
|
||||
if [ "$OUT" = "-" ]; then
|
||||
TMP=$(mktemp)
|
||||
fetch_job "$JOB" "$TMP"
|
||||
cat "$TMP"
|
||||
rm -f "$TMP"
|
||||
exit 0
|
||||
fi
|
||||
if [ -e "$OUT" ] && [ -z "${CI_FETCH_LOG_FORCE:-}" ]; then
|
||||
die "Refusing to overwrite existing ${OUT} (set CI_FETCH_LOG_FORCE=1 or pass an output path)."
|
||||
fi
|
||||
fetch_job "$JOB" "$OUT"
|
||||
printf '%s\t%s\n' "$OUT" "${NAME:-$JOB}"
|
||||
exit 0
|
||||
fi
|
||||
|
||||
# Build-wide mode: fetch finished jobs matching $SCOPE.
|
||||
[ -z "$OUT" ] || die "[output_file] is only valid when fetching a single job."
|
||||
|
||||
case "$SCOPE" in
|
||||
failed) FILTER='$3 == "True" && $4 == "False" && $5 == "True"' ;;
|
||||
soft) FILTER='$3 == "True" && $5 == "True"' ;;
|
||||
all) FILTER='$5 == "True"' ;;
|
||||
esac
|
||||
|
||||
if [ "$SCOPE" = "failed" ]; then
|
||||
SOFT=$(awk -F'\t' '$3 == "True" && $4 == "True"' "$JOBS_TSV" | wc -l)
|
||||
[ "$SOFT" -eq 0 ] || echo "Skipping ${SOFT} soft-failed job(s); use --soft to include them." >&2
|
||||
fi
|
||||
|
||||
FOUND=0
|
||||
EMITTED=" "
|
||||
while IFS=$'\t' read -r job_id _ _ _ _ slug name; do
|
||||
FOUND=$((FOUND + 1))
|
||||
out="ci-${BUILD}-${slug:-${job_id:0:13}}.log"
|
||||
# Retries share a name with the original job; disambiguate by uuid.
|
||||
case "$EMITTED" in
|
||||
*" $out "*) out="ci-${BUILD}-${slug:-job}-${job_id:0:13}.log" ;;
|
||||
esac
|
||||
EMITTED="${EMITTED}${out} "
|
||||
if [ -e "$out" ] && [ -z "${CI_FETCH_LOG_FORCE:-}" ]; then
|
||||
echo "Keeping existing ${out} (set CI_FETCH_LOG_FORCE=1 to refetch)." >&2
|
||||
elif ! fetch_job "$job_id" "$out"; then
|
||||
echo "Failed to download log for job ${job_id} (${name})." >&2
|
||||
continue
|
||||
fi
|
||||
printf '%s\t%s\n' "$out" "$name"
|
||||
done < <(awk -F'\t' "$FILTER" "$JOBS_TSV")
|
||||
|
||||
if [ "$FOUND" -eq 0 ]; then
|
||||
echo "No matching jobs in build ${BUILD} (scope: ${SCOPE})." >&2
|
||||
fi
|
||||
echo "$OUT"
|
||||
|
||||
@@ -28,18 +28,32 @@
|
||||
###############################################################################
|
||||
set -o pipefail
|
||||
|
||||
# Export Python path for commands that run directly on the host. Containerized
|
||||
# tests set this to /vllm-workspace below so spawned Python processes do not
|
||||
# depend on their current working directory.
|
||||
export PYTHONPATH="${PYTHONPATH:-..}"
|
||||
# Export Python path
|
||||
export PYTHONPATH=".."
|
||||
|
||||
###############################################################################
|
||||
# Helper Functions
|
||||
###############################################################################
|
||||
|
||||
report_docker_usage() {
|
||||
echo "--- Docker usage"
|
||||
docker system df || true
|
||||
cleanup_docker() {
|
||||
# Get Docker's root directory
|
||||
docker_root=$(docker info -f '{{.DockerRootDir}}')
|
||||
if [ -z "$docker_root" ]; then
|
||||
echo "Failed to determine Docker root directory."
|
||||
exit 1
|
||||
fi
|
||||
echo "Docker root directory: $docker_root"
|
||||
|
||||
disk_usage=$(df "$docker_root" | tail -1 | awk '{print $5}' | sed 's/%//')
|
||||
threshold=70
|
||||
if [ "$disk_usage" -gt "$threshold" ]; then
|
||||
echo "Disk usage is above $threshold%. Cleaning up Docker images and volumes..."
|
||||
docker image prune -f
|
||||
docker volume prune -f && docker system prune --force --filter "until=72h" --all
|
||||
echo "Docker images and volumes cleanup completed."
|
||||
else
|
||||
echo "Disk usage is below $threshold%. No cleanup needed."
|
||||
fi
|
||||
}
|
||||
|
||||
cleanup_network() {
|
||||
@@ -54,108 +68,6 @@ cleanup_network() {
|
||||
fi
|
||||
}
|
||||
|
||||
prepare_artifact_image() {
|
||||
if [[ "${VLLM_CI_USE_ARTIFACTS:-0}" != "1" ]]; then
|
||||
return 1
|
||||
fi
|
||||
if ! command -v buildkite-agent >/dev/null 2>&1; then
|
||||
echo "buildkite-agent not found; cannot download ROCm wheel artifact"
|
||||
return 1
|
||||
fi
|
||||
|
||||
local artifact_glob="${VLLM_CI_ARTIFACT_GLOB:-artifacts/vllm-rocm-install/vllm-rocm-install.tar.gz}"
|
||||
local archive=""
|
||||
local metadata_file=""
|
||||
local base_image="${VLLM_CI_BASE_IMAGE:-rocm/vllm-dev:ci_base}"
|
||||
local artifact_image=""
|
||||
local artifact_key=""
|
||||
local base_digest=""
|
||||
local wheel_dir=""
|
||||
local context_dir=""
|
||||
local workspace_dir=""
|
||||
|
||||
artifact_work_dir=$(mktemp -d -t vllm-rocm-artifact.XXXXXX)
|
||||
wheel_dir="${artifact_work_dir}/wheels"
|
||||
context_dir="${artifact_work_dir}/context"
|
||||
workspace_dir="${context_dir}/workspace"
|
||||
mkdir -p "${wheel_dir}" "${context_dir}/wheels" "${workspace_dir}"
|
||||
|
||||
echo "--- Downloading ROCm wheel artifact"
|
||||
if ! buildkite-agent artifact download "${artifact_glob}" "${artifact_work_dir}"; then
|
||||
echo "Failed to download ${artifact_glob}"
|
||||
return 1
|
||||
fi
|
||||
buildkite-agent artifact download \
|
||||
"artifacts/vllm-rocm-install/ci-base-image.txt" \
|
||||
"${artifact_work_dir}" >/dev/null 2>&1 || true
|
||||
|
||||
archive=$(find "${artifact_work_dir}" -name "vllm-rocm-install.tar.gz" -type f | head -1)
|
||||
if [[ -z "${archive}" || ! -f "${archive}" ]]; then
|
||||
echo "ROCm wheel artifact archive was not found"
|
||||
return 1
|
||||
fi
|
||||
|
||||
metadata_file=$(find "${artifact_work_dir}" -name "ci-base-image.txt" -type f | head -1)
|
||||
if [[ -n "${metadata_file}" && -s "${metadata_file}" ]]; then
|
||||
base_image=$(tr -d '[:space:]' < "${metadata_file}")
|
||||
fi
|
||||
|
||||
echo "--- Preparing local ROCm test image"
|
||||
echo "Base image: ${base_image}"
|
||||
docker pull "${base_image}" || return 1
|
||||
base_digest=$(
|
||||
docker image inspect \
|
||||
--format='{{if .RepoDigests}}{{index .RepoDigests 0}}{{else}}{{.Id}}{{end}}' \
|
||||
"${base_image}" 2>/dev/null || printf '%s' "${base_image}"
|
||||
)
|
||||
|
||||
artifact_key=$(
|
||||
{
|
||||
printf 'base-image:%s\n' "${base_digest}"
|
||||
sha256sum "${archive}"
|
||||
} | sha256sum | cut -c1-24
|
||||
)
|
||||
artifact_image="rocm/vllm-ci-artifact:${artifact_key}"
|
||||
|
||||
if docker image inspect "${artifact_image}" >/dev/null 2>&1; then
|
||||
echo "Using existing local ROCm artifact image: ${artifact_image}"
|
||||
image_name="${artifact_image}"
|
||||
return 0
|
||||
fi
|
||||
|
||||
tar -xzf "${archive}" -C "${wheel_dir}" || return 1
|
||||
if ! ls "${wheel_dir}"/*.whl >/dev/null 2>&1; then
|
||||
echo "ROCm wheel artifact did not contain a wheel"
|
||||
return 1
|
||||
fi
|
||||
if [[ ! -d "${wheel_dir}/tests" ]]; then
|
||||
echo "ROCm wheel artifact did not contain the test workspace"
|
||||
return 1
|
||||
fi
|
||||
|
||||
cp "${wheel_dir}"/*.whl "${context_dir}/wheels/" || return 1
|
||||
tar -C "${wheel_dir}" --exclude='*.whl' -cf - . \
|
||||
| tar -C "${workspace_dir}" -xf - || return 1
|
||||
cat > "${context_dir}/Dockerfile" <<'EOF'
|
||||
ARG BASE_IMAGE
|
||||
FROM ${BASE_IMAGE}
|
||||
COPY wheels/ /tmp/vllm-wheels/
|
||||
COPY workspace/ /vllm-workspace/
|
||||
RUN python3 -m pip install --no-deps --force-reinstall /tmp/vllm-wheels/*.whl \
|
||||
&& rm -rf /tmp/vllm-wheels
|
||||
WORKDIR /vllm-workspace
|
||||
EOF
|
||||
|
||||
echo "--- Building local ROCm test image"
|
||||
docker build \
|
||||
--pull=false \
|
||||
--build-arg "BASE_IMAGE=${base_image}" \
|
||||
-t "${artifact_image}" \
|
||||
"${context_dir}" || return 1
|
||||
image_name="${artifact_image}"
|
||||
return 0
|
||||
}
|
||||
|
||||
is_multi_node() {
|
||||
local cmds="$1"
|
||||
# Primary signal: NUM_NODES environment variable set by the pipeline
|
||||
@@ -342,36 +254,20 @@ re_quote_pytest_markers() {
|
||||
echo "--- ROCm info"
|
||||
rocminfo
|
||||
|
||||
# --- Docker status ---
|
||||
report_docker_usage
|
||||
# --- Docker housekeeping ---
|
||||
cleanup_docker
|
||||
|
||||
# --- Pull test image ---
|
||||
echo "--- Pulling container"
|
||||
image_name="${VLLM_CI_FALLBACK_IMAGE:-rocm/vllm-ci:${BUILDKITE_COMMIT:-local}}"
|
||||
artifact_work_dir=""
|
||||
image_name="rocm/vllm-ci:${BUILDKITE_COMMIT}"
|
||||
container_name="rocm_${BUILDKITE_COMMIT}_$(tr -dc A-Za-z0-9 < /dev/urandom | head -c 10; echo)"
|
||||
docker pull "${image_name}"
|
||||
|
||||
remove_docker_container() {
|
||||
if docker container inspect "${container_name}" >/dev/null 2>&1; then
|
||||
docker rm -f "${container_name}" || true
|
||||
fi
|
||||
if [[ "${VLLM_CI_REMOVE_TEST_IMAGE:-0}" == "1" ]]; then
|
||||
docker image rm -f "${image_name}" || true
|
||||
else
|
||||
# Keep images by default so later jobs on the same AMD node can reuse layers.
|
||||
echo "Keeping ROCm test image locally: ${image_name}"
|
||||
fi
|
||||
if [[ -n "${artifact_work_dir}" ]]; then
|
||||
rm -rf "${artifact_work_dir}"
|
||||
fi
|
||||
docker rm -f "${container_name}" || docker image rm -f "${image_name}" || true
|
||||
}
|
||||
trap remove_docker_container EXIT
|
||||
|
||||
if ! prepare_artifact_image; then
|
||||
echo "Using full ROCm CI image: ${image_name}"
|
||||
docker pull "${image_name}" || exit 1
|
||||
fi
|
||||
|
||||
# --- Prepare commands ---
|
||||
echo "--- Running container"
|
||||
|
||||
@@ -379,14 +275,6 @@ HF_CACHE="$(realpath ~)/huggingface"
|
||||
mkdir -p "${HF_CACHE}"
|
||||
HF_MOUNT="/root/.cache/huggingface"
|
||||
|
||||
# Hugging Face Hub defaults to 10s request/download timeouts, while the ROCm
|
||||
# CI image currently raises downloads to 60s. AMD model-test jobs routinely
|
||||
# start from a cold or partially-populated shared cache, and the 60s read cap
|
||||
# has still timed out before pytest reached the vLLM behavior under test.
|
||||
# Keep the CI default explicit and overridable from the Buildkite environment.
|
||||
: "${HF_HUB_DOWNLOAD_TIMEOUT:=300}"
|
||||
: "${HF_HUB_ETAG_TIMEOUT:=60}"
|
||||
|
||||
# ---- Command source selection ----
|
||||
# Prefer VLLM_TEST_COMMANDS (preserves all inner quoting intact).
|
||||
# Fall back to $* for backward compatibility, but warn that inner
|
||||
@@ -426,14 +314,7 @@ fi
|
||||
|
||||
echo "Final commands: $commands"
|
||||
|
||||
MYPYTHONPATH="/vllm-workspace"
|
||||
|
||||
container_job_id="${BUILDKITE_JOB_ID:-${BUILDKITE_PARALLEL_JOB:-0}}"
|
||||
container_job_id="${container_job_id//[^A-Za-z0-9_.-]/_}"
|
||||
container_job_id_short="${container_job_id:0:8}"
|
||||
CONTAINER_TMPDIR="/tmp/vllm-${container_job_id_short}"
|
||||
CONTAINER_CACHE_ROOT="/tmp/vllm-buildkite-${container_job_id}/cache"
|
||||
CONTAINER_PREFLIGHT="mkdir -p \"\$TMPDIR\" \"\$TORCHINDUCTOR_CACHE_DIR\" \"\$TRITON_CACHE_DIR\" \"\$VLLM_CACHE_ROOT\" \"\$XDG_CACHE_HOME\" && python -c \"import encodings, importlib.metadata as im, importlib.util as iu; [im.version(d) for d in ('transformers', 'torch', 'ray', 'sympy', 'markupsafe', 'vllm')]; missing=[m for m in ('torch.utils.model_zoo', 'transformers.models.nomic_bert', 'ray.dag', 'sympy.physics', 'markupsafe._speedups') if iu.find_spec(m) is None]; assert not missing, missing\""
|
||||
MYPYTHONPATH=".."
|
||||
|
||||
# Verify GPU access
|
||||
render_gid=$(getent group render | cut -d: -f3)
|
||||
@@ -510,8 +391,6 @@ else
|
||||
--group-add "$render_gid" \
|
||||
--rm \
|
||||
-e HF_TOKEN \
|
||||
-e "HF_HUB_DOWNLOAD_TIMEOUT=${HF_HUB_DOWNLOAD_TIMEOUT}" \
|
||||
-e "HF_HUB_ETAG_TIMEOUT=${HF_HUB_ETAG_TIMEOUT}" \
|
||||
-e AWS_ACCESS_KEY_ID \
|
||||
-e AWS_SECRET_ACCESS_KEY \
|
||||
-e BUILDKITE_PARALLEL_JOB \
|
||||
@@ -519,15 +398,10 @@ else
|
||||
-v "${HF_CACHE}:${HF_MOUNT}" \
|
||||
-e "HF_HOME=${HF_MOUNT}" \
|
||||
-e "PYTHONPATH=${MYPYTHONPATH}" \
|
||||
-e "TMPDIR=${CONTAINER_TMPDIR}/tmp" \
|
||||
-e "TORCHINDUCTOR_CACHE_DIR=${CONTAINER_CACHE_ROOT}/torchinductor" \
|
||||
-e "TRITON_CACHE_DIR=${CONTAINER_CACHE_ROOT}/triton" \
|
||||
-e "VLLM_CACHE_ROOT=${CONTAINER_CACHE_ROOT}/vllm" \
|
||||
-e "XDG_CACHE_HOME=${CONTAINER_CACHE_ROOT}/xdg" \
|
||||
-e "PYTORCH_ROCM_ARCH=" \
|
||||
--name "${container_name}" \
|
||||
"${image_name}" \
|
||||
/bin/bash -c "${CONTAINER_PREFLIGHT} && ${commands}"
|
||||
/bin/bash -c "${commands}"
|
||||
|
||||
exit_code=$?
|
||||
handle_pytest_exit "$exit_code"
|
||||
|
||||
@@ -37,8 +37,7 @@ function cpu_tests() {
|
||||
pytest -x -v -s tests/kernels/test_onednn.py
|
||||
pytest -x -v -s tests/kernels/attention/test_cpu_attn.py
|
||||
pytest -x -v -s tests/kernels/core/test_cpu_activation.py
|
||||
pytest -x -v -s tests/kernels/moe/test_moe.py -k test_cpu_fused_moe_basic
|
||||
pytest -x -v -s tests/kernels/mamba/cpu/test_cpu_gdn_ops.py"
|
||||
pytest -x -v -s tests/kernels/moe/test_moe.py -k test_cpu_fused_moe_basic"
|
||||
|
||||
# skip tests requiring model downloads if HF_TOKEN is not set
|
||||
# due to rate-limits
|
||||
|
||||
@@ -324,6 +324,23 @@ IMAGE="${IMAGE_TAG_XPU:-${image_name}}"
|
||||
|
||||
echo "Using image: ${IMAGE}"
|
||||
|
||||
if docker image inspect "${IMAGE}" >/dev/null 2>&1; then
|
||||
echo "Image already exists locally, skipping pull"
|
||||
else
|
||||
echo "Image not found locally, waiting for lock..."
|
||||
|
||||
flock /tmp/docker-pull.lock bash -c "
|
||||
if docker image inspect '${IMAGE}' >/dev/null 2>&1; then
|
||||
echo 'Image already pulled by another runner'
|
||||
else
|
||||
echo 'Pulling image...'
|
||||
timeout 900 docker pull '${IMAGE}'
|
||||
fi
|
||||
"
|
||||
|
||||
echo "Pull step completed"
|
||||
fi
|
||||
|
||||
remove_docker_container() {
|
||||
docker rm -f "${container_name}" || true
|
||||
}
|
||||
@@ -335,36 +352,17 @@ if [[ -z "${ZE_AFFINITY_MASK:-}" ]]; then
|
||||
echo "Warning: ZE_AFFINITY_MASK is not set. Proceeding without device affinity." >&2
|
||||
fi
|
||||
|
||||
export CMDS="${commands}"
|
||||
export HF_TOKEN ZE_AFFINITY_MASK
|
||||
|
||||
{
|
||||
flock 9
|
||||
if docker image inspect "${IMAGE}" >/dev/null 2>&1; then
|
||||
echo "Image already exists locally, skipping pull"
|
||||
else
|
||||
echo "Image not found locally, pulling image..."
|
||||
timeout 900 docker pull "${IMAGE}"
|
||||
echo "Pull step completed"
|
||||
fi
|
||||
|
||||
docker create \
|
||||
docker run \
|
||||
--device /dev/dri:/dev/dri \
|
||||
--net=host \
|
||||
--ipc=host \
|
||||
--privileged \
|
||||
-v /dev/dri/by-path:/dev/dri/by-path \
|
||||
-v "${HOME}/.cache/huggingface:/root/.cache/huggingface" \
|
||||
--entrypoint='' \
|
||||
-e HF_TOKEN \
|
||||
-e ZE_AFFINITY_MASK \
|
||||
-e BUILDKITE_PARALLEL_JOB \
|
||||
-e BUILDKITE_PARALLEL_JOB_COUNT \
|
||||
-e CMDS \
|
||||
-v ${HOME}/.cache/huggingface:/root/.cache/huggingface \
|
||||
--entrypoint="" \
|
||||
-e "HF_TOKEN=${HF_TOKEN:-}" \
|
||||
-e "ZE_AFFINITY_MASK=${ZE_AFFINITY_MASK:-}" \
|
||||
-e "CMDS=${commands}" \
|
||||
--name "${container_name}" \
|
||||
"${IMAGE}" \
|
||||
bash -c 'set -e; echo "ZE_AFFINITY_MASK is ${ZE_AFFINITY_MASK:-}"; eval "$CMDS"' \
|
||||
>/dev/null
|
||||
} 9>/tmp/docker-pull.lock
|
||||
|
||||
docker start -a "${container_name}"
|
||||
"${image_name}" \
|
||||
bash -c 'set -e; echo "ZE_AFFINITY_MASK is ${ZE_AFFINITY_MASK:-}"; eval "$CMDS"'
|
||||
|
||||
@@ -1,39 +0,0 @@
|
||||
#!/bin/bash
|
||||
# SPDX-License-Identifier: Apache-2.0
|
||||
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
set -euo pipefail
|
||||
|
||||
REQUIREMENTS_FILE="${KV_CONNECTORS_REQUIREMENTS:-/vllm-workspace/requirements/kv_connectors.txt}"
|
||||
|
||||
uv pip install --system -r "${REQUIREMENTS_FILE}"
|
||||
|
||||
NIXL_METADATA=$(python3 - <<'PY'
|
||||
import importlib.metadata as metadata
|
||||
|
||||
import torch
|
||||
|
||||
cuda_version = torch.version.cuda
|
||||
if cuda_version is None:
|
||||
raise SystemExit("torch.version.cuda is not set")
|
||||
|
||||
print(cuda_version.split(".", 1)[0], metadata.version("nixl"))
|
||||
PY
|
||||
)
|
||||
read -r CUDA_MAJOR NIXL_VERSION <<<"${NIXL_METADATA}"
|
||||
|
||||
# nixl>=1.1.0 can install multiple CUDA wheel variants. Keep only the variant
|
||||
# matching this CI image so nixl_ep_cpp links against the available libcudart.
|
||||
uv pip uninstall --system nixl-cu12 nixl-cu13 2>/dev/null || true
|
||||
uv pip install --system --no-deps "nixl-cu${CUDA_MAJOR}==${NIXL_VERSION}"
|
||||
|
||||
python3 - <<'PY'
|
||||
import importlib.metadata as metadata
|
||||
|
||||
for package_name in ("nixl", "nixl-cu12", "nixl-cu13"):
|
||||
try:
|
||||
version = metadata.version(package_name)
|
||||
except metadata.PackageNotFoundError:
|
||||
version = "not installed"
|
||||
print(f"{package_name}: {version}")
|
||||
PY
|
||||
@@ -1,187 +0,0 @@
|
||||
#!/usr/bin/env bash
|
||||
set -euo pipefail
|
||||
|
||||
MODE="${1:-}"
|
||||
|
||||
if [[ "$MODE" != "style-clippy" && "$MODE" != "test" ]]; then
|
||||
echo "Usage: $0 {style-clippy|test}" >&2
|
||||
exit 2
|
||||
fi
|
||||
|
||||
ROOT_DIR="$(git rev-parse --show-toplevel)"
|
||||
cd "$ROOT_DIR"
|
||||
|
||||
export CARGO_TERM_COLOR="${CARGO_TERM_COLOR:-always}"
|
||||
export CARGO_HOME="${CARGO_HOME:-$HOME/.cargo}"
|
||||
export RUSTUP_HOME="${RUSTUP_HOME:-$HOME/.rustup}"
|
||||
export PATH="$CARGO_HOME/bin:$PATH"
|
||||
|
||||
log_section() {
|
||||
echo "--- $*"
|
||||
}
|
||||
|
||||
install_protoc() {
|
||||
if command -v protoc >/dev/null 2>&1; then
|
||||
return
|
||||
fi
|
||||
|
||||
local version="${PROTOC_VERSION:-31.1}"
|
||||
local arch
|
||||
case "$(uname -m)" in
|
||||
x86_64)
|
||||
arch="x86_64"
|
||||
;;
|
||||
aarch64|arm64)
|
||||
arch="aarch_64"
|
||||
;;
|
||||
*)
|
||||
echo "Unsupported protoc architecture: $(uname -m)" >&2
|
||||
return 1
|
||||
;;
|
||||
esac
|
||||
|
||||
local url="https://github.com/protocolbuffers/protobuf/releases/download/v${version}/protoc-${version}-linux-${arch}.zip"
|
||||
local tmp_dir
|
||||
tmp_dir="$(mktemp -d)"
|
||||
|
||||
log_section "Installing protoc ${version}"
|
||||
curl -L --proto '=https' --tlsv1.2 -sSf "$url" -o "$tmp_dir/protoc.zip"
|
||||
mkdir -p "$CARGO_HOME/bin"
|
||||
unzip -q "$tmp_dir/protoc.zip" bin/protoc 'include/*' -d "$CARGO_HOME"
|
||||
chmod +x "$CARGO_HOME/bin/protoc"
|
||||
rm -rf "$tmp_dir"
|
||||
}
|
||||
|
||||
rust_toolchain() {
|
||||
awk -F '"' '/channel[[:space:]]*=/ { print $2; exit }' rust-toolchain.toml
|
||||
}
|
||||
|
||||
install_rust_toolchain() {
|
||||
log_section "Installing Rust toolchain"
|
||||
if ! command -v rustup >/dev/null 2>&1; then
|
||||
curl --proto '=https' --tlsv1.2 -sSf https://sh.rustup.rs \
|
||||
| sh -s -- -y --profile minimal --default-toolchain none
|
||||
fi
|
||||
|
||||
local toolchain
|
||||
toolchain="$(rust_toolchain)"
|
||||
rustup toolchain install "$toolchain" --profile minimal --component rustfmt,clippy
|
||||
rustup component add --toolchain "$toolchain" rustfmt clippy
|
||||
}
|
||||
|
||||
install_cargo_binstall() {
|
||||
if command -v cargo-binstall >/dev/null 2>&1; then
|
||||
return
|
||||
fi
|
||||
|
||||
log_section "Installing cargo-binstall"
|
||||
curl -L --proto '=https' --tlsv1.2 -sSf \
|
||||
https://raw.githubusercontent.com/cargo-bins/cargo-binstall/main/install-from-binstall-release.sh \
|
||||
| bash
|
||||
}
|
||||
|
||||
install_cargo_sort() {
|
||||
if command -v cargo-sort >/dev/null 2>&1; then
|
||||
return
|
||||
fi
|
||||
|
||||
log_section "Installing cargo-sort"
|
||||
install_cargo_binstall
|
||||
cargo binstall --no-confirm cargo-sort
|
||||
}
|
||||
|
||||
install_cargo_nextest() {
|
||||
if command -v cargo-nextest >/dev/null 2>&1; then
|
||||
return
|
||||
fi
|
||||
|
||||
log_section "Installing cargo-nextest"
|
||||
install_cargo_binstall
|
||||
cargo binstall --no-confirm --secure cargo-nextest
|
||||
}
|
||||
|
||||
install_uv() {
|
||||
if command -v uv >/dev/null 2>&1; then
|
||||
return
|
||||
fi
|
||||
|
||||
log_section "Installing uv"
|
||||
curl -LsSf --proto '=https' --tlsv1.2 https://astral.sh/uv/install.sh \
|
||||
| env UV_INSTALL_DIR="$CARGO_HOME/bin" sh
|
||||
}
|
||||
|
||||
setup_pyo3_python() {
|
||||
local python_version="${PYO3_PYTHON_VERSION:-3.12}"
|
||||
|
||||
log_section "Installing Python ${python_version} for PyO3 tests"
|
||||
uv python install "$python_version"
|
||||
PYO3_PYTHON="$(uv python find \
|
||||
--managed-python \
|
||||
--no-project \
|
||||
--resolve-links \
|
||||
"$python_version")"
|
||||
export PYO3_PYTHON
|
||||
|
||||
local python_libdir
|
||||
python_libdir="$("$PYO3_PYTHON" - <<'PY'
|
||||
import pathlib
|
||||
import sysconfig
|
||||
|
||||
libdir = pathlib.Path(sysconfig.get_config_var("LIBDIR"))
|
||||
ldlibrary = sysconfig.get_config_var("LDLIBRARY")
|
||||
assert sysconfig.get_config_var("Py_ENABLE_SHARED") == 1
|
||||
assert ldlibrary
|
||||
assert (libdir / ldlibrary).exists(), libdir / ldlibrary
|
||||
print(libdir)
|
||||
PY
|
||||
)"
|
||||
|
||||
export LD_LIBRARY_PATH="${python_libdir}:${LD_LIBRARY_PATH:-}"
|
||||
export LIBRARY_PATH="${python_libdir}:${LIBRARY_PATH:-}"
|
||||
}
|
||||
|
||||
run_style_clippy() {
|
||||
install_cargo_sort
|
||||
|
||||
log_section "Checking Rust formatting"
|
||||
cargo fmt --manifest-path rust/Cargo.toml --all -- --check
|
||||
|
||||
log_section "Checking Cargo.toml ordering"
|
||||
cargo sort --workspace --check rust
|
||||
|
||||
log_section "Running clippy"
|
||||
cargo clippy \
|
||||
--manifest-path rust/Cargo.toml \
|
||||
--workspace \
|
||||
--all-targets \
|
||||
--all-features \
|
||||
--locked \
|
||||
-- \
|
||||
-D warnings
|
||||
}
|
||||
|
||||
run_tests() {
|
||||
install_uv
|
||||
setup_pyo3_python
|
||||
install_cargo_nextest
|
||||
|
||||
log_section "Running cargo nextest"
|
||||
cargo nextest run \
|
||||
--manifest-path rust/Cargo.toml \
|
||||
--workspace \
|
||||
--all-features \
|
||||
--locked \
|
||||
--no-fail-fast
|
||||
}
|
||||
|
||||
install_protoc
|
||||
install_rust_toolchain
|
||||
|
||||
case "$MODE" in
|
||||
style-clippy)
|
||||
run_style_clippy
|
||||
;;
|
||||
test)
|
||||
run_tests
|
||||
;;
|
||||
esac
|
||||
@@ -49,7 +49,6 @@ for BACK in "${BACKENDS[@]}"; do
|
||||
--data-parallel-size 2 \
|
||||
--enable-expert-parallel \
|
||||
--enable-eplb \
|
||||
--eplb-config '{"use_async": false}' \
|
||||
--trust-remote-code \
|
||||
--max-model-len 2048 \
|
||||
--all2all-backend "$BACK" \
|
||||
|
||||
@@ -48,7 +48,7 @@ for BACK in "${BACKENDS[@]}"; do
|
||||
--enforce-eager \
|
||||
--enable-eplb \
|
||||
--all2all-backend "$BACK" \
|
||||
--eplb-config '{"window_size":10, "step_interval":100, "num_redundant_experts":0, "log_balancedness":true, "use_async":false}' \
|
||||
--eplb-config '{"window_size":10, "step_interval":100, "num_redundant_experts":0, "log_balancedness":true}' \
|
||||
--tensor-parallel-size "${TENSOR_PARALLEL_SIZE}" \
|
||||
--data-parallel-size "${DATA_PARALLEL_SIZE}" \
|
||||
--enable-expert-parallel \
|
||||
|
||||
@@ -70,7 +70,7 @@ echo "============================================"
|
||||
# ---- Install bfcl-eval if missing ----
|
||||
if ! python3 -c "import bfcl_eval" 2>/dev/null; then
|
||||
echo "Installing bfcl-eval..."
|
||||
uv pip install "bfcl-eval>=2025.10.20.1,<2026"
|
||||
pip install "bfcl-eval>=2025.10.20.1,<2026"
|
||||
fi
|
||||
|
||||
# ---- Cleanup handler ----
|
||||
@@ -100,7 +100,7 @@ SERVE_ARGS=(
|
||||
--tensor-parallel-size "$TP_SIZE"
|
||||
--max-model-len "$MAX_MODEL_LEN"
|
||||
--enforce-eager
|
||||
--enable-prefix-caching
|
||||
--no-enable-prefix-caching
|
||||
)
|
||||
|
||||
# Append reasoning parser if specified
|
||||
|
||||
+239
-345
@@ -88,16 +88,16 @@
|
||||
# - Do NOT remove `VLLM_WORKER_MULTIPROC_METHOD=spawn` setting as ROCm requires this for certain models to function. #
|
||||
# * [Transformers Nightly Models]: Whisper needs `VLLM_WORKER_MULTIPROC_METHOD=spawn` to avoid deadlock. #
|
||||
# * [Plugin Tests (2 GPUs)]: #
|
||||
# - {`pytest -v -s plugins_tests/test_oot_registration_online.py`}: It needs a clean process #
|
||||
# - {`pytest -v -s plugins_tests/test_oot_registration_offline.py`}: It needs a clean process #
|
||||
# - {`pytest -v -s plugins_tests/lora_resolvers`}: Unit tests for in-tree lora resolver plugins #
|
||||
# - {`pytest -v -s entrypoints/openai/test_oot_registration.py`}: It needs a clean process #
|
||||
# - {`pytest -v -s models/test_oot_registration.py`}: It needs a clean process #
|
||||
# - {`pytest -v -s plugins/lora_resolvers`}: Unit tests for in-tree lora resolver plugins #
|
||||
# * [LoRA TP (Distributed)]: #
|
||||
# - There is some Tensor Parallelism related processing logic in LoRA that requires multi-GPU testing for validation. #
|
||||
# - {`pytest -v -s -x lora/test_gptoss_tp.py`}: Disabled for now because MXFP4 backend on non-cuda platform doesn't support #
|
||||
# LoRA yet. #
|
||||
# * [Distributed Tests (NxGPUs)(HW-TAG)]: Don't test llama model here, it seems hf implementation is buggy. See: #
|
||||
# https://github.com/vllm-project/vllm/pull/5689 #
|
||||
# * [Distributed Tests (NxGPUs)(HW-TAG)]: Some old E2E tests were removed in https://github.com/vllm-project/vllm/pull/33293 #
|
||||
# * [Distributed Tests (NxGPUs)(HW-TAG)]: Some old E2E tests were removed in https://github.com/vllm-project/vllm/pull/33293 #
|
||||
# in favor of new tests in fusions_e2e. We avoid replicating the new jobs in #
|
||||
# this file as it's deprecated. #
|
||||
# #
|
||||
@@ -139,6 +139,19 @@ steps:
|
||||
- pytest models/multimodal -v -s -m 'distributed(num_gpus=2)' --ignore models/multimodal/generation/test_whisper.py
|
||||
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest models/multimodal/generation/test_whisper.py -v -s -m 'distributed(num_gpus=2)'
|
||||
|
||||
#-------------------------------------------------------- mi250 · benchmarks ---------------------------------------------------------#
|
||||
|
||||
- label: Benchmarks # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
|
||||
agent_pool: mi250_1
|
||||
working_dir: "/vllm-workspace/.buildkite"
|
||||
source_file_dependencies:
|
||||
- benchmarks/
|
||||
- vllm/platforms/rocm.py
|
||||
commands:
|
||||
- bash scripts/run-benchmarks.sh
|
||||
|
||||
#---------------------------------------------------------- mi250 · compile ----------------------------------------------------------#
|
||||
|
||||
- label: PyTorch Compilation Unit Tests # TBD
|
||||
@@ -315,6 +328,24 @@ steps:
|
||||
- pytest -v -s distributed/test_pp_cudagraph.py
|
||||
- pytest -v -s distributed/test_pipeline_parallel.py
|
||||
|
||||
#---------------------------------------------------------- mi250 · engine -----------------------------------------------------------#
|
||||
|
||||
- label: Engine # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
|
||||
agent_pool: mi250_1
|
||||
optional: true
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/engine
|
||||
- tests/test_sequence
|
||||
- tests/test_config
|
||||
- tests/test_logger
|
||||
- tests/test_vllm_port
|
||||
commands:
|
||||
- pytest -v -s engine test_sequence.py test_config.py test_logger.py test_vllm_port.py
|
||||
|
||||
#----------------------------------------------------------- mi250 · evals -----------------------------------------------------------#
|
||||
|
||||
- label: Multi-Modal Accuracy Eval (Small Models) # TBD
|
||||
@@ -398,7 +429,7 @@ steps:
|
||||
- tests/kernels/helion/
|
||||
- vllm/platforms/rocm.py
|
||||
commands:
|
||||
- pip install helion==1.1.0
|
||||
- pip install helion==1.0.0
|
||||
- pytest -v -s kernels/helion/
|
||||
|
||||
- label: Kernels Mamba Test # TBD
|
||||
@@ -431,7 +462,30 @@ steps:
|
||||
commands:
|
||||
- pytest -v -s lora --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --ignore=lora/test_chatglm3_tp.py --ignore=lora/test_llama_tp.py --ignore=lora/test_qwen3_with_multi_loras.py --ignore=lora/test_olmoe_tp.py --ignore=lora/test_deepseekv2_tp.py --ignore=lora/test_gptoss_tp.py --ignore=lora/test_qwen3moe_tp.py --ignore=lora/test_qwen35_densemodel_lora.py
|
||||
|
||||
#------------------------------------------------------ mi250 · models / basic -------------------------------------------------------#
|
||||
#------------------------------------------------------ mi250 · model_executor -------------------------------------------------------#
|
||||
|
||||
- label: Model Executor # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
|
||||
agent_pool: mi250_1
|
||||
optional: true
|
||||
torch_nightly: true
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/engine/arg_utils.py
|
||||
- vllm/config/model.py
|
||||
- vllm/model_executor
|
||||
- tests/model_executor
|
||||
- tests/entrypoints/openai/completion/test_tensorizer_entrypoint.py
|
||||
- vllm/_aiter_ops.py
|
||||
- vllm/platforms/rocm.py
|
||||
commands:
|
||||
- apt-get update && apt-get install -y curl libsodium23
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s model_executor -m '(not slow_test)'
|
||||
- pytest -v -s entrypoints/openai/completion/test_tensorizer_entrypoint.py
|
||||
|
||||
#---------------------------------------------------------- mi250 · models -----------------------------------------------------------#
|
||||
|
||||
- label: Basic Models Test (Other CPU) # TBD
|
||||
timeout_in_minutes: 180
|
||||
@@ -492,8 +546,6 @@ steps:
|
||||
commands:
|
||||
- pytest -v -s models/test_terratorch.py models/test_transformers.py models/test_registry.py
|
||||
|
||||
#----------------------------------------------------- mi250 · models / language -----------------------------------------------------#
|
||||
|
||||
- label: Language Models Test (MTEB) # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
|
||||
@@ -538,8 +590,6 @@ steps:
|
||||
- pip freeze | grep -E 'torch'
|
||||
- pytest -v -s models/language -m 'core_model and slow_test' --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
|
||||
|
||||
#---------------------------------------------------- mi250 · models / multimodal ----------------------------------------------------#
|
||||
|
||||
- label: Multi-Modal Models (Extended Generation 2) # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
|
||||
@@ -617,9 +667,9 @@ steps:
|
||||
- pytest -v -s plugins_tests/test_scheduler_plugins.py
|
||||
- pip install -e ./plugins/vllm_add_dummy_model
|
||||
- pytest -v -s distributed/test_distributed_oot.py
|
||||
- pytest -v -s plugins_tests/test_oot_registration_online.py # it needs a clean process
|
||||
- pytest -v -s plugins_tests/test_oot_registration_offline.py # it needs a clean process
|
||||
- pytest -v -s plugins_tests/lora_resolvers # unit tests for in-tree lora resolver plugins
|
||||
- pytest -v -s entrypoints/openai/chat_completion/test_oot_registration.py
|
||||
- pytest -v -s models/test_oot_registration.py
|
||||
- pytest -v -s plugins/lora_resolvers
|
||||
|
||||
#------------------------------------------------------------ mi250 · v1 -------------------------------------------------------------#
|
||||
|
||||
@@ -804,7 +854,7 @@ steps:
|
||||
- vllm/platforms/rocm.py
|
||||
commands:
|
||||
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
|
||||
- ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/spec_decode_acceptance_test.sh
|
||||
- ATTENTION_BACKEND=ROCM_ATTN bash v1/kv_connector/nixl_integration/spec_decode_acceptance_test.sh
|
||||
|
||||
- label: V1 e2e (2 GPUs) # TBD
|
||||
timeout_in_minutes: 180
|
||||
@@ -830,7 +880,7 @@ steps:
|
||||
- vllm/platforms/rocm.py
|
||||
commands:
|
||||
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
|
||||
- ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
|
||||
- ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
|
||||
|
||||
#------------------------------------------------------------- mi250 · misc ------------------------------------------------------------#
|
||||
|
||||
@@ -890,10 +940,10 @@ steps:
|
||||
- vllm/
|
||||
- tests/basic_correctness/test_basic_correctness
|
||||
- tests/basic_correctness/test_cpu_offload
|
||||
- tests/basic_correctness/test_mem.py
|
||||
- tests/basic_correctness/test_cumem.py
|
||||
commands:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s basic_correctness/test_mem.py
|
||||
- pytest -v -s basic_correctness/test_cumem.py
|
||||
- pytest -v -s basic_correctness/test_basic_correctness.py
|
||||
- pytest -v -s basic_correctness/test_cpu_offload.py
|
||||
|
||||
@@ -926,6 +976,18 @@ steps:
|
||||
|
||||
#-------------------------------------------------------- mi300 · benchmarks ---------------------------------------------------------#
|
||||
|
||||
- label: Benchmarks # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
|
||||
agent_pool: mi300_1
|
||||
optional: true
|
||||
working_dir: "/vllm-workspace/.buildkite"
|
||||
source_file_dependencies:
|
||||
- benchmarks/
|
||||
- vllm/platforms/rocm.py
|
||||
commands:
|
||||
- bash scripts/run-benchmarks.sh
|
||||
|
||||
- label: Benchmarks CLI Test # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
|
||||
@@ -1187,88 +1249,7 @@ steps:
|
||||
|
||||
#-------------------------------------------------------- mi300 · entrypoints --------------------------------------------------------#
|
||||
|
||||
- label: Entrypoints Unit Tests # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
|
||||
agent_pool: mi300_1
|
||||
fast_check: true
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/entrypoints
|
||||
- tests/entrypoints/unit_tests
|
||||
- tests/entrypoints/weight_transfer
|
||||
- vllm/platforms/rocm.py
|
||||
commands:
|
||||
- pytest -v -s entrypoints/unit_tests
|
||||
- pytest -v -s entrypoints/weight_transfer
|
||||
|
||||
- label: Entrypoints Integration (LLM) # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
|
||||
agent_pool: mi300_1
|
||||
optional: true
|
||||
fast_check: true
|
||||
torch_nightly: true
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/entrypoints/llm
|
||||
commands:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s entrypoints/llm --ignore=entrypoints/llm/test_generate.py --ignore=entrypoints/llm/test_collective_rpc.py --ignore=entrypoints/llm/offline_mode
|
||||
- pytest -v -s entrypoints/llm/test_generate.py # it needs a clean process
|
||||
- pytest -v -s entrypoints/llm/offline_mode # Needs to avoid interference with other tests
|
||||
|
||||
- label: Entrypoints Integration (API Server) # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
|
||||
agent_pool: mi300_1
|
||||
optional: true
|
||||
fast_check: true
|
||||
torch_nightly: true
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/entrypoints/serve
|
||||
commands:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s entrypoints/serve --ignore=entrypoints/serve/dev/rpc
|
||||
- PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/serve/dev/rpc
|
||||
|
||||
- label: Entrypoints Integration (API Server OpenAI - Part 1) # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
|
||||
agent_pool: mi300_1
|
||||
fast_check: true
|
||||
torch_nightly: true
|
||||
optional: true
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/entrypoints/openai
|
||||
- tests/entrypoints/test_chat_utils
|
||||
commands:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s entrypoints/openai/ --ignore=entrypoints/openai/completion --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/correctness
|
||||
|
||||
- label: Entrypoints Integration (API Server OpenAI - Part 2) # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
|
||||
agent_pool: mi300_1
|
||||
fast_check: true
|
||||
torch_nightly: true
|
||||
optional: true
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/entrypoints/openai
|
||||
- tests/entrypoints/test_chat_utils
|
||||
commands:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s entrypoints/openai/chat_completion
|
||||
- pytest -v -s entrypoints/openai/completion --ignore=entrypoints/openai/completion/test_tensorizer_entrypoint.py
|
||||
|
||||
- label: Entrypoints Integration (API Server Generate) # TBD
|
||||
- label: Entrypoints Integration (API Server 2) # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
|
||||
agent_pool: mi300_1
|
||||
@@ -1278,30 +1259,63 @@ steps:
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/entrypoints/rpc
|
||||
- tests/entrypoints/serve/instrumentator
|
||||
- tests/tool_use
|
||||
- tests/entrypoints/tool_parsers
|
||||
- tests/entrypoints/anthropic
|
||||
- tests/entrypoints/generate
|
||||
commands:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s entrypoints/serve/instrumentator
|
||||
- PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/rpc
|
||||
- pytest -v -s tool_use
|
||||
- pytest -v -s entrypoints/tool_parsers
|
||||
- pytest -v -s entrypoints/generate
|
||||
- pytest -v -s entrypoints/anthropic
|
||||
|
||||
- label: Entrypoints Integration (Responses API) # TBD
|
||||
- label: Entrypoints Integration (API Server openai - Part 1) # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
|
||||
agent_pool: mi300_1
|
||||
fast_check: true
|
||||
torch_nightly: true
|
||||
optional: true
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/entrypoints/openai/responses
|
||||
- tests/entrypoints/openai
|
||||
- tests/entrypoints/test_chat_utils
|
||||
commands:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s entrypoints/openai/responses
|
||||
- pytest -v -s entrypoints/openai/chat_completion --ignore=entrypoints/openai/chat_completion/test_chat_with_tool_reasoning.py --ignore=entrypoints/openai/chat_completion/test_oot_registration.py
|
||||
|
||||
- label: Entrypoints Integration (API Server openai - Part 2) # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
|
||||
agent_pool: mi300_1
|
||||
optional: true
|
||||
fast_check: true
|
||||
torch_nightly: true
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/entrypoints/openai
|
||||
- tests/entrypoints/test_chat_utils
|
||||
commands:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s entrypoints/openai/completion --ignore=entrypoints/openai/completion/test_tensorizer_entrypoint.py
|
||||
- pytest -v -s entrypoints/test_chat_utils.py
|
||||
|
||||
- label: Entrypoints Integration (API Server openai - Part 3) # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
|
||||
agent_pool: mi300_1
|
||||
optional: true
|
||||
fast_check: true
|
||||
torch_nightly: true
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/entrypoints/openai
|
||||
- tests/entrypoints/test_chat_utils
|
||||
commands:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s entrypoints/openai --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/completion --ignore=entrypoints/openai/correctness/ --ignore=entrypoints/openai/tool_parsers/ --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/test_multi_api_servers.py
|
||||
|
||||
- label: Entrypoints Integration (Speech to Text) # TBD
|
||||
timeout_in_minutes: 180
|
||||
@@ -1317,19 +1331,23 @@ steps:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s entrypoints/speech_to_text
|
||||
|
||||
- label: Entrypoints Integration (Multimodal)
|
||||
- label: Entrypoints Integration (LLM) # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
|
||||
agent_pool: mi300_1
|
||||
optional: true
|
||||
fast_check: true
|
||||
torch_nightly: true
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/entrypoints/multimodal
|
||||
- tests/entrypoints/llm
|
||||
- tests/entrypoints/offline_mode
|
||||
commands:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s entrypoints/multimodal
|
||||
- pytest -v -s entrypoints/llm --ignore=entrypoints/llm/test_generate.py --ignore=entrypoints/llm/test_collective_rpc.py
|
||||
- pytest -v -s entrypoints/llm/test_generate.py
|
||||
- pytest -v -s entrypoints/offline_mode
|
||||
|
||||
- label: Entrypoints Integration (Pooling) # TBD
|
||||
timeout_in_minutes: 180
|
||||
@@ -1345,15 +1363,43 @@ steps:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s entrypoints/pooling
|
||||
|
||||
- label: Entrypoints Integration (Responses API) # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
|
||||
agent_pool: mi300_1
|
||||
fast_check: true
|
||||
torch_nightly: true
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/entrypoints/openai/responses
|
||||
commands:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s entrypoints/openai/responses
|
||||
|
||||
- label: Entrypoints Unit Tests # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
|
||||
agent_pool: mi300_1
|
||||
fast_check: true
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/entrypoints
|
||||
- tests/entrypoints/
|
||||
- vllm/platforms/rocm.py
|
||||
commands:
|
||||
- pytest -v -s entrypoints/openai/tool_parsers
|
||||
- pytest -v -s entrypoints/ --ignore=entrypoints/llm --ignore=entrypoints/rpc --ignore=entrypoints/sleep --ignore=entrypoints/serve/instrumentator --ignore=entrypoints/openai --ignore=entrypoints/offline_mode --ignore=entrypoints/test_chat_utils.py --ignore=entrypoints/pooling
|
||||
|
||||
- label: OpenAI API correctness # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
|
||||
agent_pool: mi300_1
|
||||
optional: true
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- csrc/
|
||||
- vllm/entrypoints/openai/
|
||||
- vllm/model_executor/models/whisper.py
|
||||
- vllm/model_executor/layers/
|
||||
- vllm/v1/attention/backends/
|
||||
- vllm/v1/attention/selector.py
|
||||
@@ -1407,7 +1453,6 @@ steps:
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
|
||||
agent_pool: mi300_1
|
||||
optional: true
|
||||
working_dir: "/vllm-workspace/.buildkite/lm-eval-harness"
|
||||
source_file_dependencies:
|
||||
- csrc/
|
||||
@@ -1460,7 +1505,7 @@ steps:
|
||||
commands:
|
||||
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-mi3xx-fp8-and-mixed.txt
|
||||
|
||||
- label: DeepSeek V2-Lite Sync EPLB Accuracy (4xH100-4xMI300) # TBD
|
||||
- label: DeepSeek V2-Lite Accuracy (4xH100-4xMI300) # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
|
||||
agent_pool: mi300_4
|
||||
@@ -1502,7 +1547,7 @@ steps:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-large.txt --tp-size=4
|
||||
|
||||
- label: Qwen3-30B-A3B-FP8-block Sync EPLB Accuracy (4xH100-4xMI300) # TBD
|
||||
- label: Qwen3-30B-A3B-FP8-block Accuracy (4xH100-4xMI300) # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
|
||||
agent_pool: mi300_4
|
||||
@@ -1714,30 +1759,7 @@ steps:
|
||||
- pytest -v -s -x lora/test_gptoss_tp.py
|
||||
- pytest -v -s -x lora/test_qwen35_densemodel_lora.py
|
||||
|
||||
#------------------------------------------------------ mi300 · model_executor -------------------------------------------------------#
|
||||
|
||||
- label: Model Executor # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
|
||||
agent_pool: mi300_1
|
||||
optional: true
|
||||
torch_nightly: true
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/engine/arg_utils.py
|
||||
- vllm/config/model.py
|
||||
- vllm/model_executor
|
||||
- tests/model_executor
|
||||
- tests/entrypoints/openai/completion/test_tensorizer_entrypoint.py
|
||||
- vllm/_aiter_ops.py
|
||||
- vllm/platforms/rocm.py
|
||||
commands:
|
||||
- apt-get update && apt-get install -y curl libsodium23
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s model_executor -m '(not slow_test)'
|
||||
- pytest -v -s entrypoints/openai/completion/test_tensorizer_entrypoint.py
|
||||
|
||||
#----------------------------------------------------- mi300 · models / language -----------------------------------------------------#
|
||||
#---------------------------------------------------------- mi300 · models -----------------------------------------------------------#
|
||||
|
||||
- label: Language Models Test (Extended Pooling) # TBD
|
||||
timeout_in_minutes: 180
|
||||
@@ -1765,8 +1787,6 @@ steps:
|
||||
- pip freeze | grep -E 'torch'
|
||||
- pytest -v -s models/language -m 'core_model and (not slow_test)'
|
||||
|
||||
#---------------------------------------------------- mi300 · models / multimodal ----------------------------------------------------#
|
||||
|
||||
- label: Multi-Modal Models (Extended Generation 1) # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
|
||||
@@ -1872,11 +1892,10 @@ steps:
|
||||
- pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
|
||||
- pytest -v -s models/multimodal/processing/test_tensor_schema.py
|
||||
|
||||
- label: Multi-Modal Processor (CPU) %N # TBD
|
||||
- label: Multi-Modal Processor (CPU) # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
|
||||
agent_pool: mi300_1
|
||||
parallelism: 4
|
||||
no_gpu: true
|
||||
optional: true
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
@@ -1886,9 +1905,7 @@ steps:
|
||||
- tests/models/registry.py
|
||||
commands:
|
||||
- pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
|
||||
- pytest -v -s models/multimodal/processing --ignore models/multimodal/processing/test_tensor_schema.py --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
|
||||
|
||||
#----------------------------------------------------- mi300 · models / quantized -----------------------------------------------------#
|
||||
- pytest -v -s models/multimodal/processing --ignore models/multimodal/processing/test_tensor_schema.py
|
||||
|
||||
- label: Quantized Models Test # TBD
|
||||
timeout_in_minutes: 180
|
||||
@@ -1904,31 +1921,7 @@ steps:
|
||||
commands:
|
||||
- pytest -v -s models/quantization
|
||||
|
||||
#-------------------------------------------------- mi300 · models / transformers ---------------------------------------------------#
|
||||
|
||||
- label: Transformers Nightly Models (Shardable) %N # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
|
||||
agent_pool: mi300_1
|
||||
parallelism: 4
|
||||
optional: true
|
||||
working_dir: "/vllm-workspace/"
|
||||
source_file_dependencies:
|
||||
- vllm/model_executor/models/
|
||||
- vllm/model_executor/model_loader/
|
||||
- vllm/multimodal/
|
||||
- vllm/model_executor/layers/
|
||||
- vllm/v1/attention/backends/
|
||||
- vllm/v1/attention/selector.py
|
||||
- vllm/_aiter_ops.py
|
||||
- vllm/platforms/rocm.py
|
||||
- tests/models/
|
||||
commands:
|
||||
- pip install --upgrade git+https://github.com/huggingface/transformers
|
||||
- pytest -v -s tests/models/test_initialization.py --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
|
||||
- pytest -v -s tests/models/multimodal/processing/ --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
|
||||
|
||||
- label: Transformers Nightly Models (Single) # TBD
|
||||
- label: Transformers Nightly Models # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
|
||||
agent_pool: mi300_1
|
||||
@@ -1947,7 +1940,9 @@ steps:
|
||||
- examples/
|
||||
commands:
|
||||
- pip install --upgrade git+https://github.com/huggingface/transformers
|
||||
- pytest -v -s tests/models/test_initialization.py
|
||||
- pytest -v -s tests/models/test_transformers.py
|
||||
- pytest -v -s tests/models/multimodal/processing/
|
||||
- pytest -v -s tests/models/multimodal/test_mapping.py
|
||||
- python3 examples/basic/offline_inference/chat.py
|
||||
- python3 examples/generate/multimodal/vision_language_offline.py --model-type qwen2_5_vl
|
||||
@@ -2008,7 +2003,7 @@ steps:
|
||||
- vllm/model_executor/layers
|
||||
- vllm/sampling_metadata.py
|
||||
- vllm/v1/sample/
|
||||
- vllm/entrypoints/generate/beam_search/
|
||||
- vllm/beam_search.py
|
||||
- tests/samplers
|
||||
- tests/conftest.py
|
||||
- vllm/_aiter_ops.py
|
||||
@@ -2182,72 +2177,10 @@ steps:
|
||||
commands:
|
||||
- pytest -v -s v1/e2e/spec_decode -k "speculators or mtp_correctness"
|
||||
|
||||
- label: Speculators Correctness # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
|
||||
agent_pool: mi300_1
|
||||
optional: true
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/config/speculative.py
|
||||
- vllm/engine/arg_utils.py
|
||||
- vllm/transformers_utils/config.py
|
||||
- vllm/transformers_utils/configs/speculators/
|
||||
- vllm/v1/spec_decode/
|
||||
- vllm/v1/worker/gpu/spec_decode/
|
||||
- vllm/v1/worker/gpu_model_runner.py
|
||||
- vllm/v1/sample/
|
||||
- vllm/v1/attention/backends/
|
||||
- vllm/v1/attention/selector.py
|
||||
- vllm/model_executor/model_loader/
|
||||
- vllm/model_executor/layers/
|
||||
- vllm/model_executor/models/llama_eagle3.py
|
||||
- vllm/model_executor/models/qwen3.py
|
||||
- vllm/model_executor/models/qwen3_dflash.py
|
||||
- vllm/model_executor/models/registry.py
|
||||
- vllm/_aiter_ops.py
|
||||
- tests/evals/gsm8k/
|
||||
- tests/v1/spec_decode/test_speculators_correctness.py
|
||||
- vllm/platforms/rocm.py
|
||||
commands:
|
||||
- export VLLM_ALLOW_INSECURE_SERIALIZATION=1
|
||||
- pytest -v -s v1/spec_decode/test_speculators_correctness.py -m slow_test
|
||||
|
||||
- label: Extract Hidden States Integration # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
|
||||
agent_pool: mi300_1
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/config/speculative.py
|
||||
- vllm/distributed/kv_transfer/kv_connector/
|
||||
- vllm/model_executor/layers/attention/
|
||||
- vllm/model_executor/layers/mamba/
|
||||
- vllm/model_executor/model_loader/
|
||||
- vllm/model_executor/models/extract_hidden_states.py
|
||||
- vllm/model_executor/models/llama.py
|
||||
- vllm/model_executor/models/qwen3_5.py
|
||||
- vllm/model_executor/models/qwen3_next.py
|
||||
- vllm/model_executor/models/registry.py
|
||||
- vllm/transformers_utils/configs/extract_hidden_states.py
|
||||
- vllm/transformers_utils/configs/qwen3_5.py
|
||||
- vllm/v1/attention/backends/
|
||||
- vllm/v1/attention/selector.py
|
||||
- vllm/v1/kv_cache_interface.py
|
||||
- vllm/v1/spec_decode/extract_hidden_states.py
|
||||
- vllm/v1/worker/gpu_model_runner.py
|
||||
- vllm/_aiter_ops.py
|
||||
- tests/v1/kv_connector/extract_hidden_states_integration/
|
||||
- vllm/platforms/rocm.py
|
||||
commands:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s v1/kv_connector/extract_hidden_states_integration
|
||||
|
||||
- label: V1 attention (H100-MI300) # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
|
||||
agent_pool: mi300_1
|
||||
optional: true
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/config/attention.py
|
||||
@@ -2406,7 +2339,7 @@ steps:
|
||||
- vllm/platforms/rocm.py
|
||||
commands:
|
||||
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
|
||||
- CROSS_LAYERS_BLOCKS=True ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
|
||||
- CROSS_LAYERS_BLOCKS=True ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
|
||||
|
||||
- label: Distributed DP Tests (4 GPUs) # TBD
|
||||
timeout_in_minutes: 180
|
||||
@@ -2442,7 +2375,7 @@ steps:
|
||||
- vllm/platforms/rocm.py
|
||||
commands:
|
||||
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
|
||||
- ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
|
||||
- ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
|
||||
|
||||
- label: DP EP Distributed NixlConnector PD accuracy tests (4 GPUs) # TBD
|
||||
timeout_in_minutes: 180
|
||||
@@ -2456,9 +2389,9 @@ steps:
|
||||
- vllm/platforms/rocm.py
|
||||
commands:
|
||||
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
|
||||
- DP_EP=1 ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
|
||||
- DP_EP=1 ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
|
||||
|
||||
- label: Hybrid SSM NixlConnector PD accuracy tests (4 GPUs) # TBD
|
||||
- label: Hyrbid SSM NixlConnector PD accuracy tests (4 GPUs) # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
|
||||
agent_pool: mi300_4
|
||||
@@ -2470,7 +2403,7 @@ steps:
|
||||
- vllm/platforms/rocm.py
|
||||
commands:
|
||||
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
|
||||
- HYBRID_SSM=1 ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
|
||||
- HYBRID_SSM=1 ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
|
||||
|
||||
- label: V1 e2e (4 GPUs) # TBD
|
||||
timeout_in_minutes: 180
|
||||
@@ -2606,7 +2539,6 @@ steps:
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
|
||||
agent_pool: mi325_1
|
||||
optional: true
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
@@ -2616,7 +2548,7 @@ steps:
|
||||
- tests/test_logger
|
||||
- tests/test_vllm_port
|
||||
commands:
|
||||
- pytest -v -s engine test_sequence.py test_config.py test_logger.py test_vllm_port.py test_jit_monitor.py
|
||||
- pytest -v -s engine test_sequence.py test_config.py test_logger.py test_vllm_port.py
|
||||
|
||||
#----------------------------------------------------------- mi325 · evals -----------------------------------------------------------#
|
||||
|
||||
@@ -2661,7 +2593,7 @@ steps:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-large-rocm.txt --tp-size=8
|
||||
|
||||
#----------------------------------------------------- mi325 · models / language -----------------------------------------------------#
|
||||
#---------------------------------------------------------- mi325 · models -----------------------------------------------------------#
|
||||
|
||||
- label: Language Models Test (Extended Generation) # TBD
|
||||
timeout_in_minutes: 180
|
||||
@@ -2692,13 +2624,10 @@ steps:
|
||||
- uv pip install --system --no-build-isolation 'git+https://github.com/Dao-AILab/causal-conv1d@v1.6.0'
|
||||
- pytest -v -s models/language/generation -m hybrid_model --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
|
||||
|
||||
#---------------------------------------------------- mi325 · models / multimodal ----------------------------------------------------#
|
||||
|
||||
- label: Multi-Modal Models (Extended Pooling) # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
|
||||
agent_pool: mi325_1
|
||||
optional: true
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
@@ -2766,39 +2695,23 @@ steps:
|
||||
optional: true
|
||||
working_dir: "/vllm-workspace/"
|
||||
source_file_dependencies:
|
||||
- csrc/custom_quickreduce.cu
|
||||
- csrc/ops.h
|
||||
- csrc/torch_bindings.cpp
|
||||
- vllm/distributed/
|
||||
- vllm/model_executor/layers/
|
||||
- vllm/entrypoints/llm.py
|
||||
- vllm/config/parallel.py
|
||||
- vllm/model_executor/layers/fused_moe/
|
||||
- vllm/v1/engine/
|
||||
- vllm/v1/executor/
|
||||
- vllm/v1/worker/
|
||||
- vllm/v1/distributed/
|
||||
- vllm/model_executor/layers/fused_moe/
|
||||
- vllm/v1/attention/backends/
|
||||
- vllm/v1/attention/selector.py
|
||||
- vllm/_aiter_ops.py
|
||||
- vllm/_custom_ops.py
|
||||
- vllm/platforms/rocm.py
|
||||
- vllm/envs.py
|
||||
- examples/offline_inference/data_parallel.py
|
||||
- tests/distributed/test_context_parallel.py
|
||||
- tests/distributed/test_rocm_quick_reduce.py
|
||||
- tests/distributed/test_quick_all_reduce.py
|
||||
- tests/v1/distributed/test_dbo.py
|
||||
- tests/utils.py
|
||||
- examples/features/data_parallel/data_parallel_offline.py
|
||||
- vllm/_aiter_ops.py
|
||||
- vllm/platforms/rocm.py
|
||||
commands:
|
||||
- pytest -v -s tests/distributed/test_context_parallel.py
|
||||
- pytest -v -s tests/v1/distributed/test_dbo.py
|
||||
- pytest -v -s tests/distributed/test_rocm_quick_reduce.py
|
||||
- pytest -v -s tests/distributed/test_quick_all_reduce.py
|
||||
|
||||
#-------------------------------------------------------- mi355 · entrypoints --------------------------------------------------------#
|
||||
|
||||
- label: Entrypoints Integration (API Server) # TBD
|
||||
- label: Entrypoints Integration (API Server 2) # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
|
||||
agent_pool: mi355_1
|
||||
@@ -2808,65 +2721,62 @@ steps:
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/entrypoints/serve
|
||||
commands:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s entrypoints/serve --ignore=entrypoints/serve/dev/rpc
|
||||
- PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/serve/dev/rpc
|
||||
|
||||
- label: Entrypoints Integration (API Server OpenAI - Part 1) # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
|
||||
agent_pool: mi355_1
|
||||
fast_check: true
|
||||
torch_nightly: true
|
||||
optional: true
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/entrypoints/openai
|
||||
- tests/entrypoints/test_chat_utils
|
||||
commands:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s entrypoints/openai --ignore=entrypoints/openai/completion --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/correctness
|
||||
|
||||
- label: Entrypoints Integration (API Server OpenAI - Part 2) # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
|
||||
agent_pool: mi355_1
|
||||
fast_check: true
|
||||
torch_nightly: true
|
||||
optional: true
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/entrypoints/openai
|
||||
- tests/entrypoints/test_chat_utils
|
||||
commands:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s entrypoints/openai/chat_completion
|
||||
- pytest -v -s entrypoints/openai/completion --ignore=entrypoints/openai/completion/test_tensorizer_entrypoint.py
|
||||
|
||||
- label: Entrypoints Integration (API Server Generate) # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
|
||||
agent_pool: mi355_1
|
||||
fast_check: true
|
||||
torch_nightly: true
|
||||
optional: true
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/entrypoints/rpc
|
||||
- tests/entrypoints/serve/instrumentator
|
||||
- tests/tool_use
|
||||
- tests/entrypoints/tool_parsers
|
||||
- tests/entrypoints/anthropic
|
||||
- tests/entrypoints/generate
|
||||
commands:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s entrypoints/serve/instrumentator
|
||||
- PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/rpc
|
||||
- pytest -v -s tool_use
|
||||
- pytest -v -s entrypoints/tool_parsers
|
||||
- pytest -v -s entrypoints/generate
|
||||
- pytest -v -s entrypoints/anthropic
|
||||
|
||||
- label: Entrypoints Integration (API Server openai - Part 1) # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
|
||||
agent_pool: mi355_1
|
||||
fast_check: true
|
||||
torch_nightly: true
|
||||
optional: true
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/entrypoints/openai
|
||||
- tests/entrypoints/test_chat_utils
|
||||
commands:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s entrypoints/openai/chat_completion --ignore=entrypoints/openai/chat_completion/test_chat_with_tool_reasoning.py --ignore=entrypoints/openai/chat_completion/test_oot_registration.py
|
||||
|
||||
- label: Entrypoints Integration (API Server openai - Part 2) # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
|
||||
agent_pool: mi355_1
|
||||
fast_check: true
|
||||
torch_nightly: true
|
||||
optional: true
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/entrypoints/openai
|
||||
- tests/entrypoints/test_chat_utils
|
||||
commands:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s entrypoints/openai/completion --ignore=entrypoints/openai/completion/test_tensorizer_entrypoint.py
|
||||
- pytest -v -s entrypoints/test_chat_utils.py
|
||||
|
||||
- label: Entrypoints Integration (API Server openai - Part 3) # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
|
||||
agent_pool: mi355_1
|
||||
fast_check: true
|
||||
torch_nightly: true
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/entrypoints/openai
|
||||
- tests/entrypoints/test_chat_utils
|
||||
commands:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s entrypoints/openai --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/completion --ignore=entrypoints/openai/correctness/ --ignore=entrypoints/openai/tool_parsers/ --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/test_multi_api_servers.py
|
||||
|
||||
- label: Entrypoints Integration (Speech to Text) # TBD
|
||||
timeout_in_minutes: 180
|
||||
@@ -2882,20 +2792,6 @@ steps:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s entrypoints/speech_to_text
|
||||
|
||||
- label: Entrypoints Integration (Multimodal)
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi355]
|
||||
agent_pool: mi355_1
|
||||
fast_check: true
|
||||
torch_nightly: true
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/entrypoints/multimodal
|
||||
commands:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s entrypoints/multimodal
|
||||
|
||||
- label: Entrypoints Integration (Pooling) # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
|
||||
@@ -2946,6 +2842,7 @@ steps:
|
||||
- vllm/model_executor/models/qwen3_5_mtp.py
|
||||
- vllm/transformers_utils/configs/qwen3_5.py
|
||||
- vllm/transformers_utils/configs/qwen3_5_moe.py
|
||||
- vllm/model_executor/models/qwen.py
|
||||
- vllm/model_executor/models/qwen2.py
|
||||
- vllm/model_executor/models/qwen3.py
|
||||
- vllm/model_executor/models/qwen3_next.py
|
||||
@@ -2974,7 +2871,7 @@ steps:
|
||||
commands:
|
||||
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-mi3xx-fp8-and-mixed.txt
|
||||
|
||||
- label: Qwen3-30B-A3B-FP8-block Sync EPLB Accuracy (B200-MI355) # TBD
|
||||
- label: Qwen3-30B-A3B-FP8-block Accuracy (B200-MI355) # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
|
||||
agent_pool: mi355_2
|
||||
@@ -3072,7 +2969,7 @@ steps:
|
||||
- vllm/_aiter_ops.py
|
||||
commands:
|
||||
- rocm-smi
|
||||
- python3 examples/basic/offline_inference/chat.py --attention-backend TRITON_ATTN
|
||||
- python3 examples/basic/offline_inference/chat.py
|
||||
- pytest -v -s tests/kernels/attention/test_attention_selector.py
|
||||
|
||||
- label: Kernels Attention Test %N # TBD
|
||||
@@ -3146,7 +3043,7 @@ steps:
|
||||
commands:
|
||||
- pytest -v -s kernels/moe/test_deepep_moe.py
|
||||
|
||||
#----------------------------------------------------- mi355 · models / language -----------------------------------------------------#
|
||||
#---------------------------------------------------------- mi355 · models -----------------------------------------------------------#
|
||||
|
||||
- label: Language Models Test (Extended Generation) # TBD
|
||||
timeout_in_minutes: 180
|
||||
@@ -3183,6 +3080,7 @@ steps:
|
||||
- vllm/model_executor/models/qwen3_5_mtp.py
|
||||
- vllm/transformers_utils/configs/qwen3_5.py
|
||||
- vllm/transformers_utils/configs/qwen3_5_moe.py
|
||||
- vllm/model_executor/models/qwen.py
|
||||
- vllm/model_executor/models/qwen2.py
|
||||
- vllm/model_executor/models/qwen3.py
|
||||
- vllm/model_executor/models/qwen3_next.py
|
||||
@@ -3213,8 +3111,6 @@ steps:
|
||||
- pip freeze | grep -E 'torch'
|
||||
- pytest -v -s models/language -m 'core_model and (not slow_test)'
|
||||
|
||||
#---------------------------------------------------- mi355 · models / multimodal ----------------------------------------------------#
|
||||
|
||||
- label: Multi-Modal Models (Extended Generation 1) # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
|
||||
@@ -3286,8 +3182,6 @@ steps:
|
||||
- pytest -v -s models/multimodal/generation/test_memory_leak.py -m core_model
|
||||
- cd .. && VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s tests/models/multimodal/generation/test_whisper.py -m core_model
|
||||
|
||||
#----------------------------------------------------- mi355 · models / quantized -----------------------------------------------------#
|
||||
|
||||
- label: Quantized Models Test # TBD
|
||||
timeout_in_minutes: 180
|
||||
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
|
||||
@@ -3435,7 +3329,7 @@ steps:
|
||||
- vllm/platforms/rocm.py
|
||||
commands:
|
||||
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
|
||||
- ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/spec_decode_acceptance_test.sh
|
||||
- ATTENTION_BACKEND=ROCM_ATTN bash v1/kv_connector/nixl_integration/spec_decode_acceptance_test.sh
|
||||
|
||||
- label: Distributed NixlConnector PD accuracy (4 GPUs) # TBD
|
||||
timeout_in_minutes: 180
|
||||
@@ -3450,7 +3344,7 @@ steps:
|
||||
- vllm/platforms/rocm.py
|
||||
commands:
|
||||
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
|
||||
- ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
|
||||
- ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
|
||||
|
||||
- label: DP EP Distributed NixlConnector PD accuracy tests (4 GPUs) # TBD
|
||||
timeout_in_minutes: 180
|
||||
@@ -3465,7 +3359,7 @@ steps:
|
||||
- vllm/platforms/rocm.py
|
||||
commands:
|
||||
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
|
||||
- DP_EP=1 ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
|
||||
- DP_EP=1 ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
|
||||
|
||||
#------------------------------------------------------ mi355 · weight_loading -------------------------------------------------------#
|
||||
|
||||
|
||||
@@ -2,8 +2,8 @@ group: Attention
|
||||
depends_on:
|
||||
- image-build
|
||||
steps:
|
||||
- label: V1 attention (H100-MI300)
|
||||
key: v1-attention-h100-mi300
|
||||
- label: V1 attention (H100)
|
||||
key: v1-attention-h100
|
||||
timeout_in_minutes: 30
|
||||
device: h100
|
||||
source_file_dependencies:
|
||||
@@ -13,20 +13,6 @@ steps:
|
||||
- tests/v1/attention
|
||||
commands:
|
||||
- pytest -v -s v1/attention
|
||||
mirror:
|
||||
amd:
|
||||
device: mi325_1
|
||||
timeout_in_minutes: 70
|
||||
depends_on:
|
||||
- image-build-amd
|
||||
source_file_dependencies:
|
||||
- vllm/config/attention.py
|
||||
- vllm/model_executor/layers/attention
|
||||
- vllm/v1/attention
|
||||
- tests/v1/attention
|
||||
- vllm/_aiter_ops.py
|
||||
- vllm/envs.py
|
||||
- vllm/platforms/rocm.py
|
||||
|
||||
- label: V1 attention (B200)
|
||||
key: v1-attention-b200
|
||||
|
||||
@@ -10,9 +10,9 @@ steps:
|
||||
- vllm/
|
||||
- tests/basic_correctness/test_basic_correctness
|
||||
- tests/basic_correctness/test_cpu_offload
|
||||
- tests/basic_correctness/test_mem.py
|
||||
- tests/basic_correctness/test_cumem.py
|
||||
commands:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s basic_correctness/test_mem.py
|
||||
- pytest -v -s basic_correctness/test_cumem.py
|
||||
- pytest -v -s basic_correctness/test_basic_correctness.py
|
||||
- pytest -v -s basic_correctness/test_cpu_offload.py
|
||||
|
||||
@@ -11,7 +11,7 @@ steps:
|
||||
- vllm/distributed/kv_transfer/kv_connector/v1/nixl/
|
||||
- tests/v1/kv_connector/nixl_integration/
|
||||
commands:
|
||||
- bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
|
||||
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
|
||||
- bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
|
||||
- label: Distributed FlashInfer NixlConnector PD accuracy (4 GPUs)
|
||||
key: distributed-flashinfer-nixlconnector-pd-accuracy-4-gpus
|
||||
@@ -22,7 +22,7 @@ steps:
|
||||
- vllm/distributed/kv_transfer/kv_connector/v1/nixl/
|
||||
- tests/v1/kv_connector/nixl_integration/
|
||||
commands:
|
||||
- bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
|
||||
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
|
||||
- FLASHINFER=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
|
||||
|
||||
- label: DP EP Distributed NixlConnector PD accuracy tests (4 GPUs)
|
||||
@@ -34,7 +34,7 @@ steps:
|
||||
- vllm/distributed/kv_transfer/kv_connector/v1/nixl/
|
||||
- tests/v1/kv_connector/nixl_integration/
|
||||
commands:
|
||||
- bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
|
||||
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
|
||||
- DP_EP=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
|
||||
|
||||
- label: CrossLayer KV layout Distributed NixlConnector PD accuracy tests (4 GPUs)
|
||||
@@ -46,7 +46,7 @@ steps:
|
||||
- vllm/distributed/kv_transfer/kv_connector/v1/nixl/
|
||||
- tests/v1/kv_connector/nixl_integration/
|
||||
commands:
|
||||
- bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
|
||||
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
|
||||
- CROSS_LAYERS_BLOCKS=True bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
|
||||
|
||||
- label: Hybrid SSM NixlConnector PD accuracy tests (4 GPUs)
|
||||
@@ -58,23 +58,9 @@ steps:
|
||||
- vllm/distributed/kv_transfer/kv_connector/v1/nixl/
|
||||
- tests/v1/kv_connector/nixl_integration/
|
||||
commands:
|
||||
- bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
|
||||
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
|
||||
- HYBRID_SSM=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
|
||||
|
||||
- label: Hybrid SSM NixlConnector PD prefix cache test (2 GPUs)
|
||||
key: hybrid-ssm-nixlconnector-pd-prefix-cache-2-gpus
|
||||
timeout_in_minutes: 25
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
num_devices: 2
|
||||
source_file_dependencies:
|
||||
- vllm/distributed/kv_transfer/kv_connector/v1/nixl/
|
||||
- vllm/v1/core/sched/
|
||||
- vllm/v1/core/kv_cache_coordinator.py
|
||||
- tests/v1/kv_connector/nixl_integration/
|
||||
commands:
|
||||
- bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
|
||||
- bash v1/kv_connector/nixl_integration/run_mamba_prefix_cache_test.sh
|
||||
|
||||
- label: MultiConnector (Nixl+Offloading) PD accuracy (2 GPUs)
|
||||
key: multiconnector-nixl-offloading-pd-accuracy-2-gpus
|
||||
timeout_in_minutes: 30
|
||||
@@ -87,7 +73,7 @@ steps:
|
||||
- vllm/distributed/kv_transfer/kv_connector/v1/offloading/
|
||||
- tests/v1/kv_connector/nixl_integration/
|
||||
commands:
|
||||
- bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
|
||||
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
|
||||
- bash v1/kv_connector/nixl_integration/run_multi_connector_accuracy_test.sh
|
||||
|
||||
- label: NixlConnector PD + Spec Decode acceptance (2 GPUs)
|
||||
@@ -101,8 +87,8 @@ steps:
|
||||
- vllm/v1/worker/kv_connector_model_runner_mixin.py
|
||||
- tests/v1/kv_connector/nixl_integration/
|
||||
commands:
|
||||
- bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
|
||||
- bash v1/kv_connector/nixl_integration/config_sweep_spec_decode_test.sh
|
||||
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
|
||||
- bash v1/kv_connector/nixl_integration/spec_decode_acceptance_test.sh
|
||||
|
||||
- label: MultiConnector (Nixl+Offloading) PD edge cases (2 GPUs)
|
||||
key: multiconnector-nixl-offloading-pd-edge-cases-2-gpus
|
||||
@@ -116,5 +102,5 @@ steps:
|
||||
- vllm/distributed/kv_transfer/kv_connector/v1/offloading/
|
||||
- tests/v1/kv_connector/nixl_integration/
|
||||
commands:
|
||||
- bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
|
||||
- bash v1/kv_connector/nixl_integration/run_multi_connector_edge_case_test.sh
|
||||
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
|
||||
- bash v1/kv_connector/nixl_integration/run_multi_connector_edge_case_test.sh
|
||||
@@ -2,8 +2,8 @@ group: E2E Integration
|
||||
depends_on:
|
||||
- image-build
|
||||
steps:
|
||||
- label: DeepSeek V2-Lite Sync EPLB Accuracy
|
||||
key: deepseek-v2-lite-sync-eplb-accuracy
|
||||
- label: DeepSeek V2-Lite Accuracy
|
||||
key: deepseek-v2-lite-accuracy
|
||||
timeout_in_minutes: 60
|
||||
device: h100
|
||||
optional: true
|
||||
@@ -12,8 +12,8 @@ steps:
|
||||
commands:
|
||||
- bash .buildkite/scripts/scheduled_integration_test/deepseek_v2_lite_ep_eplb.sh 0.25 200 8010
|
||||
|
||||
- label: Qwen3-30B-A3B-FP8-block Sync EPLB Accuracy
|
||||
key: qwen3-30b-a3b-fp8-block-sync-eplb-accuracy
|
||||
- label: Qwen3-30B-A3B-FP8-block Accuracy
|
||||
key: qwen3-30b-a3b-fp8-block-accuracy
|
||||
timeout_in_minutes: 60
|
||||
device: h100
|
||||
optional: true
|
||||
@@ -22,8 +22,8 @@ steps:
|
||||
commands:
|
||||
- bash .buildkite/scripts/scheduled_integration_test/qwen30b_a3b_fp8_block_ep_eplb.sh 0.8 200 8020
|
||||
|
||||
- label: Qwen3-30B-A3B-FP8-block Sync EPLB Accuracy (B200)
|
||||
key: qwen3-30b-a3b-fp8-block-sync-eplb-accuracy-b200
|
||||
- label: Qwen3-30B-A3B-FP8-block Accuracy (B200)
|
||||
key: qwen3-30b-a3b-fp8-block-accuracy-b200
|
||||
timeout_in_minutes: 60
|
||||
device: b200-k8s
|
||||
optional: true
|
||||
|
||||
@@ -26,12 +26,6 @@ steps:
|
||||
- tests/test_jit_monitor.py
|
||||
commands:
|
||||
- pytest -v -s engine test_sequence.py test_config.py test_logger.py test_vllm_port.py test_jit_monitor.py
|
||||
mirror:
|
||||
amd:
|
||||
device: mi325_1
|
||||
timeout_in_minutes: 60
|
||||
depends_on:
|
||||
- image-build-amd
|
||||
|
||||
- label: Engine (1 GPU)
|
||||
key: engine-1-gpu
|
||||
@@ -44,7 +38,7 @@ steps:
|
||||
- pytest -v -s v1/engine --ignore v1/engine/test_preprocess_error_handling.py
|
||||
mirror:
|
||||
amd:
|
||||
device: mi325_1
|
||||
device: mi300_1
|
||||
timeout_in_minutes: 40
|
||||
depends_on:
|
||||
- image-build-amd
|
||||
@@ -66,7 +60,6 @@ steps:
|
||||
- image-build-amd
|
||||
|
||||
- label: e2e Core (1 GPU)
|
||||
device: h200_35gb
|
||||
key: e2e-core-1-gpu
|
||||
timeout_in_minutes: 30
|
||||
source_file_dependencies:
|
||||
|
||||
@@ -8,11 +8,10 @@ steps:
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/entrypoints
|
||||
- tests/entrypoints/unit_tests
|
||||
- tests/entrypoints/weight_transfer
|
||||
- tests/entrypoints/
|
||||
commands:
|
||||
- pytest -v -s entrypoints/unit_tests
|
||||
- pytest -v -s entrypoints/weight_transfer
|
||||
- pytest -v -s entrypoints/openai/tool_parsers
|
||||
- pytest -v -s entrypoints/ --ignore=entrypoints/llm --ignore=entrypoints/rpc --ignore=entrypoints/sleep --ignore=entrypoints/serve/instrumentator --ignore=entrypoints/openai --ignore=entrypoints/offline_mode --ignore=entrypoints/test_chat_utils.py --ignore=entrypoints/pooling --ignore=entrypoints/speech_to_text
|
||||
|
||||
- label: Entrypoints Integration (LLM)
|
||||
key: entrypoints-integration-llm
|
||||
@@ -21,36 +20,19 @@ steps:
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/entrypoints/llm
|
||||
- tests/entrypoints/offline_mode
|
||||
commands:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s entrypoints/llm --ignore=entrypoints/llm/test_generate.py --ignore=entrypoints/llm/test_collective_rpc.py --ignore=entrypoints/llm/offline_mode
|
||||
- pytest -v -s entrypoints/llm --ignore=entrypoints/llm/test_generate.py --ignore=entrypoints/llm/test_collective_rpc.py
|
||||
- pytest -v -s entrypoints/llm/test_generate.py # it needs a clean process
|
||||
- pytest -v -s entrypoints/llm/offline_mode # Needs to avoid interference with other tests
|
||||
- pytest -v -s entrypoints/offline_mode # Needs to avoid interference with other tests
|
||||
mirror:
|
||||
amd:
|
||||
device: mi325_1
|
||||
device: mi300_1
|
||||
depends_on:
|
||||
- image-build-amd
|
||||
|
||||
- label: Entrypoints Integration (API Server)
|
||||
key: entrypoints-integration-api-server
|
||||
device: h200_35gb
|
||||
timeout_in_minutes: 130
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/entrypoints/serve
|
||||
commands:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s entrypoints/serve --ignore=entrypoints/serve/dev/rpc
|
||||
- PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/serve/dev/rpc
|
||||
mirror:
|
||||
amd:
|
||||
device: mi325_1
|
||||
depends_on:
|
||||
- image-build-amd
|
||||
|
||||
- label: Entrypoints Integration (API Server OpenAI - Part 1)
|
||||
- label: Entrypoints Integration (API Server openai - Part 1)
|
||||
key: entrypoints-integration-api-server-openai-part-1
|
||||
timeout_in_minutes: 50
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
@@ -60,15 +42,15 @@ steps:
|
||||
- tests/entrypoints/test_chat_utils
|
||||
commands:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s entrypoints/openai --ignore=entrypoints/openai/completion --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/correctness
|
||||
- pytest -v -s entrypoints/openai/chat_completion --ignore=entrypoints/openai/chat_completion/test_chat_with_tool_reasoning.py --ignore=entrypoints/openai/chat_completion/test_oot_registration.py
|
||||
mirror:
|
||||
amd:
|
||||
device: mi325_1
|
||||
device: mi300_1
|
||||
timeout_in_minutes: 80
|
||||
depends_on:
|
||||
- image-build-amd
|
||||
|
||||
- label: Entrypoints Integration (API Server OpenAI - Part 2)
|
||||
- label: Entrypoints Integration (API Server openai - Part 2)
|
||||
key: entrypoints-integration-api-server-openai-part-2
|
||||
timeout_in_minutes: 50
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
@@ -77,50 +59,55 @@ steps:
|
||||
- tests/entrypoints/openai
|
||||
- tests/entrypoints/test_chat_utils
|
||||
commands:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s entrypoints/openai/chat_completion
|
||||
- pytest -v -s entrypoints/openai/completion --ignore=entrypoints/openai/completion/test_tensorizer_entrypoint.py
|
||||
- pytest -v -s entrypoints/test_chat_utils.py
|
||||
mirror:
|
||||
amd:
|
||||
device: mi325_1
|
||||
timeout_in_minutes: 80
|
||||
depends_on:
|
||||
- image-build-amd
|
||||
|
||||
- label: Entrypoints Integration (API Server Generate)
|
||||
key: entrypoints-integration-api-server-generate
|
||||
timeout_in_minutes: 50
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/tool_use
|
||||
- tests/entrypoints/tool_parsers
|
||||
- tests/entrypoints/anthropic
|
||||
- tests/entrypoints/generate
|
||||
commands:
|
||||
- pytest -v -s tool_use
|
||||
- pytest -v -s entrypoints/tool_parsers
|
||||
- pytest -v -s entrypoints/generate
|
||||
- pytest -v -s entrypoints/anthropic
|
||||
mirror:
|
||||
amd:
|
||||
device: mi325_1
|
||||
device: mi300_1
|
||||
timeout_in_minutes: 60
|
||||
depends_on:
|
||||
- image-build-amd
|
||||
|
||||
- label: Entrypoints Integration (Responses API)
|
||||
key: entrypoints-integration-responses-api
|
||||
- label: Entrypoints Integration (API Server openai - Part 3)
|
||||
key: entrypoints-integration-api-server-openai-part-3
|
||||
timeout_in_minutes: 50
|
||||
device: h200_18gb
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/entrypoints/openai/responses
|
||||
- tests/entrypoints/openai
|
||||
- tests/entrypoints/test_chat_utils
|
||||
commands:
|
||||
- pytest -v -s entrypoints/openai/responses
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s entrypoints/openai --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/completion --ignore=entrypoints/openai/correctness/ --ignore=entrypoints/openai/tool_parsers/ --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/test_multi_api_servers.py
|
||||
mirror:
|
||||
amd:
|
||||
device: mi300_1
|
||||
timeout_in_minutes: 60
|
||||
depends_on:
|
||||
- image-build-amd
|
||||
|
||||
- label: Entrypoints Integration (API Server 2)
|
||||
key: entrypoints-integration-api-server-2
|
||||
timeout_in_minutes: 130
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/entrypoints/rpc
|
||||
- tests/entrypoints/serve/instrumentator
|
||||
- tests/tool_use
|
||||
commands:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s entrypoints/serve/instrumentator
|
||||
- PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/rpc
|
||||
- pytest -v -s tool_use
|
||||
mirror:
|
||||
amd:
|
||||
device: mi300_1
|
||||
depends_on:
|
||||
- image-build-amd
|
||||
|
||||
- label: Entrypoints Integration (Speech to Text)
|
||||
device: h200_35gb
|
||||
key: entrypoints-integration-speech_to_text
|
||||
timeout_in_minutes: 50
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
@@ -131,18 +118,6 @@ steps:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s entrypoints/speech_to_text
|
||||
|
||||
- label: Entrypoints Integration (Multimodal)
|
||||
device: h200_35gb
|
||||
key: entrypoints-integration-multimodal
|
||||
timeout_in_minutes: 50
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/entrypoints/multimodal
|
||||
commands:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s entrypoints/multimodal
|
||||
|
||||
- label: Entrypoints Integration (Pooling)
|
||||
key: entrypoints-integration-pooling
|
||||
timeout_in_minutes: 50
|
||||
@@ -154,6 +129,16 @@ steps:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s entrypoints/pooling
|
||||
|
||||
- label: Entrypoints Integration (Responses API)
|
||||
key: entrypoints-integration-responses-api
|
||||
timeout_in_minutes: 50
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/entrypoints/openai/responses
|
||||
commands:
|
||||
- pytest -v -s entrypoints/openai/responses
|
||||
|
||||
- label: OpenAI API Correctness
|
||||
key: openai-api-correctness
|
||||
timeout_in_minutes: 30
|
||||
@@ -161,22 +146,6 @@ steps:
|
||||
source_file_dependencies:
|
||||
- csrc/
|
||||
- vllm/entrypoints/openai/
|
||||
- vllm/model_executor/models/whisper.py
|
||||
commands: # LMEval
|
||||
- pytest -s entrypoints/openai/correctness/
|
||||
mirror:
|
||||
amd:
|
||||
device: mi325_1
|
||||
depends_on:
|
||||
- image-build-amd
|
||||
source_file_dependencies:
|
||||
- csrc/
|
||||
- vllm/entrypoints/openai/
|
||||
- vllm/model_executor/layers/
|
||||
- vllm/v1/attention/backends/
|
||||
- vllm/v1/attention/selector.py
|
||||
- vllm/_aiter_ops.py
|
||||
- vllm/platforms/rocm.py
|
||||
- vllm/model_executor/model_loader/
|
||||
commands:
|
||||
- bash ../tools/install_torchcodec_rocm.sh || exit 1
|
||||
- pytest -s entrypoints/openai/correctness/
|
||||
|
||||
@@ -21,9 +21,8 @@ steps:
|
||||
- csrc/
|
||||
- tests/kernels/core
|
||||
- tests/kernels/test_concat_mla_q.py
|
||||
- tests/kernels/test_fused_qk_norm_rope_gate.py
|
||||
commands:
|
||||
- pytest -v -s kernels/core --ignore=kernels/core/test_minimax_reduce_rms.py kernels/test_concat_mla_q.py kernels/test_fused_qk_norm_rope_gate.py
|
||||
- pytest -v -s kernels/core --ignore=kernels/core/test_minimax_reduce_rms.py kernels/test_concat_mla_q.py
|
||||
|
||||
- label: Kernels MiniMax Reduce RMS Test (2 GPUs)
|
||||
key: kernels-minimax-reduce-rms-test-2-gpus
|
||||
@@ -39,28 +38,6 @@ steps:
|
||||
commands:
|
||||
- pytest -v -s kernels/core/test_minimax_reduce_rms.py
|
||||
|
||||
- label: Deepseek V4 Kernel Test (H100)
|
||||
key: deepseek-v4-kernel-test-h100
|
||||
timeout_in_minutes: 15
|
||||
device: h100
|
||||
source_file_dependencies:
|
||||
- csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu
|
||||
- vllm/models/deepseek_v4/common/ops/
|
||||
- tests/kernels/test_fused_deepseek_v4_qnorm_rope_kv_insert.py
|
||||
commands:
|
||||
- pytest -v -s kernels/test_fused_deepseek_v4_*.py
|
||||
|
||||
- label: Deepseek V4 Kernel Test (B200)
|
||||
key: deepseek-v4-kernel-test-b200
|
||||
timeout_in_minutes: 15
|
||||
device: b200-k8s
|
||||
source_file_dependencies:
|
||||
- csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu
|
||||
- vllm/models/deepseek_v4/common/ops/
|
||||
- tests/kernels/test_fused_deepseek_v4_qnorm_rope_kv_insert.py
|
||||
commands:
|
||||
- pytest -v -s kernels/test_fused_deepseek_v4_*.py
|
||||
|
||||
- label: Kernels Attention Test %N
|
||||
key: kernels-attention-test
|
||||
timeout_in_minutes: 35
|
||||
@@ -75,19 +52,6 @@ steps:
|
||||
- pytest -v -s kernels/attention --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
|
||||
parallelism: 2
|
||||
|
||||
- label: Kernels Attention DiffKV Test (H100)
|
||||
key: kernels-attention-diffkv-test-h100
|
||||
timeout_in_minutes: 20
|
||||
device: h100
|
||||
num_devices: 1
|
||||
source_file_dependencies:
|
||||
- vllm/v1/attention/ops/triton_unified_attention_diffkv.py
|
||||
- vllm/v1/attention/backends/triton_attn_diffkv.py
|
||||
- vllm/v1/attention/backends/flash_attn_diffkv.py
|
||||
- tests/kernels/attention/test_triton_unified_attention_diffkv.py
|
||||
commands:
|
||||
- pytest -v -s kernels/attention/test_triton_unified_attention_diffkv.py
|
||||
|
||||
- label: Kernels Quantization Test %N
|
||||
key: kernels-quantization-test
|
||||
timeout_in_minutes: 90
|
||||
@@ -100,7 +64,7 @@ steps:
|
||||
parallelism: 2
|
||||
mirror:
|
||||
amd:
|
||||
device: mi325_1
|
||||
device: mi300_1
|
||||
source_file_dependencies:
|
||||
- csrc/quantization/
|
||||
- vllm/model_executor/layers/quantization
|
||||
@@ -186,12 +150,9 @@ steps:
|
||||
- csrc/quantization/fp4/
|
||||
- csrc/attention/mla/
|
||||
- csrc/quantization/cutlass_w8a8/moe/
|
||||
- vllm/model_executor/layers/fused_moe/experts/cutlass_moe.py
|
||||
- vllm/model_executor/layers/fused_moe/experts/flashinfer_cutlass_moe.py
|
||||
- vllm/model_executor/layers/fused_moe/experts/trtllm_nvfp4_moe.py
|
||||
- vllm/model_executor/layers/fused_moe/oracle/nvfp4.py
|
||||
- vllm/model_executor/layers/fused_moe/prepare_finalize/flashinfer_nvlink_one_sided.py
|
||||
- vllm/model_executor/layers/fused_moe/prepare_finalize/flashinfer_nvlink_two_sided.py
|
||||
- vllm/model_executor/layers/fused_moe/cutlass_moe.py
|
||||
- vllm/model_executor/layers/fused_moe/flashinfer_cutlass_moe.py
|
||||
- vllm/model_executor/layers/fused_moe/flashinfer_a2a_prepare_finalize.py
|
||||
- vllm/model_executor/layers/quantization/utils/flashinfer_utils.py
|
||||
- vllm/v1/attention/backends/flashinfer.py
|
||||
- vllm/v1/attention/backends/mla/cutlass_mla.py
|
||||
@@ -224,7 +185,6 @@ steps:
|
||||
- pytest -v -s tests/kernels/moe/test_ocp_mx_moe.py
|
||||
- pytest -v -s tests/kernels/moe/test_flashinfer.py
|
||||
- pytest -v -s tests/kernels/moe/test_flashinfer_moe.py
|
||||
- pytest -v -s tests/kernels/moe/test_trtllm_nvfp4_moe.py
|
||||
- pytest -v -s tests/kernels/moe/test_cutedsl_moe.py
|
||||
# e2e
|
||||
- pytest -v -s tests/models/quantization/test_nvfp4.py
|
||||
@@ -237,7 +197,7 @@ steps:
|
||||
- vllm/utils/import_utils.py
|
||||
- tests/kernels/helion/
|
||||
commands:
|
||||
- pip install helion==1.1.0
|
||||
- pip install helion==1.0.0
|
||||
- pytest -v -s kernels/helion/
|
||||
|
||||
|
||||
@@ -313,4 +273,3 @@ steps:
|
||||
- vllm/config
|
||||
commands:
|
||||
- pytest -v -s kernels/moe/test_moe_layer.py
|
||||
- pytest -v -s kernels/moe/test_deepep_v2_moe.py
|
||||
|
||||
@@ -3,7 +3,6 @@ depends_on:
|
||||
- image-build
|
||||
steps:
|
||||
- label: LM Eval Small Models
|
||||
device: h200_35gb
|
||||
key: lm-eval-small-models
|
||||
timeout_in_minutes: 75
|
||||
source_file_dependencies:
|
||||
@@ -12,21 +11,6 @@ steps:
|
||||
autorun_on_main: true
|
||||
commands:
|
||||
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-small.txt
|
||||
mirror:
|
||||
amd:
|
||||
device: mi325_1
|
||||
timeout_in_minutes: 55
|
||||
depends_on:
|
||||
- image-build-amd
|
||||
source_file_dependencies:
|
||||
- csrc/
|
||||
- vllm/model_executor/layers/quantization
|
||||
- vllm/model_executor/models/
|
||||
- vllm/model_executor/model_loader/
|
||||
- vllm/v1/attention/backends/
|
||||
- vllm/v1/attention/selector.py
|
||||
- vllm/_aiter_ops.py
|
||||
- vllm/platforms/rocm.py
|
||||
|
||||
# - label: LM Eval Large Models (4 GPUs)(A100)
|
||||
# device: a100
|
||||
@@ -168,7 +152,6 @@ steps:
|
||||
- pytest -s -v evals/gpt_oss/test_gpqa_correctness.py --config-list-file=configs/models-b200.txt
|
||||
|
||||
- label: MRCR Eval Small Models
|
||||
device: h200_35gb
|
||||
timeout_in_minutes: 30
|
||||
source_file_dependencies:
|
||||
- tests/evals/mrcr/
|
||||
|
||||
@@ -3,7 +3,6 @@ depends_on:
|
||||
- image-build
|
||||
steps:
|
||||
- label: LoRA %N
|
||||
device: h200_35gb
|
||||
key: lora
|
||||
timeout_in_minutes: 30
|
||||
source_file_dependencies:
|
||||
|
||||
@@ -3,7 +3,6 @@ depends_on:
|
||||
- image-build
|
||||
steps:
|
||||
- label: V1 Spec Decode
|
||||
device: h200_35gb
|
||||
key: v1-spec-decode
|
||||
timeout_in_minutes: 30
|
||||
source_file_dependencies:
|
||||
@@ -52,7 +51,7 @@ steps:
|
||||
- pytest -v -s v1/test_outputs.py
|
||||
mirror:
|
||||
amd:
|
||||
device: mi325_1
|
||||
device: mi300_1
|
||||
depends_on:
|
||||
- image-build-amd
|
||||
|
||||
@@ -86,7 +85,7 @@ steps:
|
||||
- tests/v1/metrics
|
||||
- tests/entrypoints/openai/correctness/test_lmeval.py
|
||||
commands:
|
||||
- bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
|
||||
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
# split the test to avoid interference
|
||||
- pytest -v -s -m 'not cpu_test' v1/core
|
||||
@@ -138,26 +137,11 @@ steps:
|
||||
- vllm/v1/spec_decode/extract_hidden_states.py
|
||||
- vllm/model_executor/models/extract_hidden_states.py
|
||||
- vllm/transformers_utils/configs/extract_hidden_states.py
|
||||
- vllm/distributed/kv_transfer/kv_connector/v1/example_hidden_states_connector.py
|
||||
- tests/v1/kv_connector/extract_hidden_states_integration
|
||||
commands:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s v1/kv_connector/extract_hidden_states_integration
|
||||
|
||||
- label: Extract Hidden States Integration (2 GPUs)
|
||||
key: extract-hidden-states-integration-2-gpus
|
||||
timeout_in_minutes: 20
|
||||
num_devices: 2
|
||||
source_file_dependencies:
|
||||
- vllm/v1/spec_decode/extract_hidden_states.py
|
||||
- vllm/model_executor/models/extract_hidden_states.py
|
||||
- vllm/transformers_utils/configs/extract_hidden_states.py
|
||||
- vllm/distributed/kv_transfer/kv_connector/v1/example_hidden_states_connector.py
|
||||
- tests/v1/kv_connector/extract_hidden_states_integration
|
||||
commands:
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s -m 'distributed' v1/kv_connector/extract_hidden_states_integration
|
||||
|
||||
- label: Regression
|
||||
key: regression
|
||||
timeout_in_minutes: 20
|
||||
@@ -181,7 +165,6 @@ steps:
|
||||
working_dir: "/vllm-workspace/tests" # optional
|
||||
|
||||
- label: Examples
|
||||
device: h200_35gb
|
||||
key: examples
|
||||
timeout_in_minutes: 45
|
||||
working_dir: "/vllm-workspace/examples"
|
||||
@@ -252,7 +235,6 @@ steps:
|
||||
- bash standalone_tests/python_only_compile.sh
|
||||
|
||||
- label: Async Engine, Inputs, Utils, Worker
|
||||
device: h200_35gb
|
||||
key: async-engine-inputs-utils-worker
|
||||
timeout_in_minutes: 50
|
||||
source_file_dependencies:
|
||||
@@ -296,7 +278,6 @@ steps:
|
||||
- vllm/model_executor/layers/quantization/quark/
|
||||
- vllm/multimodal/
|
||||
- vllm/outputs.py
|
||||
- vllm/parser/
|
||||
- vllm/platforms/
|
||||
- vllm/pooling_params.py
|
||||
- vllm/ray/
|
||||
@@ -308,7 +289,6 @@ steps:
|
||||
- vllm/transformers_utils/
|
||||
- vllm/utils/
|
||||
- vllm/v1/
|
||||
- tests/test_envs.py
|
||||
- tests/test_inputs.py
|
||||
- tests/test_outputs.py
|
||||
- tests/test_pooling_params.py
|
||||
@@ -316,25 +296,24 @@ steps:
|
||||
- tests/multimodal
|
||||
- tests/renderers
|
||||
- tests/standalone_tests/lazy_imports.py
|
||||
- tests/tokenizers_
|
||||
- tests/reasoning
|
||||
- tests/tool_parsers
|
||||
- tests/tokenizers_
|
||||
- tests/parser
|
||||
- tests/transformers_utils
|
||||
- tests/config
|
||||
device: cpu-small
|
||||
commands:
|
||||
- python3 standalone_tests/lazy_imports.py
|
||||
- pytest -v -s test_envs.py
|
||||
- pytest -v -s test_inputs.py
|
||||
- pytest -v -s test_outputs.py
|
||||
- pytest -v -s test_pooling_params.py
|
||||
- pytest -v -s test_ray_env.py
|
||||
- pytest -v -s -m 'cpu_test' multimodal
|
||||
- pytest -v -s renderers
|
||||
- pytest -v -s tokenizers_
|
||||
- pytest -v -s reasoning --ignore=reasoning/test_seedoss_reasoning_parser.py --ignore=reasoning/test_glm4_moe_reasoning_parser.py
|
||||
- pytest -v -s tool_parsers
|
||||
- pytest -v -s tokenizers_
|
||||
- pytest -v -s parser
|
||||
- pytest -v -s transformers_utils
|
||||
- pytest -v -s config
|
||||
@@ -385,10 +364,8 @@ steps:
|
||||
- VLLM_TEST_MODEL=deepseek-ai/DeepSeek-V2-Lite-Chat pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle[TRITON_MLA]
|
||||
- VLLM_TEST_MODEL=Qwen/Qwen3-30B-A3B-Thinking-2507-FP8 pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle[FLASH_ATTN]
|
||||
- pytest -v -s v1/determinism/test_nvfp4_batch_invariant.py
|
||||
- pytest -v -s v1/determinism/test_nvfp4_batch_invariant_scaled_mm.py
|
||||
|
||||
- label: Acceptance Length Test (Large Models) # optional
|
||||
device: h200_35gb
|
||||
key: acceptance-length-test-large-models
|
||||
timeout_in_minutes: 25
|
||||
gpu: h100
|
||||
|
||||
@@ -14,12 +14,5 @@ steps:
|
||||
commands:
|
||||
- apt-get update && apt-get install -y curl libsodium23
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
# Dump tracebacks of all threads if a test hangs, so a wedged GPU/CUDA
|
||||
# init surfaces a stack instead of silently stalling.
|
||||
- export PYTHONFAULTHANDLER=1
|
||||
# Per-test watchdog: a single hung test (e.g. stuck during engine/CUDA
|
||||
# init) fails fast with a traceback instead of running until the global
|
||||
# build timeout. The `thread` method also handles hangs inside C/CUDA
|
||||
# calls that the signal method cannot interrupt.
|
||||
- pytest -v -s model_executor -m '(not slow_test)' --timeout=900 --timeout-method=thread
|
||||
- pytest -v -s entrypoints/openai/completion/test_tensorizer_entrypoint.py --timeout=900 --timeout-method=thread
|
||||
- pytest -v -s model_executor -m '(not slow_test)'
|
||||
- pytest -v -s entrypoints/openai/completion/test_tensorizer_entrypoint.py
|
||||
|
||||
@@ -3,7 +3,6 @@ depends_on:
|
||||
- image-build
|
||||
steps:
|
||||
- label: Model Runner V2 Core Tests
|
||||
device: h200_35gb
|
||||
key: model-runner-v2-core-tests
|
||||
timeout_in_minutes: 45
|
||||
source_file_dependencies:
|
||||
@@ -27,7 +26,6 @@ steps:
|
||||
- pytest -v -s entrypoints/llm/test_struct_output_generate.py -k "xgrammar and not speculative_config6 and not speculative_config7 and not speculative_config8 and not speculative_config0"
|
||||
|
||||
- label: Model Runner V2 Examples
|
||||
device: h200_35gb
|
||||
key: model-runner-v2-examples
|
||||
timeout_in_minutes: 45
|
||||
working_dir: "/vllm-workspace/examples"
|
||||
@@ -94,16 +92,13 @@ steps:
|
||||
- vllm/v1/worker/gpu_worker.py
|
||||
- tests/distributed/test_pipeline_parallel.py
|
||||
- tests/distributed/test_pp_cudagraph.py
|
||||
- tests/v1/distributed/test_pp_dp_v2.py
|
||||
commands:
|
||||
- set -x
|
||||
- export VLLM_USE_V2_MODEL_RUNNER=1
|
||||
- pytest -v -s distributed/test_pipeline_parallel.py -k "not ray and not Jamba"
|
||||
- pytest -v -s distributed/test_pp_cudagraph.py -k "not ray"
|
||||
- pytest -v -s v1/distributed/test_pp_dp_v2.py
|
||||
|
||||
- label: Model Runner V2 Spec Decode
|
||||
device: h200_35gb
|
||||
key: model-runner-v2-spec-decode
|
||||
timeout_in_minutes: 30
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
|
||||
@@ -18,7 +18,6 @@ steps:
|
||||
torch_nightly: {}
|
||||
|
||||
- label: Basic Models Tests (Extra Initialization) %N
|
||||
device: h200_35gb
|
||||
key: basic-models-tests-extra-initialization
|
||||
timeout_in_minutes: 45
|
||||
source_file_dependencies:
|
||||
@@ -35,7 +34,6 @@ steps:
|
||||
torch_nightly: {}
|
||||
|
||||
- label: Basic Models Tests (Other)
|
||||
device: h200_35gb
|
||||
key: basic-models-tests-other
|
||||
timeout_in_minutes: 45
|
||||
source_file_dependencies:
|
||||
@@ -58,3 +56,35 @@ steps:
|
||||
device: cpu-small
|
||||
commands:
|
||||
- pytest -v -s models/test_utils.py models/test_vision.py
|
||||
|
||||
- label: Transformers Nightly Models
|
||||
key: transformers-nightly-models
|
||||
working_dir: "/vllm-workspace/"
|
||||
optional: true
|
||||
soft_fail: true
|
||||
commands:
|
||||
- pip install --upgrade git+https://github.com/huggingface/transformers
|
||||
- pytest -v -s tests/models/test_initialization.py
|
||||
- pytest -v -s tests/models/test_transformers.py
|
||||
- pytest -v -s tests/models/multimodal/processing/
|
||||
- pytest -v -s tests/models/multimodal/test_mapping.py
|
||||
- python3 examples/basic/offline_inference/chat.py
|
||||
- python3 examples/generate/multimodal/vision_language_offline.py --model-type qwen2_5_vl
|
||||
# Whisper needs spawn method to avoid deadlock
|
||||
- VLLM_WORKER_MULTIPROC_METHOD=spawn python3 examples/generate/multimodal/audio_language_offline.py --model-type whisper
|
||||
|
||||
- label: Transformers Backward Compatibility Models Test
|
||||
key: transformers-backward-compatibility-models-test
|
||||
working_dir: "/vllm-workspace/"
|
||||
optional: true
|
||||
soft_fail: true
|
||||
commands:
|
||||
- pip install transformers==4.57.5
|
||||
- pytest -v -s tests/models/test_initialization.py
|
||||
- pytest -v -s tests/models/test_transformers.py
|
||||
- pytest -v -s tests/models/multimodal/processing/
|
||||
- pytest -v -s tests/models/multimodal/test_mapping.py
|
||||
- python3 examples/basic/offline_inference/chat.py
|
||||
- python3 examples/generate/multimodal/vision_language_offline.py --model-type qwen2_5_vl
|
||||
# Whisper needs spawn method to avoid deadlock
|
||||
- VLLM_WORKER_MULTIPROC_METHOD=spawn python3 examples/generate/multimodal/audio_language_offline.py --model-type whisper
|
||||
|
||||
@@ -50,8 +50,7 @@ steps:
|
||||
mirror:
|
||||
torch_nightly: {}
|
||||
amd:
|
||||
device: mi325_1
|
||||
timeout_in_minutes: 90
|
||||
device: mi300_1
|
||||
depends_on:
|
||||
- image-build-amd
|
||||
commands:
|
||||
@@ -60,7 +59,6 @@ steps:
|
||||
- pytest -v -s models/language/generation -m hybrid_model --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
|
||||
|
||||
- label: Language Models Test (Extended Generation) # 80min
|
||||
device: h200_35gb
|
||||
key: language-models-test-extended-generation
|
||||
timeout_in_minutes: 110
|
||||
optional: true
|
||||
@@ -86,7 +84,6 @@ steps:
|
||||
- pytest -v -s models/language/generation_ppl_test
|
||||
|
||||
- label: Language Models Test (Extended Pooling) # 36min
|
||||
device: h200_35gb
|
||||
key: language-models-test-extended-pooling
|
||||
timeout_in_minutes: 50
|
||||
optional: true
|
||||
@@ -97,7 +94,7 @@ steps:
|
||||
- pytest -v -s models/language/pooling -m 'not core_model'
|
||||
mirror:
|
||||
amd:
|
||||
device: mi325_1
|
||||
device: mi300_1
|
||||
timeout_in_minutes: 100
|
||||
depends_on:
|
||||
- image-build-amd
|
||||
|
||||
@@ -15,7 +15,7 @@ steps:
|
||||
- pytest -v -s models/multimodal/generation/test_ultravox.py -m core_model
|
||||
mirror:
|
||||
amd:
|
||||
device: mi325_1
|
||||
device: mi300_1
|
||||
depends_on:
|
||||
- image-build-amd
|
||||
|
||||
@@ -33,12 +33,11 @@ steps:
|
||||
- pytest -v -s models/multimodal/generation/test_vit_cudagraph.py -m core_model
|
||||
mirror:
|
||||
amd:
|
||||
device: mi325_1
|
||||
device: mi300_1
|
||||
depends_on:
|
||||
- image-build-amd
|
||||
|
||||
- label: "Multi-Modal Models (Standard) 3: llava + qwen2_vl"
|
||||
device: h200_35gb
|
||||
key: multi-modal-models-standard-3-llava-qwen2-vl
|
||||
timeout_in_minutes: 45
|
||||
source_file_dependencies:
|
||||
@@ -50,12 +49,11 @@ steps:
|
||||
- pytest -v -s models/multimodal/generation/test_qwen2_vl.py -m core_model
|
||||
mirror:
|
||||
amd:
|
||||
device: mi325_1
|
||||
device: mi300_1
|
||||
depends_on:
|
||||
- image-build-amd
|
||||
|
||||
- label: "Multi-Modal Models (Standard) 4: other + whisper"
|
||||
device: h200_35gb
|
||||
key: multi-modal-models-standard-4-other-whisper
|
||||
timeout_in_minutes: 45
|
||||
source_file_dependencies:
|
||||
@@ -94,7 +92,6 @@ steps:
|
||||
- pytest -v -s models/multimodal/processing/test_tensor_schema.py
|
||||
|
||||
- label: Multi-Modal Accuracy Eval (Small Models) # 50min
|
||||
device: h200_35gb
|
||||
key: multi-modal-accuracy-eval-small-models
|
||||
timeout_in_minutes: 70
|
||||
working_dir: "/vllm-workspace/.buildkite/lm-eval-harness"
|
||||
@@ -118,12 +115,11 @@ steps:
|
||||
- pytest -v -s models/multimodal/test_mapping.py
|
||||
mirror:
|
||||
amd:
|
||||
device: mi325_1
|
||||
device: mi300_1
|
||||
depends_on:
|
||||
- image-build-amd
|
||||
|
||||
- label: Multi-Modal Models (Extended Generation 2)
|
||||
device: h200_35gb
|
||||
key: multi-modal-models-extended-generation-2
|
||||
optional: true
|
||||
source_file_dependencies:
|
||||
@@ -134,7 +130,6 @@ steps:
|
||||
- pytest -v -s models/multimodal/generation/test_common.py -m 'split(group=0) and not core_model'
|
||||
|
||||
- label: Multi-Modal Models (Extended Generation 3)
|
||||
device: h200_35gb
|
||||
key: multi-modal-models-extended-generation-3
|
||||
optional: true
|
||||
source_file_dependencies:
|
||||
@@ -153,12 +148,3 @@ steps:
|
||||
- tests/models/multimodal/pooling
|
||||
commands:
|
||||
- pytest -v -s models/multimodal/pooling -m 'not core_model'
|
||||
mirror:
|
||||
amd:
|
||||
device: mi325_1
|
||||
timeout_in_minutes: 60
|
||||
depends_on:
|
||||
- image-build-amd
|
||||
source_file_dependencies:
|
||||
- vllm/
|
||||
- tests/models/multimodal/pooling
|
||||
|
||||
@@ -37,20 +37,6 @@ steps:
|
||||
- pytest -v -s plugins_tests/test_scheduler_plugins.py
|
||||
- pip install -e ./plugins/vllm_add_dummy_model
|
||||
- pytest -v -s distributed/test_distributed_oot.py
|
||||
- pytest -v -s plugins_tests/test_oot_registration_online.py # it needs a clean process
|
||||
- pytest -v -s plugins_tests/test_oot_registration_offline.py # it needs a clean process
|
||||
- pytest -v -s plugins_tests/lora_resolvers # unit tests for in-tree lora resolver plugins
|
||||
|
||||
|
||||
- label: GGUF Plugin
|
||||
key: gguf-plugin
|
||||
device: h200_18gb
|
||||
timeout_in_minutes: 30
|
||||
soft_fail: true
|
||||
optional: true
|
||||
source_file_dependencies:
|
||||
- vllm/model_executor/layers/quantization
|
||||
- tests/plugins_tests/test_gguf_plugin.py
|
||||
commands:
|
||||
- pip install "vllm-gguf-plugin >= 0.0.2"
|
||||
- pytest -v -s plugins_tests/gguf
|
||||
- pytest -v -s entrypoints/openai/chat_completion/test_oot_registration.py # it needs a clean process
|
||||
- pytest -v -s models/test_oot_registration.py # it needs a clean process
|
||||
- pytest -v -s plugins/lora_resolvers # unit tests for in-tree lora resolver plugins
|
||||
|
||||
@@ -3,7 +3,6 @@ depends_on:
|
||||
- image-build
|
||||
steps:
|
||||
- label: PyTorch Compilation Unit Tests
|
||||
device: h200_35gb
|
||||
key: pytorch-compilation-unit-tests
|
||||
timeout_in_minutes: 10
|
||||
source_file_dependencies:
|
||||
|
||||
@@ -21,18 +21,6 @@ steps:
|
||||
- uv pip install --system conch-triton-kernels
|
||||
- VLLM_TEST_FORCE_LOAD_FORMAT=auto pytest -v -s quantization/ --ignore quantization/test_blackwell_moe.py
|
||||
|
||||
- label: Quantized Fusions
|
||||
key: quantized-fusions
|
||||
timeout_in_minutes: 30
|
||||
source_file_dependencies:
|
||||
- tests/fusion
|
||||
- vllm/model_executor/layers/fusion
|
||||
- vllm/model_executor/kernels/linear
|
||||
- vllm/model_executor/layers/quantization/compressed_tensors
|
||||
- vllm/model_executor/layers/quantization/modelopt.py
|
||||
commands:
|
||||
- pytest -v -s fusion/
|
||||
|
||||
- label: Quantized MoE Test (B200)
|
||||
key: quantized-moe-test-b200
|
||||
timeout_in_minutes: 60
|
||||
|
||||
@@ -1,107 +0,0 @@
|
||||
group: Rust Frontend E2E
|
||||
depends_on:
|
||||
- image-build
|
||||
steps:
|
||||
- label: Rust Frontend OpenAI Coverage
|
||||
timeout_in_minutes: 90
|
||||
device: h200_18gb
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- rust/
|
||||
- vllm/benchmarks/
|
||||
- vllm/entrypoints/openai/
|
||||
- vllm/entrypoints/serve/
|
||||
- vllm/v1/sample/
|
||||
- tests/utils.py
|
||||
- tests/benchmarks/test_serve_cli.py
|
||||
- tests/entrypoints/openai/chat_completion/test_chat_completion.py
|
||||
# - tests/entrypoints/openai/chat_completion/test_chat_logit_bias_validation.py
|
||||
|
||||
# - tests/entrypoints/openai/completion/test_prompt_validation.py
|
||||
- tests/entrypoints/openai/completion/test_shutdown.py
|
||||
# - tests/entrypoints/openai/test_return_token_ids.py
|
||||
# - tests/entrypoints/openai/test_uds.py
|
||||
- tests/v1/sample/test_logprobs_e2e.py
|
||||
commands:
|
||||
- export VLLM_USE_RUST_FRONTEND=1
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s benchmarks/test_serve_cli.py -k "not insecure and not (test_bench_serve and not test_bench_serve_chat)"
|
||||
- pytest -v -s entrypoints/openai/chat_completion/test_chat_completion.py
|
||||
# - pytest -v -s entrypoints/openai/chat_completion/test_chat_logit_bias_validation.py -k "not invalid"
|
||||
|
||||
# - pytest -v -s entrypoints/openai/completion/test_prompt_validation.py -k "not prompt_embeds"
|
||||
- pytest -v -s entrypoints/openai/completion/test_shutdown.py -k "not engine_failure and not test_abort_timeout_exits_quickly"
|
||||
# - pytest -v -s entrypoints/openai/test_return_token_ids.py
|
||||
# - pytest -v -s entrypoints/openai/test_uds.py
|
||||
- pytest -v -s v1/sample/test_logprobs_e2e.py -k "test_prompt_logprobs_e2e_server"
|
||||
|
||||
- label: Rust Frontend Serve/Admin Coverage
|
||||
timeout_in_minutes: 60
|
||||
device: h200_18gb
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- rust/
|
||||
- vllm/entrypoints/openai/
|
||||
- vllm/entrypoints/serve/
|
||||
- vllm/v1/engine/
|
||||
- tests/utils.py
|
||||
# - tests/entrypoints/serve/dev/rpc/test_collective_rpc.py
|
||||
- tests/entrypoints/serve/disagg/test_serving_tokens.py
|
||||
- tests/entrypoints/serve/instrumentator/test_basic.py
|
||||
- tests/entrypoints/serve/instrumentator/test_metrics.py
|
||||
# - tests/entrypoints/serve/dev/test_sleep.py
|
||||
commands:
|
||||
- export VLLM_USE_RUST_FRONTEND=1
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
# - pytest -v -s entrypoints/serve/dev/rpc/test_collective_rpc.py
|
||||
- pytest -v -s entrypoints/serve/instrumentator/test_basic.py -k "not show_version and not server_load"
|
||||
- pytest -v -s entrypoints/serve/disagg/test_serving_tokens.py -k "not stream and not lora and not test_generate_logprobs and not stop_string_workflow"
|
||||
- pytest -v -s entrypoints/serve/instrumentator/test_metrics.py -k "text and not show and not run_batch and not test_metrics_counts and not test_metrics_exist"
|
||||
# - pytest -v -s entrypoints/serve/dev/test_sleep.py
|
||||
|
||||
- label: Rust Frontend Core Correctness
|
||||
timeout_in_minutes: 30
|
||||
device: h200_18gb
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- rust/
|
||||
- vllm/entrypoints/openai/
|
||||
- tests/utils.py
|
||||
- tests/entrypoints/openai/correctness/test_lmeval.py
|
||||
commands:
|
||||
- export VLLM_USE_RUST_FRONTEND=1
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine
|
||||
|
||||
- label: Rust Frontend Tool Use
|
||||
timeout_in_minutes: 60
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- rust/
|
||||
- vllm/entrypoints/openai/
|
||||
- vllm/tool_parsers/
|
||||
- tests/utils.py
|
||||
- tests/tool_use/
|
||||
commands:
|
||||
- export VLLM_USE_RUST_FRONTEND=1
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- pytest -v -s tool_use --ignore=tool_use/mistral --models llama3.2 -k "not test_response_format_with_tool_choice_required and not test_parallel_tool_calls_false and not test_tool_call_and_choice"
|
||||
|
||||
- label: Rust Frontend Distributed
|
||||
timeout_in_minutes: 30
|
||||
num_devices: 4
|
||||
working_dir: "/vllm-workspace/tests"
|
||||
source_file_dependencies:
|
||||
- rust/
|
||||
- vllm/distributed/
|
||||
- vllm/engine/
|
||||
- vllm/executor/
|
||||
- vllm/v1/engine/
|
||||
- vllm/v1/worker/
|
||||
- tests/utils.py
|
||||
- tests/v1/distributed/test_internal_lb_dp.py
|
||||
commands:
|
||||
- export VLLM_USE_RUST_FRONTEND=1
|
||||
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
|
||||
- export NCCL_CUMEM_HOST_ENABLE=0
|
||||
- TP_SIZE=1 DP_SIZE=4 pytest -v -s v1/distributed/test_internal_lb_dp.py -k "not 4 and not server_info"
|
||||
@@ -1,30 +0,0 @@
|
||||
group: Rust Frontend Cargo
|
||||
depends_on: []
|
||||
steps:
|
||||
- label: Rust Frontend Cargo Style + Clippy
|
||||
key: rust-frontend-cargo-style-clippy
|
||||
depends_on: []
|
||||
timeout_in_minutes: 30
|
||||
device: cpu-medium
|
||||
no_plugin: true
|
||||
source_file_dependencies:
|
||||
- rust/
|
||||
- rust-toolchain.toml
|
||||
- .buildkite/test_areas/rust_frontend_cargo.yaml
|
||||
- .buildkite/scripts/run-rust-frontend-cargo-ci.sh
|
||||
commands:
|
||||
- .buildkite/scripts/run-rust-frontend-cargo-ci.sh style-clippy
|
||||
|
||||
- label: Rust Frontend Cargo Tests
|
||||
key: rust-frontend-cargo-tests
|
||||
depends_on: []
|
||||
timeout_in_minutes: 30
|
||||
device: cpu-medium
|
||||
no_plugin: true
|
||||
source_file_dependencies:
|
||||
- rust/
|
||||
- rust-toolchain.toml
|
||||
- .buildkite/test_areas/rust_frontend_cargo.yaml
|
||||
- .buildkite/scripts/run-rust-frontend-cargo-ci.sh
|
||||
commands:
|
||||
- .buildkite/scripts/run-rust-frontend-cargo-ci.sh test
|
||||
@@ -3,7 +3,6 @@ depends_on:
|
||||
- image-build
|
||||
steps:
|
||||
- label: Samplers Test
|
||||
device: h200_35gb
|
||||
key: samplers-test
|
||||
timeout_in_minutes: 75
|
||||
source_file_dependencies:
|
||||
@@ -11,7 +10,6 @@ steps:
|
||||
- vllm/sampling_metadata.py
|
||||
- tests/samplers
|
||||
- tests/conftest.py
|
||||
- vllm/entrypoints/generate/beam_search
|
||||
commands:
|
||||
# VLLM_USE_FLASHINFER_SAMPLER defaults to 1 now, so we need to pin both
|
||||
# values explicitly to still cover the PyTorch-native (Triton) path.
|
||||
|
||||
@@ -32,26 +32,10 @@ steps:
|
||||
source_file_dependencies:
|
||||
- vllm/v1/spec_decode/
|
||||
- vllm/v1/worker/gpu/spec_decode/
|
||||
- vllm/v1/attention/backends/
|
||||
- vllm/transformers_utils/configs/speculators/
|
||||
- tests/v1/e2e/spec_decode/
|
||||
commands:
|
||||
- pytest -v -s v1/e2e/spec_decode -k "speculators or mtp_correctness"
|
||||
mirror:
|
||||
amd:
|
||||
device: mi325_1
|
||||
timeout_in_minutes: 65
|
||||
depends_on:
|
||||
- image-build-amd
|
||||
source_file_dependencies:
|
||||
- vllm/v1/spec_decode/
|
||||
- vllm/v1/worker/gpu/spec_decode/
|
||||
- vllm/model_executor/model_loader/
|
||||
- vllm/v1/sample/
|
||||
- vllm/model_executor/layers/
|
||||
- vllm/transformers_utils/configs/speculators/
|
||||
- tests/v1/e2e/spec_decode/
|
||||
- vllm/platforms/rocm.py
|
||||
|
||||
- label: Spec Decode Speculators + MTP Nightly B200
|
||||
key: spec-decode-speculators-mtp-nightly-b200
|
||||
@@ -76,20 +60,6 @@ steps:
|
||||
- tests/v1/e2e/spec_decode/
|
||||
commands:
|
||||
- pytest -v -s v1/e2e/spec_decode -k "ngram or suffix"
|
||||
mirror:
|
||||
amd:
|
||||
device: mi325_1
|
||||
timeout_in_minutes: 65
|
||||
depends_on:
|
||||
- image-build-amd
|
||||
source_file_dependencies:
|
||||
- vllm/v1/spec_decode/
|
||||
- vllm/v1/worker/gpu/spec_decode/
|
||||
- vllm/model_executor/model_loader/
|
||||
- vllm/v1/sample/
|
||||
- vllm/model_executor/layers/
|
||||
- tests/v1/e2e/spec_decode/
|
||||
- vllm/platforms/rocm.py
|
||||
|
||||
- label: Spec Decode Draft Model
|
||||
key: spec-decode-draft-model
|
||||
@@ -101,20 +71,6 @@ steps:
|
||||
- tests/v1/e2e/spec_decode/
|
||||
commands:
|
||||
- pytest -v -s v1/e2e/spec_decode -k "draft_model or no_sync or batch_inference"
|
||||
mirror:
|
||||
amd:
|
||||
device: mi325_1
|
||||
timeout_in_minutes: 50
|
||||
depends_on:
|
||||
- image-build-amd
|
||||
source_file_dependencies:
|
||||
- vllm/v1/spec_decode/
|
||||
- vllm/v1/worker/gpu/spec_decode/
|
||||
- vllm/model_executor/model_loader/
|
||||
- vllm/v1/sample/
|
||||
- vllm/model_executor/layers/
|
||||
- tests/v1/e2e/spec_decode/
|
||||
- vllm/platforms/rocm.py
|
||||
|
||||
- label: Spec Decode Draft Model Nightly B200
|
||||
key: spec-decode-draft-model-nightly-b200
|
||||
|
||||
@@ -2,7 +2,6 @@
|
||||
/build
|
||||
dist
|
||||
vllm/*.so
|
||||
vllm/vllm-rs
|
||||
|
||||
# Byte-compiled / optimized / DLL files
|
||||
__pycache__/
|
||||
@@ -32,4 +31,3 @@ share/python-wheels/
|
||||
.installed.cfg
|
||||
*.egg
|
||||
MANIFEST
|
||||
rust/target/
|
||||
|
||||
+26
-38
@@ -23,14 +23,9 @@
|
||||
|
||||
# Any change to the VllmConfig changes can have a large user-facing impact,
|
||||
# so spam a lot of people
|
||||
/vllm/config @WoosukKwon @youkaichao @robertgshaw2-redhat @mgoin @tlrmchlsmth @houseroad @yewentao256 @ProExpertProg
|
||||
/vllm/config @WoosukKwon @youkaichao @robertgshaw2-redhat @mgoin @tlrmchlsmth @houseroad @hmellor @yewentao256 @ProExpertProg
|
||||
/vllm/config/cache.py @heheda12345
|
||||
|
||||
# Config utils
|
||||
/vllm/config/utils.py @hmellor
|
||||
/vllm/engine/arg_utils.py @hmellor
|
||||
/vllm/utils/argparse_utils.py
|
||||
|
||||
# Entrypoints
|
||||
/vllm/entrypoints/anthropic @mgoin @DarkLight1337
|
||||
/vllm/entrypoints/cli @hmellor @mgoin @DarkLight1337 @russellb
|
||||
@@ -39,19 +34,12 @@
|
||||
/vllm/entrypoints/speech_to_text/realtime @njhill
|
||||
/vllm/entrypoints/speech_to_text @NickLucche
|
||||
/vllm/entrypoints/pooling @noooop
|
||||
/vllm/entrypoints/serve/sagemaker @DarkLight1337
|
||||
/vllm/entrypoints/sagemaker @DarkLight1337
|
||||
/vllm/entrypoints/serve @njhill
|
||||
/vllm/entrypoints/*.py @njhill
|
||||
/vllm/entrypoints/chat_utils.py @DarkLight1337
|
||||
/vllm/entrypoints/offline_utils.py @DarkLight1337
|
||||
/vllm/entrypoints/llm.py @DarkLight1337
|
||||
|
||||
# Rust Frontend
|
||||
/rust/ @BugenZhao @njhill
|
||||
/build_rust.sh @BugenZhao @njhill
|
||||
/rust-toolchain.toml @BugenZhao @njhill
|
||||
/.buildkite/test_areas/rust* @BugenZhao @njhill
|
||||
|
||||
# Input/Output Processing
|
||||
/vllm/sampling_params.py @njhill @NickLucche
|
||||
/vllm/pooling_params.py @noooop @DarkLight1337
|
||||
@@ -80,27 +68,25 @@
|
||||
/vllm/v1/worker/kv_connector_model_runner_mixin.py @orozery @NickLucche
|
||||
|
||||
# Model runner V2
|
||||
/vllm/v1/worker/gpu @WoosukKwon @njhill @yewentao256
|
||||
/vllm/v1/worker/gpu @WoosukKwon @njhill
|
||||
/vllm/v1/worker/gpu/kv_connector.py @orozery
|
||||
|
||||
# CI & building
|
||||
/.buildkite @Harry-Chen @khluu
|
||||
/docker/Dockerfile @Harry-Chen @khluu
|
||||
/pyproject.toml @khluu
|
||||
/setup.py @khluu
|
||||
/.buildkite @Harry-Chen
|
||||
/docker/Dockerfile @Harry-Chen
|
||||
|
||||
# Test ownership
|
||||
/.buildkite/lm-eval-harness @mgoin
|
||||
/.buildkite/lm-eval-harness @mgoin
|
||||
/tests/distributed/test_multi_node_assignment.py @youkaichao
|
||||
/tests/distributed/test_pipeline_parallel.py @youkaichao
|
||||
/tests/distributed/test_same_node.py @youkaichao
|
||||
/tests/entrypoints @DarkLight1337 @robertgshaw2-redhat @aarnphm @NickLucche @AndreasKaratzas
|
||||
/tests/evals @mgoin @vadiklyutiy @AndreasKaratzas
|
||||
/tests/kernels @mgoin @tlrmchlsmth @WoosukKwon @yewentao256 @zyongye @AndreasKaratzas
|
||||
/tests/entrypoints @DarkLight1337 @robertgshaw2-redhat @aarnphm @NickLucche
|
||||
/tests/evals @mgoin @vadiklyutiy
|
||||
/tests/kernels @mgoin @tlrmchlsmth @WoosukKwon @yewentao256 @zyongye
|
||||
/tests/kernels/ir @ProExpertProg @tjtanaa
|
||||
/tests/models @DarkLight1337 @ywang96 @AndreasKaratzas
|
||||
/tests/models @DarkLight1337 @ywang96
|
||||
/tests/multimodal @DarkLight1337 @ywang96 @NickLucche
|
||||
/tests/quantization @mgoin @robertgshaw2-redhat @yewentao256 @pavanimajety @zyongye @AndreasKaratzas
|
||||
/tests/quantization @mgoin @robertgshaw2-redhat @yewentao256 @pavanimajety @zyongye
|
||||
/tests/test_inputs.py @DarkLight1337 @ywang96
|
||||
/tests/entrypoints/llm/test_struct_output_generate.py @mgoin @russellb @aarnphm
|
||||
/tests/v1/structured_output @mgoin @russellb @aarnphm
|
||||
@@ -167,7 +153,9 @@ mkdocs.yaml @hmellor
|
||||
/vllm/model_executor/models/deepseek_mtp.py @luccafong
|
||||
|
||||
# DeepseekV4-specific files
|
||||
/vllm/models/deepseek_v4 @zyongye
|
||||
/vllm/v1/attention/ops/deepseek_v4_ops @zyongye
|
||||
/vllm/model_executor/layers/deepseek_compressor.py @zyongye
|
||||
/vllm/model_executor/layers/deepseek_v4_attention.py @zyongye
|
||||
/vllm/model_executor/layers/sparse_attn_indexer.py @zyongye
|
||||
|
||||
# Mistral-specific files
|
||||
@@ -184,21 +172,21 @@ mkdocs.yaml @hmellor
|
||||
/vllm/model_executor/layers/fla @ZJY0516 @vadiklyutiy
|
||||
|
||||
# ROCm related: specify owner with write access to notify AMD folks for careful code review
|
||||
/vllm/**/*rocm* @tjtanaa @dllehr-amd
|
||||
/docker/Dockerfile.rocm* @tjtanaa @dllehr-amd @AndreasKaratzas
|
||||
/vllm/v1/attention/backends/rocm*.py @tjtanaa @dllehr-amd
|
||||
/vllm/v1/attention/backends/mla/rocm*.py @tjtanaa @dllehr-amd
|
||||
/vllm/v1/attention/ops/rocm*.py @tjtanaa @dllehr-amd
|
||||
/vllm/model_executor/layers/fused_moe/rocm*.py @tjtanaa @dllehr-amd
|
||||
/csrc/rocm @tjtanaa @dllehr-amd
|
||||
/requirements/*rocm* @tjtanaa @AndreasKaratzas
|
||||
/tests/**/*rocm* @tjtanaa @AndreasKaratzas
|
||||
/vllm/**/*rocm* @tjtanaa
|
||||
/docker/Dockerfile.rocm* @gshtras @tjtanaa
|
||||
/vllm/v1/attention/backends/rocm*.py @gshtras @tjtanaa
|
||||
/vllm/v1/attention/backends/mla/rocm*.py @gshtras @tjtanaa
|
||||
/vllm/v1/attention/ops/rocm*.py @gshtras @tjtanaa
|
||||
/vllm/model_executor/layers/fused_moe/rocm*.py @gshtras @tjtanaa
|
||||
/csrc/rocm @gshtras @tjtanaa
|
||||
/requirements/*rocm* @tjtanaa
|
||||
/tests/**/*rocm* @tjtanaa
|
||||
/docs/**/*rocm* @tjtanaa
|
||||
/vllm/**/*quark* @tjtanaa
|
||||
/tests/**/*quark* @tjtanaa @AndreasKaratzas
|
||||
/tests/**/*quark* @tjtanaa
|
||||
/docs/**/*quark* @tjtanaa
|
||||
/vllm/**/*aiter* @tjtanaa @AndreasKaratzas
|
||||
/tests/**/*aiter* @tjtanaa @AndreasKaratzas
|
||||
/vllm/**/*aiter* @tjtanaa
|
||||
/tests/**/*aiter* @tjtanaa
|
||||
|
||||
# TPU
|
||||
/vllm/v1/worker/tpu* @NickLucche
|
||||
|
||||
@@ -21,6 +21,7 @@ updates:
|
||||
- dependency-name: "torchvision"
|
||||
- dependency-name: "xformers"
|
||||
- dependency-name: "lm-format-enforcer"
|
||||
- dependency-name: "gguf"
|
||||
- dependency-name: "compressed-tensors"
|
||||
- dependency-name: "ray[cgraph]" # Ray Compiled Graph
|
||||
- dependency-name: "lm-eval"
|
||||
|
||||
+16
-20
@@ -21,9 +21,6 @@ pull_request_rules:
|
||||
- check-failure=pre-commit
|
||||
- -closed
|
||||
- -draft
|
||||
- or:
|
||||
- label=ready
|
||||
- label=verified
|
||||
actions:
|
||||
comment:
|
||||
message: |
|
||||
@@ -39,6 +36,18 @@ pull_request_rules:
|
||||
|
||||
For future commits, `pre-commit` will run automatically on changed files before each commit.
|
||||
|
||||
> [!TIP]
|
||||
> <details>
|
||||
> <summary>Is <code>mypy</code> failing?</summary>
|
||||
> <br/>
|
||||
> <code>mypy</code> is run differently in CI. If the failure is related to this check, please use the following command to run it locally:
|
||||
>
|
||||
> ```bash
|
||||
> # For mypy (substitute "3.10" with the failing version if needed)
|
||||
> pre-commit run --hook-stage manual mypy-3.10
|
||||
> ```
|
||||
> </details>
|
||||
|
||||
- name: comment-dco-failure
|
||||
description: Comment on PR when DCO check fails
|
||||
conditions:
|
||||
@@ -94,19 +103,6 @@ pull_request_rules:
|
||||
add:
|
||||
- frontend
|
||||
|
||||
- name: label-rust
|
||||
description: Automatically apply rust label
|
||||
conditions:
|
||||
- label != stale
|
||||
- or:
|
||||
- files~=(?i)rust
|
||||
- title~=(?i)rust
|
||||
- title~=(?i)vllm-rs
|
||||
actions:
|
||||
label:
|
||||
add:
|
||||
- rust
|
||||
|
||||
- name: label-llama
|
||||
description: Automatically apply llama label
|
||||
conditions:
|
||||
@@ -144,12 +140,12 @@ pull_request_rules:
|
||||
- label != stale
|
||||
- or:
|
||||
- files~=^examples/.*mistral.*\.py
|
||||
- files~=^tests/.*(?:mistral|voxtral|mixtral|pixtral).*\.py
|
||||
- files~=^vllm/model_executor/models/.*(?:mistral|voxtral|mixtral|pixtral).*\.py
|
||||
- files~=^tests/.*mistral.*\.py
|
||||
- files~=^vllm/model_executor/models/.*mistral.*\.py
|
||||
- files~=^vllm/reasoning/.*mistral.*\.py
|
||||
- files~=^vllm/tool_parsers/.*mistral.*\.py
|
||||
- files~=^vllm/transformers_utils/.*(?:mistral|voxtral|pixtral).*\.py
|
||||
- title~=(?i)(?:mistral|ministral|voxtral|mixtral|pixtral)
|
||||
- files~=^vllm/transformers_utils/.*mistral.*\.py
|
||||
- title~=(?i)Mistral
|
||||
actions:
|
||||
label:
|
||||
add:
|
||||
|
||||
@@ -10,7 +10,7 @@ jobs:
|
||||
runs-on: ubuntu-latest
|
||||
steps:
|
||||
- name: Add label
|
||||
uses: actions/github-script@3a2844b7e9c422d3c10d287c895573f7108da1b3 # v9.0.0
|
||||
uses: actions/github-script@ed597411d8f924073f98dfc5c65a23a2325f34cd # v8.0.0
|
||||
with:
|
||||
script: |
|
||||
github.rest.issues.addLabels({
|
||||
|
||||
@@ -14,7 +14,7 @@ jobs:
|
||||
steps:
|
||||
- name: Label issues based on keywords
|
||||
id: label-step
|
||||
uses: actions/github-script@3a2844b7e9c422d3c10d287c895573f7108da1b3 # v9.0.0
|
||||
uses: actions/github-script@ed597411d8f924073f98dfc5c65a23a2325f34cd # v8.0.0
|
||||
with:
|
||||
script: |
|
||||
// Configuration: Add new labels and keywords here
|
||||
@@ -315,7 +315,7 @@ jobs:
|
||||
|
||||
- name: CC users for labeled issues
|
||||
if: steps.label-step.outputs.labels_added != '[]'
|
||||
uses: actions/github-script@3a2844b7e9c422d3c10d287c895573f7108da1b3 # v9.0.0
|
||||
uses: actions/github-script@ed597411d8f924073f98dfc5c65a23a2325f34cd # v8.0.0
|
||||
with:
|
||||
script: |
|
||||
// Configuration: Map labels to GitHub users to CC
|
||||
@@ -392,7 +392,7 @@ jobs:
|
||||
|
||||
- name: Request missing ROCm info from issue author
|
||||
if: contains(steps.label-step.outputs.labels_added, 'rocm') && contains(toJSON(github.event.issue.labels.*.name), 'bug')
|
||||
uses: actions/github-script@3a2844b7e9c422d3c10d287c895573f7108da1b3 # v9.0.0
|
||||
uses: actions/github-script@ed597411d8f924073f98dfc5c65a23a2325f34cd # v8.0.0
|
||||
with:
|
||||
script: |
|
||||
const body = (context.payload.issue.body || '').toLowerCase();
|
||||
|
||||
@@ -12,7 +12,7 @@ jobs:
|
||||
runs-on: ubuntu-latest
|
||||
steps:
|
||||
- name: Update PR description
|
||||
uses: actions/github-script@3a2844b7e9c422d3c10d287c895573f7108da1b3 # v9.0.0
|
||||
uses: actions/github-script@ed597411d8f924073f98dfc5c65a23a2325f34cd # v8.0.0
|
||||
with:
|
||||
script: |
|
||||
const { owner, repo } = context.repo;
|
||||
@@ -55,7 +55,7 @@ jobs:
|
||||
runs-on: ubuntu-latest
|
||||
steps:
|
||||
- name: Post welcome comment for first-time contributors
|
||||
uses: actions/github-script@3a2844b7e9c422d3c10d287c895573f7108da1b3 # v9.0.0
|
||||
uses: actions/github-script@ed597411d8f924073f98dfc5c65a23a2325f34cd # v8.0.0
|
||||
with:
|
||||
script: |
|
||||
const { owner, repo } = context.repo;
|
||||
|
||||
@@ -20,7 +20,7 @@ jobs:
|
||||
runs-on: ubuntu-latest
|
||||
steps:
|
||||
- name: Check PR label and author merge count
|
||||
uses: actions/github-script@3a2844b7e9c422d3c10d287c895573f7108da1b3 # v9.0.0
|
||||
uses: actions/github-script@ed597411d8f924073f98dfc5c65a23a2325f34cd # v8.0.0
|
||||
with:
|
||||
script: |
|
||||
const { data: pr } = await github.rest.pulls.get({
|
||||
|
||||
@@ -9,7 +9,7 @@ PATH=${cuda_home}/bin:$PATH
|
||||
LD_LIBRARY_PATH=${cuda_home}/lib64:$LD_LIBRARY_PATH
|
||||
|
||||
# Install requirements
|
||||
if [ "$(echo "$2" | cut -d. -f1)" = "12" ]; then
|
||||
if [ "$(echo $2 | cut -d. -f1)" = "12" ]; then
|
||||
sed -i 's/^nvidia-cutlass-dsl\[cu13\]>=/nvidia-cutlass-dsl>=/' requirements/cuda.txt
|
||||
fi
|
||||
$python_executable -m pip install -r requirements/build/cuda.txt -r requirements/cuda.txt
|
||||
@@ -17,10 +17,7 @@ $python_executable -m pip install -r requirements/build/cuda.txt -r requirements
|
||||
# Limit the number of parallel jobs to avoid OOM
|
||||
export MAX_JOBS=1
|
||||
# Make sure release wheels are built for the following architectures
|
||||
# Do not add +PTX here: vLLM filters torch's top-level PTX flag when it
|
||||
# converts global gencode flags into per-kernel arch lists. If a specific
|
||||
# kernel needs PTX, add +PTX to that kernel's CMake arch list instead.
|
||||
export TORCH_CUDA_ARCH_LIST="7.5 8.0 8.6 8.9 9.0 10.0 12.0"
|
||||
export TORCH_CUDA_ARCH_LIST="7.5 8.0 8.6 8.9 9.0 10.0 12.0+PTX"
|
||||
|
||||
bash tools/check_repo.sh
|
||||
|
||||
|
||||
@@ -15,7 +15,7 @@ jobs:
|
||||
actions: write
|
||||
runs-on: ubuntu-latest
|
||||
steps:
|
||||
- uses: actions/stale@eb5cf3af3ac0a1aa4c9c45633dd1ae542a27a899 # v10.3.0
|
||||
- uses: actions/stale@997185467fa4f803885201cee163a9f38240193d # v10.1.1
|
||||
with:
|
||||
# Increasing this value ensures that changes to this workflow
|
||||
# propagate to all issues and PRs in days rather than months
|
||||
|
||||
+1
-7
@@ -15,9 +15,6 @@ vllm/third_party/flashmla/flash_mla_interface.py
|
||||
# DeepGEMM vendored package built from source
|
||||
vllm/third_party/deep_gemm/
|
||||
|
||||
# fmha_sm100 vendored package built from source
|
||||
vllm/third_party/fmha_sm100/
|
||||
|
||||
# triton jit
|
||||
.triton
|
||||
|
||||
@@ -29,9 +26,6 @@ __pycache__/
|
||||
# C extensions
|
||||
*.so
|
||||
|
||||
# Rust binaries
|
||||
vllm/vllm-rs
|
||||
|
||||
# Distribution / packaging
|
||||
.Python
|
||||
build/
|
||||
@@ -236,7 +230,7 @@ actionlint
|
||||
shellcheck*/
|
||||
|
||||
# Ignore moe/marlin_moe gen code
|
||||
csrc/libtorch_stable/moe/marlin_moe_wna16/kernel_*
|
||||
csrc/moe/marlin_moe_wna16/kernel_*
|
||||
|
||||
# Ignore ep_kernels_workspace folder
|
||||
ep_kernels_workspace/
|
||||
|
||||
+14
-40
@@ -21,7 +21,7 @@ repos:
|
||||
rev: v21.1.2
|
||||
hooks:
|
||||
- id: clang-format
|
||||
exclude: 'csrc/libtorch_stable/moe/topk_softmax_kernels.cu|vllm/third_party/.*'
|
||||
exclude: 'csrc/(moe/topk_softmax_kernels.cu|quantization/gguf/(ggml-common.h|dequantize.cuh|vecdotq.cuh|mmq.cuh|mmvq.cuh))|vllm/third_party/.*'
|
||||
types_or: [c++, cuda]
|
||||
args: [--style=file, --verbose]
|
||||
- repo: https://github.com/DavidAnson/markdownlint-cli2
|
||||
@@ -148,27 +148,33 @@ repos:
|
||||
language: python
|
||||
entry: python tools/pre_commit/generate_nightly_torch_test.py
|
||||
files: ^requirements/test/cuda\.(in|txt)$
|
||||
- id: mypy-3.10 # TODO: Use https://github.com/pre-commit/mirrors-mypy when mypy setup is less awkward
|
||||
name: Run mypy for Python 3.10
|
||||
entry: python tools/pre_commit/mypy.py "3.10"
|
||||
- id: mypy-local
|
||||
name: Run mypy locally for lowest supported Python version
|
||||
entry: python tools/pre_commit/mypy.py 0 "3.10"
|
||||
stages: [pre-commit] # Don't run in CI
|
||||
<<: &mypy_common
|
||||
language: python
|
||||
types_or: [python, pyi]
|
||||
require_serial: true
|
||||
additional_dependencies: ["mypy==1.20.2", regex, types-cachetools, types-setuptools, types-PyYAML, types-requests, types-torch, pydantic]
|
||||
additional_dependencies: ["mypy[faster-cache]==1.19.1", regex, types-cachetools, types-setuptools, types-PyYAML, types-requests, types-torch, pydantic]
|
||||
- id: mypy-3.10 # TODO: Use https://github.com/pre-commit/mirrors-mypy when mypy setup is less awkward
|
||||
name: Run mypy for Python 3.10
|
||||
entry: python tools/pre_commit/mypy.py 1 "3.10"
|
||||
<<: *mypy_common
|
||||
stages: [manual] # Only run in CI
|
||||
- id: mypy-3.11 # TODO: Use https://github.com/pre-commit/mirrors-mypy when mypy setup is less awkward
|
||||
name: Run mypy for Python 3.11
|
||||
entry: python tools/pre_commit/mypy.py "3.11"
|
||||
entry: python tools/pre_commit/mypy.py 1 "3.11"
|
||||
<<: *mypy_common
|
||||
stages: [manual] # Only run in CI
|
||||
- id: mypy-3.12 # TODO: Use https://github.com/pre-commit/mirrors-mypy when mypy setup is less awkward
|
||||
name: Run mypy for Python 3.12
|
||||
entry: python tools/pre_commit/mypy.py "3.12"
|
||||
entry: python tools/pre_commit/mypy.py 1 "3.12"
|
||||
<<: *mypy_common
|
||||
stages: [manual] # Only run in CI
|
||||
- id: mypy-3.13 # TODO: Use https://github.com/pre-commit/mirrors-mypy when mypy setup is less awkward
|
||||
name: Run mypy for Python 3.13
|
||||
entry: python tools/pre_commit/mypy.py "3.13"
|
||||
entry: python tools/pre_commit/mypy.py 1 "3.13"
|
||||
<<: *mypy_common
|
||||
stages: [manual] # Only run in CI
|
||||
- id: shellcheck
|
||||
@@ -216,12 +222,6 @@ repos:
|
||||
name: Update Dockerfile dependency graph
|
||||
entry: tools/pre_commit/update-dockerfile-graph.sh
|
||||
language: script
|
||||
- id: test-nonroot-entrypoint
|
||||
name: Test non-root entrypoint wrapper
|
||||
entry: bash docker/entrypoints/test_vllm_nonroot_entrypoint.sh
|
||||
language: system
|
||||
pass_filenames: false
|
||||
files: ^docker/entrypoints/(vllm-nonroot-entrypoint|test_vllm_nonroot_entrypoint)\.sh$
|
||||
- id: check-forbidden-imports
|
||||
name: Check for forbidden imports
|
||||
entry: python tools/pre_commit/check_forbidden_imports.py
|
||||
@@ -256,32 +256,6 @@ repos:
|
||||
entry: python tools/pre_commit/check_boolean_context_manager.py
|
||||
language: python
|
||||
types: [python]
|
||||
# Rust hooks. These shell out to `cargo`; tools/pre_commit/rust-check.sh
|
||||
# skips with a warning when cargo is not installed.
|
||||
- id: rust-cargo-autoinherit
|
||||
name: Rust - Normalize Cargo manifests with autoinherit
|
||||
entry: tools/pre_commit/rust-check.sh autoinherit --prefer-simple-dotted
|
||||
language: script
|
||||
pass_filenames: false
|
||||
require_serial: true
|
||||
stages: [pre-commit] # Only run locally as Buildkite will cover this
|
||||
files: ^rust/(Cargo\.toml|src/.*/Cargo\.toml)$
|
||||
- id: rust-cargo-sort
|
||||
name: Rust - Sort Cargo manifest sections
|
||||
entry: tools/pre_commit/rust-check.sh sort --workspace
|
||||
language: script
|
||||
pass_filenames: false
|
||||
require_serial: true
|
||||
stages: [pre-commit] # Only run locally as Buildkite will cover this
|
||||
files: ^rust/(Cargo\.toml|src/.*/Cargo\.toml)$
|
||||
- id: rust-cargo-fmt
|
||||
name: Rust - Format code
|
||||
entry: tools/pre_commit/rust-check.sh fmt
|
||||
language: script
|
||||
pass_filenames: false
|
||||
require_serial: true
|
||||
stages: [pre-commit] # Only run locally as Buildkite will cover this
|
||||
files: ^rust/.*(\.rs|Cargo\.toml|rustfmt\.toml)$
|
||||
# Keep `suggestion` last
|
||||
- id: suggestion
|
||||
name: Suggestion
|
||||
|
||||
+1
-1
@@ -9,8 +9,8 @@ build:
|
||||
python: "3.12"
|
||||
jobs:
|
||||
post_checkout:
|
||||
- git fetch origin main --unshallow --no-tags --filter=blob:none || true
|
||||
- bash docs/pre_run_check.sh
|
||||
- git fetch origin main --unshallow --no-tags --filter=blob:none || true
|
||||
pre_create_environment:
|
||||
- pip install uv
|
||||
create_environment:
|
||||
|
||||
@@ -98,31 +98,7 @@ pre-commit run --all-files
|
||||
pre-commit run ruff-check --all-files
|
||||
|
||||
# Run mypy as it is in CI:
|
||||
pre-commit run mypy-3.12 --all-files --hook-stage manual
|
||||
```
|
||||
|
||||
The line length limit for Python code is 88 characters. If you are not sure, use pre-commit to check.
|
||||
|
||||
Use [Google-style docstrings](https://google.github.io/styleguide/pyguide.html#38-comments-and-docstrings) (`Args:`/`Returns:`/`Raises:` sections), not reStructuredText/Sphinx fields (`:param:`, `:return:`, `:rtype:`).
|
||||
|
||||
### Coding style guidelines
|
||||
|
||||
Follow these rules for all code changes in this repository:
|
||||
|
||||
- Try to match existing code style.
|
||||
- Code should be self-documenting and self-explanatory.
|
||||
- Keep comments and docstrings minimal and concise.
|
||||
- Assume the reader is familiar with vLLM.
|
||||
|
||||
### Diagnosing CI failures
|
||||
|
||||
Buildkite logs are public; no login needed. Details: [docs/contributing/ci/failures.md](docs/contributing/ci/failures.md).
|
||||
|
||||
```bash
|
||||
# All failed-job logs for a PR's latest build (current branch's PR if omitted):
|
||||
.buildkite/scripts/ci-fetch-log.sh --pr <PR>
|
||||
# Any Buildkite build or job URL also works:
|
||||
.buildkite/scripts/ci-fetch-log.sh "<buildkite_url>"
|
||||
pre-commit run mypy-3.10 --all-files --hook-stage manual
|
||||
```
|
||||
|
||||
### Commit messages
|
||||
|
||||
+397
-520
File diff suppressed because it is too large
Load Diff
@@ -4,7 +4,6 @@ include requirements/cuda.txt
|
||||
include requirements/rocm.txt
|
||||
include requirements/cpu.txt
|
||||
include CMakeLists.txt
|
||||
include tools/build_rust.py
|
||||
|
||||
recursive-include cmake *
|
||||
recursive-include csrc *
|
||||
|
||||
@@ -34,15 +34,6 @@ Vulnerabilities that cause denial of service or partial disruption, but do not a
|
||||
|
||||
Minor issues such as informational disclosures, logging errors, non-exploitable flaws, or weaknesses that require local or high-privilege access and offer negligible impact. Examples include side channel attacks or hash collisions. These issues often have CVSS scores less than 4.0
|
||||
|
||||
## Fix disclosure policy
|
||||
|
||||
When a security report is accepted, the fix process depends on the severity:
|
||||
|
||||
* **CRITICAL and HIGH severity**: Fixes are developed in a private security fork and coordinated with the prenotification group before public disclosure.
|
||||
* **MODERATE and LOW severity**: Fixes are developed and submitted as public pull requests. These issues do not require embargo since they do not enable arbitrary code execution or significant data breach, and public visibility accelerates community review and adoption of the fix.
|
||||
|
||||
The vulnerability management team reserves the right to adjust the disclosure approach on a case-by-case basis, taking into account factors such as active exploitation, unusual attack surface, or coordination requirements with downstream vendors.
|
||||
|
||||
## Prenotification policy
|
||||
|
||||
For certain security issues of CRITICAL, HIGH, or MODERATE severity level, we may prenotify certain organizations or vendors that ship vLLM. The purpose of this prenotification is to allow for a coordinated release of fixes for severe issues.
|
||||
|
||||
@@ -108,6 +108,7 @@ python benchmark.py \
|
||||
--backends flash triton flashinfer \
|
||||
--batch-specs "q2k" "8q1s1k" "2q2k_32q1s1k" \
|
||||
--num-layers 10 \
|
||||
--repeats 5 \
|
||||
--output-csv results.csv
|
||||
```
|
||||
|
||||
@@ -163,17 +164,14 @@ python benchmark.py \
|
||||
# Model configuration
|
||||
--num-layers N # Number of layers
|
||||
--head-dim N # Head dimension
|
||||
--v-head-dim N # Value head dimension (defaults to --head-dim)
|
||||
--num-q-heads N # Query heads
|
||||
--num-kv-heads N # KV heads
|
||||
--block-size N # Block size
|
||||
--kv-lora-rank N # MLA KV LoRA rank
|
||||
--qk-nope-head-dim N # MLA non-RoPE QK head dim
|
||||
--qk-rope-head-dim N # MLA RoPE QK head dim
|
||||
|
||||
# Benchmark settings
|
||||
--device DEVICE # Device (default: cuda:0)
|
||||
--warmup-ms N # Warmup window in ms for triton do_bench
|
||||
--repeats N # Repetitions
|
||||
--warmup-iters N # Warmup iterations
|
||||
--profile-memory # Profile memory usage
|
||||
|
||||
# Parameter sweeps
|
||||
@@ -213,6 +211,8 @@ config = BenchmarkConfig(
|
||||
num_kv_heads=1,
|
||||
block_size=128,
|
||||
device="cuda:0",
|
||||
repeats=5,
|
||||
warmup_iters=3,
|
||||
)
|
||||
|
||||
# CUTLASS MLA with specific num_kv_splits
|
||||
@@ -253,10 +253,14 @@ formatter.save_json(results, "output.json")
|
||||
|
||||
## Tips
|
||||
|
||||
**1. Save results** - Always use `--output-csv` or `--output-json`
|
||||
**1. Warmup matters** - Use `--warmup-iters 10` for stable results
|
||||
|
||||
**2. Test incrementally** - Start with `--num-layers 1`
|
||||
**2. Multiple repeats** - Use `--repeats 20` for low variance
|
||||
|
||||
**3. Extended grammar** - Leverage spec decode, chunked prefill patterns
|
||||
**3. Save results** - Always use `--output-csv` or `--output-json`
|
||||
|
||||
**4. Parameter sweeps** - Use `--sweep-param` and `--sweep-values` to find optimal values
|
||||
**4. Test incrementally** - Start with `--num-layers 1 --repeats 1`
|
||||
|
||||
**5. Extended grammar** - Leverage spec decode, chunked prefill patterns
|
||||
|
||||
**6. Parameter sweeps** - Use `--sweep-param` and `--sweep-values` to find optimal values
|
||||
|
||||
@@ -26,9 +26,6 @@ Examples:
|
||||
"""
|
||||
|
||||
import argparse
|
||||
import os
|
||||
import shutil
|
||||
import subprocess
|
||||
import sys
|
||||
from dataclasses import replace
|
||||
from pathlib import Path
|
||||
@@ -86,15 +83,13 @@ def run_benchmark(config: BenchmarkConfig, **kwargs) -> BenchmarkResult:
|
||||
else:
|
||||
return run_standard_attention_benchmark(config)
|
||||
except Exception as e:
|
||||
error_msg = str(e) or repr(e)
|
||||
return BenchmarkResult(
|
||||
config=config,
|
||||
mean_time=float("inf"),
|
||||
median_time=float("inf"),
|
||||
std_time=0,
|
||||
min_time=float("inf"),
|
||||
max_time=float("inf"),
|
||||
error=error_msg,
|
||||
error=str(e),
|
||||
)
|
||||
|
||||
|
||||
@@ -120,12 +115,9 @@ def run_model_parameter_sweep(
|
||||
"""
|
||||
all_results = []
|
||||
|
||||
sweep_desc = (
|
||||
f"{sweep.param_name} = {sweep.values}"
|
||||
if sweep.param_name
|
||||
else f"{len(sweep.values)} configurations"
|
||||
console.print(
|
||||
f"[yellow]Model sweep mode: testing {sweep.param_name} = {sweep.values}[/]"
|
||||
)
|
||||
console.print(f"[yellow]Model sweep mode: testing {sweep_desc}[/]")
|
||||
|
||||
total = len(backends) * len(batch_specs) * len(sweep.values)
|
||||
|
||||
@@ -133,9 +125,9 @@ def run_model_parameter_sweep(
|
||||
for backend in backends:
|
||||
for spec in batch_specs:
|
||||
for value in sweep.values:
|
||||
# Create config with modified model parameter(s)
|
||||
# Create config with modified model parameter
|
||||
config_args = base_config_args.copy()
|
||||
sweep.apply(config_args, value)
|
||||
config_args[sweep.param_name] = value
|
||||
|
||||
# Create config with original backend for running
|
||||
clean_config = BenchmarkConfig(
|
||||
@@ -152,21 +144,13 @@ def run_model_parameter_sweep(
|
||||
all_results.append(result)
|
||||
|
||||
if not result.success:
|
||||
err_label = (
|
||||
f"{sweep.param_name}={value}"
|
||||
if sweep.param_name
|
||||
else f"{value}"
|
||||
)
|
||||
console.print(
|
||||
f"[red]Error {backend} {spec} {err_label}"
|
||||
f": {result.error}[/]"
|
||||
f"[red]Error {backend} {spec} {sweep.param_name}="
|
||||
f"{value}: {result.error}[/]"
|
||||
)
|
||||
|
||||
pbar.update(1)
|
||||
|
||||
if base_config_args.get("ncu_profile"):
|
||||
return all_results
|
||||
|
||||
# Display sweep results - create separate table for each parameter value
|
||||
console.print("\n[bold green]Model Parameter Sweep Results:[/]")
|
||||
formatter = ResultsFormatter(console)
|
||||
@@ -200,10 +184,7 @@ def run_model_parameter_sweep(
|
||||
)
|
||||
|
||||
for param_value in sorted_param_values:
|
||||
label = (
|
||||
f"{sweep.param_name} = {param_value}" if sweep.param_name else param_value
|
||||
)
|
||||
console.print(f"\n[bold cyan]{label}[/]")
|
||||
console.print(f"\n[bold cyan]{sweep.param_name} = {param_value}[/]")
|
||||
param_results = by_param_value[param_value]
|
||||
|
||||
# Create modified results with original backend names
|
||||
@@ -219,9 +200,8 @@ def run_model_parameter_sweep(
|
||||
formatter.print_table(modified_results, backends, compare_to_fastest=True)
|
||||
|
||||
# Show optimal backend for each (param_value, batch_spec) combination
|
||||
sweep_name = sweep.param_name or "config"
|
||||
console.print(
|
||||
f"\n[bold cyan]Optimal backend for each ({sweep_name}, batch_spec):[/]"
|
||||
f"\n[bold cyan]Optimal backend for each ({sweep.param_name}, batch_spec):[/]"
|
||||
)
|
||||
|
||||
# Group by (param_value, batch_spec)
|
||||
@@ -256,10 +236,7 @@ def run_model_parameter_sweep(
|
||||
for param_value, spec in sorted_keys:
|
||||
# Print header when param value changes
|
||||
if param_value != current_param_value:
|
||||
header = (
|
||||
f"{sweep.param_name}={param_value}" if sweep.param_name else param_value
|
||||
)
|
||||
console.print(f"\n [bold]{header}:[/]")
|
||||
console.print(f"\n [bold]{sweep.param_name}={param_value}:[/]")
|
||||
current_param_value = param_value
|
||||
|
||||
results = by_param_and_spec[(param_value, spec)]
|
||||
@@ -345,9 +322,6 @@ def run_parameter_sweep(
|
||||
|
||||
pbar.update(1)
|
||||
|
||||
if base_config_args.get("ncu_profile"):
|
||||
return all_results
|
||||
|
||||
# Display sweep results
|
||||
console.print("\n[bold green]Sweep Results:[/]")
|
||||
backend_labels = [sweep.get_label(b, v) for b in backends for v in sweep_values]
|
||||
@@ -500,35 +474,11 @@ def main():
|
||||
parser.add_argument("--num-q-heads", type=int, default=32, help="Query heads")
|
||||
parser.add_argument("--num-kv-heads", type=int, default=8, help="KV heads")
|
||||
parser.add_argument("--block-size", type=int, default=16, help="Block size")
|
||||
parser.add_argument(
|
||||
"--v-head-dim",
|
||||
type=int,
|
||||
default=None,
|
||||
help="Value head dimension (defaults to --head-dim if unset)",
|
||||
)
|
||||
|
||||
# MLA-specific model dimensions
|
||||
parser.add_argument(
|
||||
"--kv-lora-rank", type=int, default=None, help="MLA KV LoRA rank"
|
||||
)
|
||||
parser.add_argument(
|
||||
"--qk-nope-head-dim", type=int, default=None, help="MLA non-RoPE QK head dim"
|
||||
)
|
||||
parser.add_argument(
|
||||
"--qk-rope-head-dim", type=int, default=None, help="MLA RoPE QK head dim"
|
||||
)
|
||||
|
||||
# Benchmark settings
|
||||
parser.add_argument("--device", default="cuda:0", help="Device")
|
||||
parser.add_argument(
|
||||
"--warmup-ms",
|
||||
type=int,
|
||||
default=None,
|
||||
help=(
|
||||
"Warmup window in ms for triton's do_bench (default: triton's own). "
|
||||
"Has no effect with CUDA graphs; pass --no-cuda-graphs to use it."
|
||||
),
|
||||
)
|
||||
parser.add_argument("--repeats", type=int, default=1, help="Repetitions")
|
||||
parser.add_argument("--warmup-iters", type=int, default=3, help="Warmup iterations")
|
||||
parser.add_argument("--profile-memory", action="store_true", help="Profile memory")
|
||||
parser.add_argument(
|
||||
"--kv-cache-dtype",
|
||||
@@ -541,33 +491,10 @@ def main():
|
||||
action=argparse.BooleanOptionalAction,
|
||||
default=True,
|
||||
help=(
|
||||
"Use triton do_bench_cudagraph (True) or do_bench (False) "
|
||||
"for timing. CUDA graphs eliminate CPU launch overhead "
|
||||
"(default: True)"
|
||||
"Launch kernels with CUDA graphs to eliminate CPU overhead"
|
||||
"in measurements (default: True)"
|
||||
),
|
||||
)
|
||||
parser.add_argument(
|
||||
"--num-splits",
|
||||
type=int,
|
||||
default=None,
|
||||
help="FlashAttention split-K factor (0=auto heuristic, 1=disabled, >1=force N)",
|
||||
)
|
||||
parser.add_argument(
|
||||
"--ncu-profile",
|
||||
action="store_true",
|
||||
default=False,
|
||||
help=(
|
||||
"Enable Nsight Compute profiling mode. Automatically wraps the "
|
||||
"script with ncu, capturing a profile with source correlation. "
|
||||
"Use --ncu-output to set the output file name."
|
||||
),
|
||||
)
|
||||
parser.add_argument(
|
||||
"--ncu-output",
|
||||
type=str,
|
||||
default="profile",
|
||||
help="Output file name for ncu profile (default: 'profile').",
|
||||
)
|
||||
|
||||
# Parameter sweep (use YAML config for advanced sweeps)
|
||||
parser.add_argument(
|
||||
@@ -649,28 +576,23 @@ def main():
|
||||
model = yaml_config["model"]
|
||||
args.num_layers = model.get("num_layers", args.num_layers)
|
||||
args.head_dim = model.get("head_dim", args.head_dim)
|
||||
args.v_head_dim = model.get("v_head_dim", args.v_head_dim)
|
||||
args.num_q_heads = model.get("num_q_heads", args.num_q_heads)
|
||||
args.num_kv_heads = model.get("num_kv_heads", args.num_kv_heads)
|
||||
args.block_size = model.get("block_size", args.block_size)
|
||||
# MLA-specific dimensions
|
||||
args.kv_lora_rank = model.get("kv_lora_rank", args.kv_lora_rank)
|
||||
args.qk_nope_head_dim = model.get("qk_nope_head_dim", args.qk_nope_head_dim)
|
||||
args.qk_rope_head_dim = model.get("qk_rope_head_dim", args.qk_rope_head_dim)
|
||||
|
||||
# Benchmark settings (top-level keys)
|
||||
if "device" in yaml_config:
|
||||
args.device = yaml_config["device"]
|
||||
if "warmup_ms" in yaml_config:
|
||||
args.warmup_ms = yaml_config["warmup_ms"]
|
||||
if "repeats" in yaml_config:
|
||||
args.repeats = yaml_config["repeats"]
|
||||
if "warmup_iters" in yaml_config:
|
||||
args.warmup_iters = yaml_config["warmup_iters"]
|
||||
if "profile_memory" in yaml_config:
|
||||
args.profile_memory = yaml_config["profile_memory"]
|
||||
if "kv_cache_dtype" in yaml_config:
|
||||
args.kv_cache_dtype = yaml_config["kv_cache_dtype"]
|
||||
if "cuda_graphs" in yaml_config:
|
||||
args.cuda_graphs = yaml_config["cuda_graphs"]
|
||||
if "ncu_profile" in yaml_config:
|
||||
args.ncu_profile = yaml_config["ncu_profile"]
|
||||
|
||||
# Parameter sweep configuration
|
||||
if "parameter_sweep" in yaml_config:
|
||||
@@ -690,7 +612,7 @@ def main():
|
||||
if "model_parameter_sweep" in yaml_config:
|
||||
sweep_config = yaml_config["model_parameter_sweep"]
|
||||
args.model_parameter_sweep = ModelParameterSweep(
|
||||
param_name=sweep_config.get("param_name"),
|
||||
param_name=sweep_config["param_name"],
|
||||
values=sweep_config["values"],
|
||||
label_format=sweep_config.get(
|
||||
"label_format", "{backend}_{param_name}_{value}"
|
||||
@@ -709,32 +631,6 @@ def main():
|
||||
|
||||
console.print()
|
||||
|
||||
# Re-exec under ncu if --ncu-profile and not already inside ncu. This runs
|
||||
# after YAML processing so ncu_profile set via config file is honored.
|
||||
if args.ncu_profile and "_NCU_INNER" not in os.environ:
|
||||
ncu = shutil.which("ncu")
|
||||
if ncu is None:
|
||||
print("Error: 'ncu' not found in PATH", file=sys.stderr)
|
||||
sys.exit(1)
|
||||
cmd = [
|
||||
ncu,
|
||||
"--profile-from-start",
|
||||
"off",
|
||||
"--set",
|
||||
"full",
|
||||
"--import-source",
|
||||
"yes",
|
||||
"-o",
|
||||
args.ncu_output,
|
||||
sys.executable,
|
||||
*sys.argv,
|
||||
]
|
||||
env = os.environ.copy()
|
||||
env["CUTE_DSL_LINEINFO"] = "1"
|
||||
env["_NCU_INNER"] = "1"
|
||||
print(f"Launching: {' '.join(cmd)}")
|
||||
sys.exit(subprocess.call(cmd, env=env))
|
||||
|
||||
# Handle CLI-based parameter sweep (if not from YAML)
|
||||
if (
|
||||
(not hasattr(args, "parameter_sweep") or args.parameter_sweep is None)
|
||||
@@ -759,18 +655,6 @@ def main():
|
||||
console.print(f"Batch specs: {', '.join(args.batch_specs)}")
|
||||
console.print(f"KV cache dtype: {args.kv_cache_dtype}")
|
||||
console.print(f"CUDA graphs: {args.cuda_graphs}")
|
||||
if args.warmup_ms is not None and args.cuda_graphs:
|
||||
console.print(
|
||||
"[yellow]Warning: --warmup-ms is ignored with CUDA graphs "
|
||||
"(do_bench_cudagraph warms up internally). Pass --no-cuda-graphs "
|
||||
"to use it.[/]"
|
||||
)
|
||||
if args.num_splits == 0 and args.cuda_graphs:
|
||||
console.print(
|
||||
"[yellow]Warning: --num-splits 0 (FA3 heuristic) is not CUDA-graph "
|
||||
"compatible and may fail or fall back. Pass --no-cuda-graphs or use "
|
||||
"--num-splits >=1.[/]"
|
||||
)
|
||||
console.print()
|
||||
|
||||
init_workspace_manager(args.device)
|
||||
@@ -778,15 +662,6 @@ def main():
|
||||
# Run benchmarks
|
||||
all_results = []
|
||||
|
||||
# Under ncu profiling the kernels run only to be captured by the profiler;
|
||||
# timings are placeholder zeros, so the result tables and saved metrics are
|
||||
# skipped. The Nsight Compute report (--ncu-output) holds the real data.
|
||||
if args.ncu_profile:
|
||||
console.print(
|
||||
"[dim]ncu profiling enabled: result tables and saved metrics are "
|
||||
"skipped (timings are placeholder zeros).[/]"
|
||||
)
|
||||
|
||||
# Handle special mode: decode_vs_prefill comparison
|
||||
if hasattr(args, "mode") and args.mode == "decode_vs_prefill":
|
||||
console.print("[yellow]Mode: Decode vs Prefill pipeline comparison[/]")
|
||||
@@ -833,11 +708,11 @@ def main():
|
||||
num_kv_heads=args.num_kv_heads,
|
||||
block_size=args.block_size,
|
||||
device=args.device,
|
||||
repeats=args.repeats,
|
||||
warmup_iters=args.warmup_iters,
|
||||
profile_memory=args.profile_memory,
|
||||
kv_cache_dtype=args.kv_cache_dtype,
|
||||
use_cuda_graphs=args.cuda_graphs,
|
||||
ncu_profile=args.ncu_profile,
|
||||
warmup_ms=args.warmup_ms,
|
||||
)
|
||||
|
||||
# Add decode pipeline config
|
||||
@@ -874,7 +749,6 @@ def main():
|
||||
result = BenchmarkResult(
|
||||
config=config,
|
||||
mean_time=timing["mean"],
|
||||
median_time=timing.get("median", timing["mean"]),
|
||||
std_time=timing["std"],
|
||||
min_time=timing["min"],
|
||||
max_time=timing["max"],
|
||||
@@ -896,7 +770,6 @@ def main():
|
||||
result = BenchmarkResult(
|
||||
config=config,
|
||||
mean_time=float("inf"),
|
||||
median_time=float("inf"),
|
||||
std_time=0,
|
||||
min_time=float("inf"),
|
||||
max_time=float("inf"),
|
||||
@@ -906,9 +779,6 @@ def main():
|
||||
|
||||
pbar.update(1)
|
||||
|
||||
if args.ncu_profile:
|
||||
return
|
||||
|
||||
# Display decode vs prefill results
|
||||
console.print("\n[bold green]Decode vs Prefill Results:[/]")
|
||||
|
||||
@@ -988,20 +858,15 @@ def main():
|
||||
base_config_args = {
|
||||
"num_layers": args.num_layers,
|
||||
"head_dim": args.head_dim,
|
||||
"v_head_dim": args.v_head_dim,
|
||||
"num_q_heads": args.num_q_heads,
|
||||
"num_kv_heads": args.num_kv_heads,
|
||||
"block_size": args.block_size,
|
||||
"device": args.device,
|
||||
"repeats": args.repeats,
|
||||
"warmup_iters": args.warmup_iters,
|
||||
"profile_memory": args.profile_memory,
|
||||
"kv_cache_dtype": args.kv_cache_dtype,
|
||||
"use_cuda_graphs": args.cuda_graphs,
|
||||
"ncu_profile": args.ncu_profile,
|
||||
"warmup_ms": args.warmup_ms,
|
||||
"num_splits": args.num_splits,
|
||||
"kv_lora_rank": args.kv_lora_rank,
|
||||
"qk_nope_head_dim": args.qk_nope_head_dim,
|
||||
"qk_rope_head_dim": args.qk_rope_head_dim,
|
||||
}
|
||||
all_results = run_model_parameter_sweep(
|
||||
backends,
|
||||
@@ -1017,17 +882,15 @@ def main():
|
||||
base_config_args = {
|
||||
"num_layers": args.num_layers,
|
||||
"head_dim": args.head_dim,
|
||||
"v_head_dim": args.v_head_dim,
|
||||
"num_q_heads": args.num_q_heads,
|
||||
"num_kv_heads": args.num_kv_heads,
|
||||
"block_size": args.block_size,
|
||||
"device": args.device,
|
||||
"repeats": args.repeats,
|
||||
"warmup_iters": args.warmup_iters,
|
||||
"profile_memory": args.profile_memory,
|
||||
"kv_cache_dtype": args.kv_cache_dtype,
|
||||
"use_cuda_graphs": args.cuda_graphs,
|
||||
"ncu_profile": args.ncu_profile,
|
||||
"warmup_ms": args.warmup_ms,
|
||||
"num_splits": args.num_splits,
|
||||
}
|
||||
all_results = run_parameter_sweep(
|
||||
backends, args.batch_specs, base_config_args, args.parameter_sweep, console
|
||||
@@ -1051,17 +914,15 @@ def main():
|
||||
batch_spec=spec,
|
||||
num_layers=args.num_layers,
|
||||
head_dim=args.head_dim,
|
||||
v_head_dim=getattr(args, "v_head_dim", None),
|
||||
num_q_heads=args.num_q_heads,
|
||||
num_kv_heads=args.num_kv_heads,
|
||||
block_size=args.block_size,
|
||||
device=args.device,
|
||||
repeats=args.repeats,
|
||||
warmup_iters=args.warmup_iters,
|
||||
profile_memory=args.profile_memory,
|
||||
kv_cache_dtype=args.kv_cache_dtype,
|
||||
use_cuda_graphs=args.cuda_graphs,
|
||||
ncu_profile=args.ncu_profile,
|
||||
warmup_ms=args.warmup_ms,
|
||||
num_splits=args.num_splits,
|
||||
)
|
||||
|
||||
result = run_benchmark(config)
|
||||
@@ -1074,10 +935,9 @@ def main():
|
||||
|
||||
pbar.update(1)
|
||||
|
||||
if not args.ncu_profile:
|
||||
console.print("\n[bold green]Results:[/]")
|
||||
formatter = ResultsFormatter(console)
|
||||
formatter.print_table(decode_results, backends)
|
||||
console.print("\n[bold green]Results:[/]")
|
||||
formatter = ResultsFormatter(console)
|
||||
formatter.print_table(decode_results, backends)
|
||||
|
||||
# Run prefill backend comparison
|
||||
if prefill_backends:
|
||||
@@ -1102,8 +962,9 @@ def main():
|
||||
num_kv_heads=args.num_kv_heads,
|
||||
block_size=args.block_size,
|
||||
device=args.device,
|
||||
repeats=args.repeats,
|
||||
warmup_iters=args.warmup_iters,
|
||||
profile_memory=args.profile_memory,
|
||||
warmup_ms=args.warmup_ms,
|
||||
prefill_backend=pb,
|
||||
)
|
||||
|
||||
@@ -1119,17 +980,16 @@ def main():
|
||||
|
||||
pbar.update(1)
|
||||
|
||||
if not args.ncu_profile:
|
||||
console.print("\n[bold green]Prefill Backend Results:[/]")
|
||||
formatter = ResultsFormatter(console)
|
||||
formatter.print_table(
|
||||
prefill_results, prefill_backends, compare_to_fastest=True
|
||||
)
|
||||
console.print("\n[bold green]Prefill Backend Results:[/]")
|
||||
formatter = ResultsFormatter(console)
|
||||
formatter.print_table(
|
||||
prefill_results, prefill_backends, compare_to_fastest=True
|
||||
)
|
||||
|
||||
all_results = decode_results + prefill_results
|
||||
|
||||
# Save results (skip ncu profiling runs: timings are placeholder zeros)
|
||||
if all_results and not args.ncu_profile:
|
||||
# Save results
|
||||
if all_results:
|
||||
formatter = ResultsFormatter(console)
|
||||
if args.output_csv:
|
||||
formatter.save_csv(all_results, args.output_csv)
|
||||
|
||||
@@ -15,8 +15,6 @@ from batch_spec import get_batch_type, parse_batch_spec
|
||||
from rich.console import Console
|
||||
from rich.table import Table
|
||||
|
||||
from vllm.triton_utils import triton
|
||||
|
||||
|
||||
def batch_spec_sort_key(spec: str) -> tuple[int, int, int]:
|
||||
"""
|
||||
@@ -36,30 +34,6 @@ def batch_spec_sort_key(spec: str) -> tuple[int, int, int]:
|
||||
return (0, 0, 0)
|
||||
|
||||
|
||||
def run_do_bench(
|
||||
benchmark_fn,
|
||||
use_cuda_graphs: bool,
|
||||
warmup_ms: int | None = None,
|
||||
) -> list[float]:
|
||||
kwargs: dict[str, Any] = {"return_mode": "all"}
|
||||
if use_cuda_graphs:
|
||||
result = triton.testing.do_bench_cudagraph(benchmark_fn, **kwargs)
|
||||
else:
|
||||
if warmup_ms is not None:
|
||||
kwargs["warmup"] = warmup_ms
|
||||
result = triton.testing.do_bench(benchmark_fn, **kwargs)
|
||||
return result
|
||||
|
||||
|
||||
def run_ncu_profile(benchmark_fn) -> None:
|
||||
benchmark_fn()
|
||||
torch.accelerator.synchronize()
|
||||
torch.cuda.cudart().cudaProfilerStart()
|
||||
benchmark_fn()
|
||||
torch.accelerator.synchronize()
|
||||
torch.cuda.cudart().cudaProfilerStop()
|
||||
|
||||
|
||||
# Mock classes for vLLM attention infrastructure
|
||||
|
||||
|
||||
@@ -208,37 +182,18 @@ class ParameterSweep:
|
||||
|
||||
@dataclass
|
||||
class ModelParameterSweep:
|
||||
"""Configuration for sweeping model configuration parameter(s).
|
||||
"""Configuration for sweeping a model configuration parameter."""
|
||||
|
||||
Supports two modes:
|
||||
- Single param: param_name="head_dim", values=[128, 256, 512]
|
||||
- Multi param: values=[{head_dim: 192, v_head_dim: 128}, {head_dim: 256}]
|
||||
When values are dicts, each dict's keys are applied as config overrides.
|
||||
"""
|
||||
|
||||
param_name: str | None = None
|
||||
values: list[Any] | None = None
|
||||
label_format: str = "{backend}_{param_name}_{value}"
|
||||
param_name: str # Name of the model config parameter to sweep (e.g., "num_q_heads")
|
||||
values: list[Any] # List of values to test
|
||||
label_format: str = "{backend}_{param_name}_{value}" # Result label template
|
||||
|
||||
def get_label(self, backend: str, value: Any) -> str:
|
||||
"""Generate a label for a specific parameter value."""
|
||||
if isinstance(value, dict):
|
||||
return self.label_format.format(
|
||||
backend=backend, param_name=self.param_name, value=value, **value
|
||||
)
|
||||
return self.label_format.format(
|
||||
backend=backend, param_name=self.param_name, value=value
|
||||
)
|
||||
|
||||
def apply(self, config_args: dict, value: Any) -> None:
|
||||
"""Apply a sweep value to config args."""
|
||||
if isinstance(value, dict):
|
||||
config_args.update(value)
|
||||
elif self.param_name is not None:
|
||||
config_args[self.param_name] = value
|
||||
else:
|
||||
raise ValueError("param_name must be set if sweep values are not dicts")
|
||||
|
||||
|
||||
@dataclass
|
||||
class BenchmarkConfig:
|
||||
@@ -253,10 +208,10 @@ class BenchmarkConfig:
|
||||
block_size: int
|
||||
device: str
|
||||
dtype: torch.dtype = torch.float16
|
||||
repeats: int = 1
|
||||
warmup_iters: int = 3
|
||||
profile_memory: bool = False
|
||||
use_cuda_graphs: bool = False
|
||||
ncu_profile: bool = False
|
||||
warmup_ms: int | None = None
|
||||
|
||||
# "auto" or "fp8"
|
||||
kv_cache_dtype: str = "auto"
|
||||
@@ -271,7 +226,6 @@ class BenchmarkConfig:
|
||||
# Backend-specific tuning
|
||||
num_kv_splits: int | None = None # CUTLASS MLA
|
||||
reorder_batch_threshold: int | None = None # FlashAttn MLA, FlashMLA
|
||||
num_splits: int | None = None # FlashAttention split-K (0=auto, 1=disabled)
|
||||
|
||||
|
||||
@dataclass
|
||||
@@ -280,7 +234,6 @@ class BenchmarkResult:
|
||||
|
||||
config: BenchmarkConfig
|
||||
mean_time: float # seconds
|
||||
median_time: float # seconds
|
||||
std_time: float # seconds
|
||||
min_time: float # seconds
|
||||
max_time: float # seconds
|
||||
@@ -299,7 +252,6 @@ class BenchmarkResult:
|
||||
return {
|
||||
"config": asdict(self.config),
|
||||
"mean_time": self.mean_time,
|
||||
"median_time": self.median_time,
|
||||
"std_time": self.std_time,
|
||||
"min_time": self.min_time,
|
||||
"max_time": self.max_time,
|
||||
|
||||
@@ -56,6 +56,8 @@ backends:
|
||||
- TOKENSPEED_MLA # Blackwell + R1 dims + FP8 KV (use --kv-cache-dtype fp8)
|
||||
|
||||
device: "cuda:0"
|
||||
repeats: 100
|
||||
warmup_iters: 10
|
||||
profile_memory: true
|
||||
|
||||
# Backend-specific tuning
|
||||
|
||||
@@ -51,6 +51,8 @@ backends:
|
||||
- FLASHMLA # Hopper only
|
||||
|
||||
device: "cuda:0"
|
||||
repeats: 5
|
||||
warmup_iters: 3
|
||||
profile_memory: true
|
||||
|
||||
# Analyze chunked prefill workspace size impact
|
||||
|
||||
@@ -124,3 +124,5 @@ prefill_backends:
|
||||
- tokenspeed
|
||||
|
||||
device: "cuda:0"
|
||||
repeats: 20
|
||||
warmup_iters: 5
|
||||
|
||||
@@ -53,4 +53,6 @@ backends:
|
||||
- FLASHINFER_MLA_SPARSE
|
||||
|
||||
device: "cuda:0"
|
||||
repeats: 100
|
||||
warmup_iters: 10
|
||||
profile_memory: true
|
||||
|
||||
@@ -57,4 +57,6 @@ backends:
|
||||
- FLASHINFER_MLA_SPARSE
|
||||
|
||||
device: "cuda:0"
|
||||
repeats: 10
|
||||
warmup_iters: 3
|
||||
profile_memory: true
|
||||
|
||||
@@ -63,6 +63,8 @@ model:
|
||||
|
||||
# Benchmark settings
|
||||
device: "cuda:0"
|
||||
repeats: 15 # More repeats for spec decode variance
|
||||
warmup_iters: 5
|
||||
profile_memory: false
|
||||
|
||||
# Output
|
||||
|
||||
@@ -49,6 +49,8 @@ backends:
|
||||
|
||||
# Benchmark settings
|
||||
device: "cuda:0"
|
||||
repeats: 10 # More repeats for statistical significance
|
||||
warmup_iters: 5
|
||||
profile_memory: false
|
||||
|
||||
# Test these threshold values for optimization
|
||||
|
||||
@@ -43,4 +43,6 @@ backends:
|
||||
- FLASHINFER
|
||||
|
||||
device: "cuda:0"
|
||||
repeats: 5
|
||||
warmup_iters: 3
|
||||
profile_memory: false
|
||||
|
||||
@@ -1,142 +0,0 @@
|
||||
# Standard attention decode benchmark configuration
|
||||
# Sweeps num_q_heads and num_kv_heads to isolate effects of:
|
||||
# 1. GQA ratio (fixed num_q_heads=32, vary num_kv_heads)
|
||||
# 2. Absolute head count (fixed 4:1 ratio, vary scale)
|
||||
|
||||
model:
|
||||
num_layers: 32
|
||||
num_q_heads: 32 # Base value, overridden by sweep
|
||||
num_kv_heads: 8 # Base value, overridden by sweep
|
||||
head_dim: 128
|
||||
block_size: 16
|
||||
|
||||
# Head count sweep: each entry overrides num_q_heads, num_kv_heads, and
|
||||
# head_dim where it differs from the base (128). Head counts are per-GPU
|
||||
# (i.e. after TP sharding).
|
||||
#
|
||||
# Group A — vary GQA ratio (fixed q=32, head_dim=128):
|
||||
# 32:32 (MHA), 32:8 (GQA 4:1), 32:4 (GQA 8:1), 32:1 (MQA)
|
||||
#
|
||||
# Groups B-E — real model configs at various TP degrees:
|
||||
# Model head_dim Full TP2 TP4 TP8
|
||||
# Llama 3 8B 128 32:8 16:4 8:2 4:1
|
||||
# Llama 3 70B 128 64:8 32:4 16:2 8:1
|
||||
# GPT-OSS 120B 64 64:8 32:4 16:2 8:1
|
||||
# Llama 3 405B 128 128:8 64:4 32:2 16:1
|
||||
model_parameter_sweep:
|
||||
values:
|
||||
# --- head_dim=128 (Llama 3 family) ---
|
||||
- { num_q_heads: 32, num_kv_heads: 32, head_dim: 128 } # MHA 1:1
|
||||
- { num_q_heads: 32, num_kv_heads: 1, head_dim: 128 } # MQA 32:1
|
||||
- { num_q_heads: 4, num_kv_heads: 1, head_dim: 128 } # Llama 3 8B TP8
|
||||
- { num_q_heads: 8, num_kv_heads: 2, head_dim: 128 } # Llama 3 8B TP4
|
||||
- { num_q_heads: 16, num_kv_heads: 4, head_dim: 128 } # Llama 3 8B TP2
|
||||
- { num_q_heads: 32, num_kv_heads: 8, head_dim: 128 } # Llama 3 8B TP1 / GQA 4:1
|
||||
- { num_q_heads: 8, num_kv_heads: 1, head_dim: 128 } # Llama 3 70B TP8
|
||||
- { num_q_heads: 16, num_kv_heads: 2, head_dim: 128 } # Llama 3 70B TP4
|
||||
- { num_q_heads: 32, num_kv_heads: 4, head_dim: 128 } # Llama 3 70B TP2 / GQA 8:1
|
||||
- { num_q_heads: 64, num_kv_heads: 8, head_dim: 128 } # Llama 3 70B TP1
|
||||
- { num_q_heads: 16, num_kv_heads: 1, head_dim: 128 } # Llama 3 405B TP8
|
||||
- { num_q_heads: 32, num_kv_heads: 2, head_dim: 128 } # Llama 3 405B TP4
|
||||
- { num_q_heads: 64, num_kv_heads: 4, head_dim: 128 } # Llama 3 405B TP2
|
||||
- { num_q_heads: 128, num_kv_heads: 8, head_dim: 128 } # Llama 3 405B TP1
|
||||
# --- head_dim=64 (GPT-OSS 120B) ---
|
||||
- { num_q_heads: 8, num_kv_heads: 1, head_dim: 64 } # GPT-OSS 120B TP8
|
||||
- { num_q_heads: 16, num_kv_heads: 2, head_dim: 64 } # GPT-OSS 120B TP4
|
||||
- { num_q_heads: 32, num_kv_heads: 4, head_dim: 64 } # GPT-OSS 120B TP2
|
||||
- { num_q_heads: 64, num_kv_heads: 8, head_dim: 64 } # GPT-OSS 120B TP1
|
||||
label_format: "{backend}_q{num_q_heads}kv{num_kv_heads}d{head_dim}"
|
||||
|
||||
batch_specs:
|
||||
# ---- batch_size x seq_len grid (decode: q_len=1) ----
|
||||
# Small grid for quick iteration. Uncomment for full sweep.
|
||||
|
||||
# Batch size 1
|
||||
- "q1s1k"
|
||||
- "q1s512"
|
||||
- "q1s2k"
|
||||
- "q1s4k"
|
||||
- "q1s8k"
|
||||
- "q1s16k"
|
||||
- "q1s32k"
|
||||
|
||||
# Batch size 2
|
||||
- "2q1s512"
|
||||
- "2q1s1k"
|
||||
- "2q1s2k"
|
||||
- "2q1s4k"
|
||||
- "2q1s8k"
|
||||
- "2q1s16k"
|
||||
- "2q1s32k"
|
||||
|
||||
# Batch size 4
|
||||
- "4q1s512"
|
||||
- "4q1s1k"
|
||||
- "4q1s2k"
|
||||
- "4q1s4k"
|
||||
- "4q1s8k"
|
||||
- "4q1s16k"
|
||||
- "4q1s32k"
|
||||
|
||||
# Batch size 8
|
||||
- "8q1s1k"
|
||||
- "8q1s512"
|
||||
- "8q1s2k"
|
||||
- "8q1s4k"
|
||||
- "8q1s8k"
|
||||
- "8q1s16k"
|
||||
- "8q1s32k"
|
||||
|
||||
# Batch size 16
|
||||
- "16q1s512"
|
||||
- "16q1s1k"
|
||||
- "16q1s2k"
|
||||
- "16q1s4k"
|
||||
- "16q1s8k"
|
||||
- "16q1s16k"
|
||||
- "16q1s32k"
|
||||
|
||||
# Batch size 32
|
||||
- "32q1s512"
|
||||
- "32q1s1k"
|
||||
- "32q1s2k"
|
||||
- "32q1s4k"
|
||||
- "32q1s8k"
|
||||
- "32q1s16k"
|
||||
- "32q1s32k"
|
||||
|
||||
# Batch size 64
|
||||
- "64q1s1k"
|
||||
- "64q1s512"
|
||||
- "64q1s2k"
|
||||
- "64q1s4k"
|
||||
- "64q1s8k"
|
||||
- "64q1s16k"
|
||||
- "64q1s32k"
|
||||
|
||||
# Batch size 128
|
||||
- "128q1s512"
|
||||
- "128q1s1k"
|
||||
- "128q1s2k"
|
||||
- "128q1s4k"
|
||||
- "128q1s8k"
|
||||
- "128q1s16k"
|
||||
- "128q1s32k"
|
||||
|
||||
# Batch size 256
|
||||
- "256q1s1k"
|
||||
- "256q1s512"
|
||||
- "256q1s2k"
|
||||
- "256q1s4k"
|
||||
- "256q1s8k"
|
||||
- "256q1s16k"
|
||||
- "256q1s32k"
|
||||
|
||||
# Available backends: FLASH_ATTN, TRITON_ATTN, FLASHINFER
|
||||
backends:
|
||||
- FLASH_ATTN
|
||||
- TRITON_ATTN
|
||||
- FLASHINFER
|
||||
|
||||
device: "cuda:0"
|
||||
profile_memory: false
|
||||
@@ -1,108 +0,0 @@
|
||||
# Standard attention prefill benchmark configuration
|
||||
# Sweeps num_q_heads and num_kv_heads to isolate effects of:
|
||||
# 1. GQA ratio (fixed num_q_heads=32, vary num_kv_heads)
|
||||
# 2. Absolute head count (fixed 4:1 ratio, vary scale)
|
||||
|
||||
model:
|
||||
num_layers: 32
|
||||
num_q_heads: 32 # Base value, overridden by sweep
|
||||
num_kv_heads: 8 # Base value, overridden by sweep
|
||||
head_dim: 128
|
||||
block_size: 16
|
||||
|
||||
# Head count sweep: each entry overrides num_q_heads, num_kv_heads, and
|
||||
# head_dim where it differs from the base (128). Head counts are per-GPU
|
||||
# (i.e. after TP sharding).
|
||||
#
|
||||
# Group A — vary GQA ratio (fixed q=32, head_dim=128):
|
||||
# 32:32 (MHA), 32:8 (GQA 4:1), 32:4 (GQA 8:1), 32:1 (MQA)
|
||||
#
|
||||
# Groups B-E — real model configs at various TP degrees:
|
||||
# Model head_dim Full TP2 TP4 TP8
|
||||
# Llama 3 8B 128 32:8 16:4 8:2 4:1
|
||||
# Llama 3 70B 128 64:8 32:4 16:2 8:1
|
||||
# GPT-OSS 120B 64 64:8 32:4 16:2 8:1
|
||||
# Llama 3 405B 128 128:8 64:4 32:2 16:1
|
||||
model_parameter_sweep:
|
||||
values:
|
||||
# --- head_dim=128 (Llama 3 family) ---
|
||||
- { num_q_heads: 32, num_kv_heads: 32, head_dim: 128 } # MHA 1:1
|
||||
- { num_q_heads: 32, num_kv_heads: 1, head_dim: 128 } # MQA 32:1
|
||||
- { num_q_heads: 4, num_kv_heads: 1, head_dim: 128 } # Llama 3 8B TP8
|
||||
- { num_q_heads: 8, num_kv_heads: 2, head_dim: 128 } # Llama 3 8B TP4
|
||||
- { num_q_heads: 16, num_kv_heads: 4, head_dim: 128 } # Llama 3 8B TP2
|
||||
- { num_q_heads: 32, num_kv_heads: 8, head_dim: 128 } # Llama 3 8B TP1 / GQA 4:1
|
||||
- { num_q_heads: 8, num_kv_heads: 1, head_dim: 128 } # Llama 3 70B TP8
|
||||
- { num_q_heads: 16, num_kv_heads: 2, head_dim: 128 } # Llama 3 70B TP4
|
||||
- { num_q_heads: 32, num_kv_heads: 4, head_dim: 128 } # Llama 3 70B TP2 / GQA 8:1
|
||||
- { num_q_heads: 64, num_kv_heads: 8, head_dim: 128 } # Llama 3 70B TP1
|
||||
- { num_q_heads: 16, num_kv_heads: 1, head_dim: 128 } # Llama 3 405B TP8
|
||||
- { num_q_heads: 32, num_kv_heads: 2, head_dim: 128 } # Llama 3 405B TP4
|
||||
- { num_q_heads: 64, num_kv_heads: 4, head_dim: 128 } # Llama 3 405B TP2
|
||||
- { num_q_heads: 128, num_kv_heads: 8, head_dim: 128 } # Llama 3 405B TP1
|
||||
# --- head_dim=64 (GPT-OSS 120B) ---
|
||||
- { num_q_heads: 8, num_kv_heads: 1, head_dim: 64 } # GPT-OSS 120B TP8
|
||||
- { num_q_heads: 16, num_kv_heads: 2, head_dim: 64 } # GPT-OSS 120B TP4
|
||||
- { num_q_heads: 32, num_kv_heads: 4, head_dim: 64 } # GPT-OSS 120B TP2
|
||||
- { num_q_heads: 64, num_kv_heads: 8, head_dim: 64 } # GPT-OSS 120B TP1
|
||||
label_format: "{backend}_q{num_q_heads}kv{num_kv_heads}d{head_dim}"
|
||||
|
||||
batch_specs:
|
||||
# ---- batch_size x prefill_len grid (prefill: q_len == seq_len) ----
|
||||
# Total tokens = batch_size * prefill_len, and prefill compute scales with
|
||||
# prefill_len^2, so the largest cells are expensive. Trim batch sizes or
|
||||
# lengths for quick iteration.
|
||||
|
||||
# Batch size 1
|
||||
- "q512"
|
||||
- "q1k"
|
||||
- "q2k"
|
||||
- "q4k"
|
||||
- "q8k"
|
||||
- "q16k"
|
||||
- "q32k"
|
||||
|
||||
# Batch size 2
|
||||
- "2q512"
|
||||
- "2q1k"
|
||||
- "2q2k"
|
||||
- "2q4k"
|
||||
- "2q8k"
|
||||
- "2q16k"
|
||||
- "2q32k"
|
||||
|
||||
# Batch size 4
|
||||
- "4q512"
|
||||
- "4q1k"
|
||||
- "4q2k"
|
||||
- "4q4k"
|
||||
- "4q8k"
|
||||
- "4q16k"
|
||||
- "4q32k"
|
||||
|
||||
# Batch size 8
|
||||
- "8q512"
|
||||
- "8q1k"
|
||||
- "8q2k"
|
||||
- "8q4k"
|
||||
- "8q8k"
|
||||
- "8q16k"
|
||||
- "8q32k"
|
||||
|
||||
# Batch size 16
|
||||
- "16q512"
|
||||
- "16q1k"
|
||||
- "16q2k"
|
||||
- "16q4k"
|
||||
- "16q8k"
|
||||
- "16q16k"
|
||||
- "16q32k"
|
||||
|
||||
# Available backends: FLASH_ATTN, TRITON_ATTN, FLASHINFER
|
||||
backends:
|
||||
- FLASH_ATTN
|
||||
- TRITON_ATTN
|
||||
- FLASHINFER
|
||||
|
||||
device: "cuda:0"
|
||||
profile_memory: false
|
||||
@@ -8,8 +8,6 @@ This module provides helpers for running MLA backends without
|
||||
needing full VllmConfig integration.
|
||||
"""
|
||||
|
||||
import statistics
|
||||
|
||||
import numpy as np
|
||||
import torch
|
||||
from batch_spec import parse_batch_spec
|
||||
@@ -19,8 +17,6 @@ from common import (
|
||||
MockIndexer,
|
||||
MockKVBProj,
|
||||
MockLayer,
|
||||
run_do_bench,
|
||||
run_ncu_profile,
|
||||
setup_mla_dims,
|
||||
)
|
||||
|
||||
@@ -824,7 +820,7 @@ def _run_single_benchmark(
|
||||
num_prefill, mla_dims, query_fmt, device, torch.bfloat16
|
||||
)
|
||||
|
||||
# Build forward function (runs a single decode/prefill pass)
|
||||
# Build forward function
|
||||
def forward_fn():
|
||||
results = []
|
||||
if has_decode:
|
||||
@@ -843,35 +839,44 @@ def _run_single_benchmark(
|
||||
)
|
||||
return results[0] if len(results) == 1 else tuple(results)
|
||||
|
||||
def benchmark_fn():
|
||||
for _ in range(config.num_layers):
|
||||
# Warmup
|
||||
for _ in range(config.warmup_iters):
|
||||
forward_fn()
|
||||
torch.accelerator.synchronize()
|
||||
|
||||
# Optionally capture a CUDA graph after warmup.
|
||||
# Graph replay eliminates CPU launch overhead so timings reflect pure
|
||||
# kernel time.
|
||||
if config.use_cuda_graphs:
|
||||
graph = torch.cuda.CUDAGraph()
|
||||
with torch.cuda.graph(graph):
|
||||
forward_fn()
|
||||
benchmark_fn = graph.replay
|
||||
else:
|
||||
benchmark_fn = forward_fn
|
||||
|
||||
if config.ncu_profile:
|
||||
run_ncu_profile(benchmark_fn)
|
||||
return BenchmarkResult(
|
||||
config=config,
|
||||
mean_time=0.0,
|
||||
median_time=0.0,
|
||||
std_time=0.0,
|
||||
min_time=0.0,
|
||||
max_time=0.0,
|
||||
throughput_tokens_per_sec=0.0,
|
||||
)
|
||||
# Benchmark
|
||||
times = []
|
||||
for _ in range(config.repeats):
|
||||
start = torch.cuda.Event(enable_timing=True)
|
||||
end = torch.cuda.Event(enable_timing=True)
|
||||
|
||||
all_ms = run_do_bench(benchmark_fn, config.use_cuda_graphs, config.warmup_ms)
|
||||
start.record()
|
||||
for _ in range(config.num_layers):
|
||||
benchmark_fn()
|
||||
end.record()
|
||||
|
||||
# Convert ms to seconds per layer
|
||||
times = [t / 1000.0 / config.num_layers for t in all_ms]
|
||||
mean_time = statistics.mean(times)
|
||||
torch.accelerator.synchronize()
|
||||
elapsed_ms = start.elapsed_time(end)
|
||||
times.append(elapsed_ms / 1000.0 / config.num_layers)
|
||||
|
||||
mean_time = float(np.mean(times))
|
||||
return BenchmarkResult(
|
||||
config=config,
|
||||
mean_time=mean_time,
|
||||
median_time=statistics.median(times),
|
||||
std_time=statistics.stdev(times) if len(times) > 1 else 0.0,
|
||||
min_time=min(times),
|
||||
max_time=max(times),
|
||||
std_time=float(np.std(times)),
|
||||
min_time=float(np.min(times)),
|
||||
max_time=float(np.max(times)),
|
||||
throughput_tokens_per_sec=total_q / mean_time if mean_time > 0 else 0,
|
||||
)
|
||||
|
||||
|
||||
@@ -9,20 +9,13 @@ This module provides helpers for running standard attention backends
|
||||
"""
|
||||
|
||||
import logging
|
||||
import statistics
|
||||
import types
|
||||
from contextlib import contextmanager
|
||||
|
||||
import numpy as np
|
||||
import torch
|
||||
from batch_spec import parse_batch_spec, reorder_for_flashinfer
|
||||
from common import (
|
||||
BenchmarkConfig,
|
||||
BenchmarkResult,
|
||||
MockLayer,
|
||||
get_attention_scale,
|
||||
run_do_bench,
|
||||
run_ncu_profile,
|
||||
)
|
||||
from common import BenchmarkConfig, BenchmarkResult, MockLayer, get_attention_scale
|
||||
|
||||
from vllm.config import (
|
||||
CacheConfig,
|
||||
@@ -215,13 +208,6 @@ def _create_backend_impl(
|
||||
|
||||
scale = get_attention_scale(config.head_dim)
|
||||
|
||||
# Set v_head_dim for diff-headdim backends. Always reset (defaulting to
|
||||
# head_dim) so a prior run's value doesn't leak into this one via the
|
||||
# backend's class-level state.
|
||||
if hasattr(backend_class, "set_head_size_v"):
|
||||
v_dim = config.v_head_dim if config.v_head_dim is not None else config.head_dim
|
||||
backend_class.set_head_size_v(v_dim)
|
||||
|
||||
impl = backend_class.get_impl_cls()(
|
||||
num_heads=config.num_q_heads,
|
||||
head_size=config.head_dim,
|
||||
@@ -314,7 +300,6 @@ def _create_input_tensors(
|
||||
from vllm.platforms import current_platform
|
||||
|
||||
q_dtype = current_platform.fp8_dtype()
|
||||
v_dim = config.v_head_dim if config.v_head_dim is not None else config.head_dim
|
||||
q_list = [
|
||||
torch.randn(
|
||||
total_q, config.num_q_heads, config.head_dim, device=device, dtype=dtype
|
||||
@@ -328,7 +313,9 @@ def _create_input_tensors(
|
||||
for _ in range(config.num_layers)
|
||||
]
|
||||
v_list = [
|
||||
torch.randn(total_q, config.num_kv_heads, v_dim, device=device, dtype=dtype)
|
||||
torch.randn(
|
||||
total_q, config.num_kv_heads, config.head_dim, device=device, dtype=dtype
|
||||
)
|
||||
for _ in range(config.num_layers)
|
||||
]
|
||||
return q_list, k_list, v_list
|
||||
@@ -402,17 +389,14 @@ def _run_single_benchmark(
|
||||
device: torch.device,
|
||||
dtype: torch.dtype,
|
||||
) -> tuple:
|
||||
"""Run single benchmark using triton's do_bench_cudagraph/do_bench.
|
||||
|
||||
Returns:
|
||||
(timing_stats, mem_stats) where timing_stats is a dict with
|
||||
mean/std/min/max in seconds per layer.
|
||||
"""
|
||||
"""Run single benchmark iteration with warmup and timing loop."""
|
||||
total_q = q_list[0].shape[0]
|
||||
v_dim = config.v_head_dim if config.v_head_dim is not None else config.head_dim
|
||||
out = torch.empty(total_q, config.num_q_heads, v_dim, device=device, dtype=dtype)
|
||||
out = torch.empty(
|
||||
total_q, config.num_q_heads, config.head_dim, device=device, dtype=dtype
|
||||
)
|
||||
|
||||
def benchmark_fn():
|
||||
# Warmup
|
||||
for _ in range(config.warmup_iters):
|
||||
for i in range(config.num_layers):
|
||||
impl.forward(
|
||||
layer,
|
||||
@@ -423,22 +407,52 @@ def _run_single_benchmark(
|
||||
attn_metadata,
|
||||
output=out,
|
||||
)
|
||||
torch.accelerator.synchronize()
|
||||
|
||||
if config.ncu_profile:
|
||||
run_ncu_profile(benchmark_fn)
|
||||
timing_stats = dict.fromkeys(("mean", "median", "std", "min", "max"), 0.0)
|
||||
# Optionally capture a CUDA graph after warmup.
|
||||
# Graph replay eliminates CPU launch overhead so timings reflect pure
|
||||
# kernel time.
|
||||
if config.use_cuda_graphs:
|
||||
graph = torch.cuda.CUDAGraph()
|
||||
with torch.cuda.graph(graph):
|
||||
for i in range(config.num_layers):
|
||||
impl.forward(
|
||||
layer,
|
||||
q_list[i],
|
||||
k_list[i],
|
||||
v_list[i],
|
||||
cache_list[i],
|
||||
attn_metadata,
|
||||
output=out,
|
||||
)
|
||||
benchmark_fn = graph.replay
|
||||
else:
|
||||
all_ms = run_do_bench(benchmark_fn, config.use_cuda_graphs, config.warmup_ms)
|
||||
|
||||
# Convert ms to seconds per layer
|
||||
times = [t / 1000.0 / config.num_layers for t in all_ms]
|
||||
timing_stats = {
|
||||
"mean": statistics.mean(times),
|
||||
"std": statistics.stdev(times) if len(times) > 1 else 0.0,
|
||||
"min": min(times),
|
||||
"max": max(times),
|
||||
"median": statistics.median(times),
|
||||
}
|
||||
def benchmark_fn():
|
||||
for i in range(config.num_layers):
|
||||
impl.forward(
|
||||
layer,
|
||||
q_list[i],
|
||||
k_list[i],
|
||||
v_list[i],
|
||||
cache_list[i],
|
||||
attn_metadata,
|
||||
output=out,
|
||||
)
|
||||
|
||||
# Benchmark
|
||||
times = []
|
||||
for _ in range(config.repeats):
|
||||
start = torch.cuda.Event(enable_timing=True)
|
||||
end = torch.cuda.Event(enable_timing=True)
|
||||
|
||||
start.record()
|
||||
benchmark_fn()
|
||||
end.record()
|
||||
|
||||
torch.accelerator.synchronize()
|
||||
elapsed_ms = start.elapsed_time(end)
|
||||
times.append(elapsed_ms / 1000.0 / config.num_layers) # seconds per layer
|
||||
|
||||
mem_stats = {}
|
||||
if config.profile_memory:
|
||||
@@ -447,7 +461,7 @@ def _run_single_benchmark(
|
||||
"reserved_mb": torch.accelerator.memory_reserved(device) / 1024**2,
|
||||
}
|
||||
|
||||
return timing_stats, mem_stats
|
||||
return times, mem_stats
|
||||
|
||||
|
||||
# ============================================================================
|
||||
@@ -527,12 +541,6 @@ def run_attention_benchmark(config: BenchmarkConfig) -> BenchmarkResult:
|
||||
common_attn_metadata=common_metadata,
|
||||
)
|
||||
|
||||
# Override num_splits for split-K testing (FlashAttention only)
|
||||
if config.num_splits is not None and hasattr(
|
||||
attn_metadata, "max_num_splits"
|
||||
):
|
||||
attn_metadata.max_num_splits = config.num_splits
|
||||
|
||||
# Only quantize queries when the impl supports it
|
||||
quantize_query = config.kv_cache_dtype.startswith("fp8") and getattr(
|
||||
impl, "supports_quant_query_input", False
|
||||
@@ -545,7 +553,7 @@ def run_attention_benchmark(config: BenchmarkConfig) -> BenchmarkResult:
|
||||
config, max_num_blocks, backend_class, device, dtype
|
||||
)
|
||||
|
||||
timing_stats, mem_stats = _run_single_benchmark(
|
||||
times, mem_stats = _run_single_benchmark(
|
||||
config,
|
||||
impl,
|
||||
layer,
|
||||
@@ -558,16 +566,15 @@ def run_attention_benchmark(config: BenchmarkConfig) -> BenchmarkResult:
|
||||
dtype,
|
||||
)
|
||||
|
||||
mean_time = timing_stats["mean"]
|
||||
mean_time = np.mean(times)
|
||||
throughput = total_q / mean_time if mean_time > 0 else 0
|
||||
|
||||
return BenchmarkResult(
|
||||
config=config,
|
||||
mean_time=mean_time,
|
||||
median_time=timing_stats["median"],
|
||||
std_time=timing_stats["std"],
|
||||
min_time=timing_stats["min"],
|
||||
max_time=timing_stats["max"],
|
||||
std_time=np.std(times),
|
||||
min_time=np.min(times),
|
||||
max_time=np.max(times),
|
||||
throughput_tokens_per_sec=throughput,
|
||||
memory_allocated_mb=mem_stats.get("allocated_mb"),
|
||||
memory_reserved_mb=mem_stats.get("reserved_mb"),
|
||||
|
||||
@@ -1,413 +0,0 @@
|
||||
# SPDX-License-Identifier: Apache-2.0
|
||||
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
"""
|
||||
Benchmark hidden state extraction throughput.
|
||||
|
||||
Measures two modes:
|
||||
1. Baseline: bulk inference with max_tokens=1, no extraction.
|
||||
2. Extract: async hidden state extraction via ExampleHiddenStatesConnector
|
||||
with N concurrent clients, each consuming hidden states as
|
||||
soon as their request finishes (overlapping I/O with generation).
|
||||
|
||||
Reports tokens/s and prompts/s for each mode.
|
||||
|
||||
Usage:
|
||||
python benchmarks/benchmark_hidden_state_extraction.py \
|
||||
--model Qwen/Qwen3-0.6B \
|
||||
--num-prompts 64 \
|
||||
--num-clients 8 \
|
||||
--prompt-len 8192 \
|
||||
--layers 1 2 3 4
|
||||
"""
|
||||
|
||||
import argparse
|
||||
import asyncio
|
||||
import time
|
||||
from concurrent.futures import ThreadPoolExecutor
|
||||
|
||||
import torch
|
||||
from transformers import AutoConfig
|
||||
|
||||
from vllm import LLM, SamplingParams
|
||||
from vllm.config.kv_transfer import KVTransferConfig
|
||||
from vllm.distributed.kv_transfer.kv_connector.v1 import (
|
||||
example_hidden_states_connector,
|
||||
)
|
||||
from vllm.engine.arg_utils import AsyncEngineArgs
|
||||
from vllm.sampling_params import RequestOutputKind
|
||||
from vllm.v1.engine.async_llm import AsyncLLM
|
||||
|
||||
|
||||
def _make_profiler_config(profile_dir: str) -> dict:
|
||||
"""Build a profiler_config dict for torch profiling."""
|
||||
return {
|
||||
"profiler": "torch",
|
||||
"torch_profiler_dir": profile_dir,
|
||||
"torch_profiler_with_stack": True,
|
||||
}
|
||||
|
||||
|
||||
def make_random_prompts(
|
||||
num_prompts: int, prompt_len: int, vocab_size: int, seed: int = 42
|
||||
) -> list[list[int]]:
|
||||
"""Generate lists of random token IDs."""
|
||||
# Set seed for reproducibility
|
||||
torch.manual_seed(seed)
|
||||
return [
|
||||
torch.randint(0, vocab_size, (prompt_len,)).tolist() for _ in range(num_prompts)
|
||||
]
|
||||
|
||||
|
||||
def consume_hidden_states(path: str) -> float:
|
||||
"""Load hidden states from disk and compute per-position mean.
|
||||
|
||||
Returns a single float: the grand mean of all hidden state values.
|
||||
This forces the benchmark to actually read and reduce the data.
|
||||
|
||||
Uses :func:`load_hidden_states` which acquires a shared flock,
|
||||
blocking (without polling) until the async writer releases its
|
||||
exclusive lock.
|
||||
"""
|
||||
obj = example_hidden_states_connector.load_hidden_states(path)
|
||||
hs = obj["hidden_states"]
|
||||
total = hs.mean().item()
|
||||
|
||||
example_hidden_states_connector.cleanup_hidden_states(path)
|
||||
|
||||
return total
|
||||
|
||||
|
||||
def run_baseline(
|
||||
model: str,
|
||||
prompts: list[list[int]],
|
||||
extra_args: dict,
|
||||
profile_dir: str | None = None,
|
||||
) -> dict:
|
||||
"""Baseline: bulk inference, no hidden state extraction."""
|
||||
if profile_dir:
|
||||
extra_args = {
|
||||
**extra_args,
|
||||
"profiler_config": _make_profiler_config(profile_dir),
|
||||
}
|
||||
llm = LLM(
|
||||
model=model,
|
||||
enable_prefix_caching=False,
|
||||
**extra_args,
|
||||
)
|
||||
sampling_params = SamplingParams(max_tokens=1)
|
||||
prompt_inputs = [{"prompt_token_ids": p} for p in prompts]
|
||||
|
||||
# Warmup
|
||||
llm.generate(prompt_inputs[:4], sampling_params, use_tqdm=False)
|
||||
|
||||
if profile_dir:
|
||||
llm.start_profile()
|
||||
|
||||
t0 = time.perf_counter()
|
||||
outputs = llm.generate(prompt_inputs, sampling_params, use_tqdm=True)
|
||||
elapsed = time.perf_counter() - t0
|
||||
|
||||
if profile_dir:
|
||||
llm.stop_profile()
|
||||
|
||||
total_prompt_tokens = sum(len(o.prompt_token_ids) for o in outputs)
|
||||
num_prompts = len(outputs)
|
||||
|
||||
del llm
|
||||
torch.accelerator.empty_cache()
|
||||
|
||||
return {
|
||||
"mode": "baseline",
|
||||
"elapsed_s": elapsed,
|
||||
"num_prompts": num_prompts,
|
||||
"total_prompt_tokens": total_prompt_tokens,
|
||||
"tokens_per_s": total_prompt_tokens / elapsed,
|
||||
"prompts_per_s": num_prompts / elapsed,
|
||||
}
|
||||
|
||||
|
||||
# ---- Async extraction benchmark ----
|
||||
|
||||
|
||||
async def _client_loop(
|
||||
engine: AsyncLLM,
|
||||
prompt_queue: asyncio.Queue,
|
||||
consume_pool: ThreadPoolExecutor,
|
||||
results: list[dict],
|
||||
client_id: int,
|
||||
):
|
||||
"""A single async client: pulls prompts, submits to engine, consumes
|
||||
hidden states as soon as each request finishes."""
|
||||
loop = asyncio.get_event_loop()
|
||||
while True:
|
||||
item = await prompt_queue.get()
|
||||
if item is None:
|
||||
prompt_queue.task_done()
|
||||
break
|
||||
idx, token_ids = item
|
||||
|
||||
request_id = f"req-{idx}"
|
||||
sampling_params = SamplingParams(
|
||||
max_tokens=1,
|
||||
output_kind=RequestOutputKind.FINAL_ONLY,
|
||||
)
|
||||
|
||||
final_output = None
|
||||
async for output in engine.generate(
|
||||
request_id=request_id,
|
||||
prompt={"prompt_token_ids": token_ids},
|
||||
sampling_params=sampling_params,
|
||||
):
|
||||
if output.finished:
|
||||
final_output = output
|
||||
|
||||
# Consume hidden states on a thread (disk I/O)
|
||||
path = final_output.kv_transfer_params["hidden_states_path"]
|
||||
mean_val = await loop.run_in_executor(consume_pool, consume_hidden_states, path)
|
||||
num_tokens = len(final_output.prompt_token_ids)
|
||||
|
||||
results.append(
|
||||
{
|
||||
"request_id": request_id,
|
||||
"num_prompt_tokens": num_tokens,
|
||||
"mean_hidden_value": mean_val,
|
||||
}
|
||||
)
|
||||
prompt_queue.task_done()
|
||||
|
||||
|
||||
async def _run_extraction_async(
|
||||
model: str,
|
||||
prompts: list[list[int]],
|
||||
num_clients: int,
|
||||
layers: list[int],
|
||||
tmpdir: str,
|
||||
extra_args: dict,
|
||||
profile_dir: str | None = None,
|
||||
) -> dict:
|
||||
if profile_dir:
|
||||
extra_args = {
|
||||
**extra_args,
|
||||
"profiler_config": _make_profiler_config(profile_dir),
|
||||
}
|
||||
engine_args = AsyncEngineArgs(
|
||||
model=model,
|
||||
enable_prefix_caching=False,
|
||||
max_num_batched_tokens=40960,
|
||||
max_model_len=40960,
|
||||
speculative_config={
|
||||
"method": "extract_hidden_states",
|
||||
"num_speculative_tokens": 1,
|
||||
"draft_model_config": {
|
||||
"hf_config": {
|
||||
"eagle_aux_hidden_state_layer_ids": layers,
|
||||
},
|
||||
},
|
||||
},
|
||||
kv_transfer_config=KVTransferConfig(
|
||||
kv_connector="ExampleHiddenStatesConnector",
|
||||
kv_role="kv_producer",
|
||||
kv_connector_extra_config={
|
||||
"shared_storage_path": tmpdir,
|
||||
},
|
||||
),
|
||||
**extra_args,
|
||||
)
|
||||
engine = AsyncLLM.from_engine_args(engine_args)
|
||||
|
||||
try:
|
||||
# Warmup: run a few prompts sequentially, cleaning up generated files
|
||||
for i in range(min(4, len(prompts))):
|
||||
sp = SamplingParams(max_tokens=1, output_kind=RequestOutputKind.FINAL_ONLY)
|
||||
final_output = None
|
||||
async for output in engine.generate(
|
||||
request_id=f"warmup-{i}",
|
||||
prompt={"prompt_token_ids": prompts[i]},
|
||||
sampling_params=sp,
|
||||
):
|
||||
if output.finished:
|
||||
final_output = output
|
||||
if final_output and final_output.kv_transfer_params:
|
||||
path = final_output.kv_transfer_params.get("hidden_states_path")
|
||||
if path:
|
||||
example_hidden_states_connector.cleanup_hidden_states(path)
|
||||
|
||||
if profile_dir:
|
||||
await engine.start_profile()
|
||||
|
||||
# Fill prompt queue
|
||||
prompt_queue: asyncio.Queue = asyncio.Queue()
|
||||
for idx, token_ids in enumerate(prompts):
|
||||
prompt_queue.put_nowait((idx, token_ids))
|
||||
# Sentinel per client
|
||||
for _ in range(num_clients):
|
||||
prompt_queue.put_nowait(None)
|
||||
|
||||
results: list[dict] = []
|
||||
consume_pool = ThreadPoolExecutor(max_workers=num_clients)
|
||||
|
||||
t0 = time.perf_counter()
|
||||
tasks = [
|
||||
asyncio.create_task(
|
||||
_client_loop(engine, prompt_queue, consume_pool, results, i)
|
||||
)
|
||||
for i in range(num_clients)
|
||||
]
|
||||
await asyncio.gather(*tasks)
|
||||
elapsed = time.perf_counter() - t0
|
||||
|
||||
consume_pool.shutdown(wait=True)
|
||||
|
||||
if profile_dir:
|
||||
await engine.stop_profile()
|
||||
|
||||
total_prompt_tokens = sum(r["num_prompt_tokens"] for r in results)
|
||||
num_prompts = len(results)
|
||||
mean_hidden = sum(r["mean_hidden_value"] for r in results) / max(
|
||||
len(results), 1
|
||||
)
|
||||
|
||||
return {
|
||||
"mode": "extract",
|
||||
"elapsed_s": elapsed,
|
||||
"num_prompts": num_prompts,
|
||||
"total_prompt_tokens": total_prompt_tokens,
|
||||
"tokens_per_s": total_prompt_tokens / elapsed,
|
||||
"prompts_per_s": num_prompts / elapsed,
|
||||
"mean_hidden_value": mean_hidden,
|
||||
}
|
||||
finally:
|
||||
engine.shutdown()
|
||||
|
||||
|
||||
def run_extraction(
|
||||
model: str,
|
||||
prompts: list[list[int]],
|
||||
num_clients: int,
|
||||
layers: list[int],
|
||||
extra_args: dict,
|
||||
profile_dir: str | None = None,
|
||||
) -> dict:
|
||||
return asyncio.run(
|
||||
_run_extraction_async(
|
||||
model,
|
||||
prompts,
|
||||
num_clients,
|
||||
layers,
|
||||
"/dev/shm",
|
||||
extra_args,
|
||||
profile_dir=profile_dir,
|
||||
)
|
||||
)
|
||||
|
||||
|
||||
def print_results(results: dict):
|
||||
mode = results["mode"]
|
||||
print(f"\n{'=' * 60}")
|
||||
print(f" {mode.upper()} RESULTS")
|
||||
print(f"{'=' * 60}")
|
||||
print(f" Prompts: {results['num_prompts']}")
|
||||
print(f" Total prompt tokens: {results['total_prompt_tokens']:,}")
|
||||
print(f" Wall time: {results['elapsed_s']:.2f}s")
|
||||
print(f" Tokens/s: {results['tokens_per_s']:,.0f}")
|
||||
print(f" Prompts/s: {results['prompts_per_s']:.2f}")
|
||||
if mode == "extract":
|
||||
print(f" Mean hidden value: {results['mean_hidden_value']:.6f}")
|
||||
print(f"{'=' * 60}\n")
|
||||
|
||||
|
||||
def main():
|
||||
parser = argparse.ArgumentParser(
|
||||
description="Benchmark hidden state extraction throughput"
|
||||
)
|
||||
parser.add_argument("--model", type=str, required=True)
|
||||
parser.add_argument("--num-prompts", type=int, default=64)
|
||||
parser.add_argument("--num-clients", type=int, default=8)
|
||||
parser.add_argument("--prompt-len", type=int, default=8192)
|
||||
parser.add_argument("--layers", type=int, nargs="+", default=[1, 2, 3, 4])
|
||||
parser.add_argument("--skip-baseline", action="store_true")
|
||||
parser.add_argument("--skip-extract", action="store_true")
|
||||
parser.add_argument("--gpu-memory-utilization", type=float, default=0.9)
|
||||
parser.add_argument("--max-num-batched-tokens", type=int, default=None)
|
||||
parser.add_argument("--max-cudagraph-capture-size", type=int, default=None)
|
||||
parser.add_argument("--max-model-len", type=int, default=None)
|
||||
parser.add_argument("--enforce-eager", action="store_true")
|
||||
parser.add_argument("--load-format", type=str, default=None)
|
||||
parser.add_argument(
|
||||
"--profile",
|
||||
action="store_true",
|
||||
help="Enable torch profiler for both baseline and extraction runs.",
|
||||
)
|
||||
parser.add_argument(
|
||||
"--torch-profiler-dir",
|
||||
type=str,
|
||||
default="./vllm_profile",
|
||||
help="Directory to save torch profiler traces (default: ./vllm_profile).",
|
||||
)
|
||||
parser.add_argument(
|
||||
"--enable-flashinfer-autotune",
|
||||
action="store_true",
|
||||
default=False,
|
||||
help="Enable FlashInfer autotuning (can be slow).",
|
||||
)
|
||||
args = parser.parse_args()
|
||||
|
||||
extra_args = {
|
||||
"gpu_memory_utilization": args.gpu_memory_utilization,
|
||||
}
|
||||
if args.max_model_len is not None:
|
||||
extra_args["max_model_len"] = args.max_model_len
|
||||
if args.max_num_batched_tokens is not None:
|
||||
extra_args["max_num_batched_tokens"] = args.max_num_batched_tokens
|
||||
if args.max_model_len and args.max_num_batched_tokens < args.max_model_len:
|
||||
raise ValueError(
|
||||
"max_num_batched_tokens must be >= max_model_len since chunked prefill"
|
||||
" is not supported by hidden state extraction."
|
||||
)
|
||||
if args.enforce_eager:
|
||||
extra_args["enforce_eager"] = True
|
||||
if args.load_format is not None:
|
||||
extra_args["load_format"] = args.load_format
|
||||
if args.max_cudagraph_capture_size is not None:
|
||||
extra_args["max_cudagraph_capture_size"] = args.max_cudagraph_capture_size
|
||||
extra_args["enable_flashinfer_autotune"] = args.enable_flashinfer_autotune
|
||||
|
||||
# Get vocab size from HF config without loading the full model
|
||||
hf_config = AutoConfig.from_pretrained(args.model, trust_remote_code=True)
|
||||
vocab_size = hf_config.vocab_size
|
||||
prompts = make_random_prompts(args.num_prompts, args.prompt_len, vocab_size)
|
||||
print(
|
||||
f"Generated {args.num_prompts} prompts, "
|
||||
f"{args.prompt_len} tokens each (vocab {vocab_size})"
|
||||
)
|
||||
|
||||
profile_dir = args.torch_profiler_dir if args.profile else None
|
||||
if profile_dir:
|
||||
print(f"Torch profiler enabled, traces will be saved to {profile_dir}/")
|
||||
|
||||
if not args.skip_baseline:
|
||||
baseline_profile_dir = f"{profile_dir}/baseline" if profile_dir else None
|
||||
baseline = run_baseline(
|
||||
args.model, prompts, extra_args, profile_dir=baseline_profile_dir
|
||||
)
|
||||
print_results(baseline)
|
||||
|
||||
if not args.skip_extract:
|
||||
extract_profile_dir = f"{profile_dir}/extract" if profile_dir else None
|
||||
extract = run_extraction(
|
||||
args.model,
|
||||
prompts,
|
||||
args.num_clients,
|
||||
args.layers,
|
||||
extra_args,
|
||||
profile_dir=extract_profile_dir,
|
||||
)
|
||||
print_results(extract)
|
||||
|
||||
if not args.skip_baseline and not args.skip_extract:
|
||||
slowdown = baseline["tokens_per_s"] / extract["tokens_per_s"]
|
||||
print("Extraction slowdown factor: {:.2f}x".format(slowdown))
|
||||
|
||||
|
||||
if __name__ == "__main__":
|
||||
main()
|
||||
@@ -0,0 +1,143 @@
|
||||
#!/bin/bash
|
||||
|
||||
# benchmark the overhead of disaggregated prefill.
|
||||
# methodology:
|
||||
# - send all request to prefill vLLM instance. It will buffer KV cache.
|
||||
# - then send all request to decode instance.
|
||||
# - The TTFT of decode instance is the overhead.
|
||||
|
||||
set -ex
|
||||
|
||||
kill_gpu_processes() {
|
||||
# kill all processes on GPU.
|
||||
pgrep pt_main_thread | xargs -r kill -9
|
||||
pgrep python3 | xargs -r kill -9
|
||||
# vLLM now names the process with VLLM prefix after https://github.com/vllm-project/vllm/pull/21445
|
||||
pgrep VLLM | xargs -r kill -9
|
||||
sleep 10
|
||||
|
||||
# remove vllm config file
|
||||
rm -rf ~/.config/vllm
|
||||
|
||||
# Print the GPU memory usage
|
||||
# so that we know if all GPU processes are killed.
|
||||
gpu_memory_usage=$(nvidia-smi --query-gpu=memory.used --format=csv,noheader,nounits -i 0)
|
||||
# The memory usage should be 0 MB.
|
||||
echo "GPU 0 Memory Usage: $gpu_memory_usage MB"
|
||||
}
|
||||
|
||||
wait_for_server() {
|
||||
# wait for vllm server to start
|
||||
# return 1 if vllm server crashes
|
||||
local port=$1
|
||||
timeout 1200 bash -c "
|
||||
until curl -s localhost:${port}/v1/completions > /dev/null; do
|
||||
sleep 1
|
||||
done" && return 0 || return 1
|
||||
}
|
||||
|
||||
|
||||
benchmark() {
|
||||
|
||||
export VLLM_LOGGING_LEVEL=DEBUG
|
||||
export VLLM_HOST_IP=$(hostname -I | awk '{print $1}')
|
||||
|
||||
# compare chunked prefill with disaggregated prefill
|
||||
|
||||
results_folder="./results"
|
||||
model="meta-llama/Meta-Llama-3.1-8B-Instruct"
|
||||
dataset_name="sonnet"
|
||||
dataset_path="../sonnet_4x.txt"
|
||||
num_prompts=10
|
||||
qps=$1
|
||||
prefix_len=50
|
||||
input_len=2048
|
||||
output_len=$2
|
||||
|
||||
|
||||
CUDA_VISIBLE_DEVICES=0 vllm serve $model \
|
||||
--port 8100 \
|
||||
--max-model-len 10000 \
|
||||
--gpu-memory-utilization 0.6 \
|
||||
--kv-transfer-config \
|
||||
'{"kv_connector":"P2pNcclConnector","kv_role":"kv_producer","kv_rank":0,"kv_parallel_size":2,"kv_buffer_size":5e9}' &
|
||||
|
||||
|
||||
CUDA_VISIBLE_DEVICES=1 vllm serve $model \
|
||||
--port 8200 \
|
||||
--max-model-len 10000 \
|
||||
--gpu-memory-utilization 0.6 \
|
||||
--kv-transfer-config \
|
||||
'{"kv_connector":"P2pNcclConnector","kv_role":"kv_consumer","kv_rank":1,"kv_parallel_size":2,"kv_buffer_size":5e9}' &
|
||||
|
||||
wait_for_server 8100
|
||||
wait_for_server 8200
|
||||
|
||||
# let the prefill instance finish prefill
|
||||
vllm bench serve \
|
||||
--backend vllm \
|
||||
--model $model \
|
||||
--dataset-name $dataset_name \
|
||||
--dataset-path $dataset_path \
|
||||
--sonnet-input-len $input_len \
|
||||
--sonnet-output-len "$output_len" \
|
||||
--sonnet-prefix-len $prefix_len \
|
||||
--num-prompts $num_prompts \
|
||||
--port 8100 \
|
||||
--save-result \
|
||||
--result-dir $results_folder \
|
||||
--result-filename disagg_prefill_tp1.json \
|
||||
--request-rate "inf"
|
||||
|
||||
|
||||
# send the request to decode.
|
||||
# The TTFT of this command will be the overhead of disagg prefill impl.
|
||||
vllm bench serve \
|
||||
--backend vllm \
|
||||
--model $model \
|
||||
--dataset-name $dataset_name \
|
||||
--dataset-path $dataset_path \
|
||||
--sonnet-input-len $input_len \
|
||||
--sonnet-output-len "$output_len" \
|
||||
--sonnet-prefix-len $prefix_len \
|
||||
--num-prompts $num_prompts \
|
||||
--port 8200 \
|
||||
--save-result \
|
||||
--result-dir $results_folder \
|
||||
--result-filename disagg_prefill_tp1_overhead.json \
|
||||
--request-rate "$qps"
|
||||
kill_gpu_processes
|
||||
|
||||
}
|
||||
|
||||
|
||||
main() {
|
||||
|
||||
(which wget && which curl) || (apt-get update && apt-get install -y wget curl)
|
||||
(which jq) || (apt-get -y install jq)
|
||||
(which socat) || (apt-get -y install socat)
|
||||
|
||||
pip install quart httpx datasets
|
||||
|
||||
cd "$(dirname "$0")"
|
||||
|
||||
cd ..
|
||||
# create sonnet-4x.txt
|
||||
echo "" > sonnet_4x.txt
|
||||
for _ in {1..4}
|
||||
do
|
||||
cat sonnet.txt >> sonnet_4x.txt
|
||||
done
|
||||
cd disagg_benchmarks
|
||||
|
||||
rm -rf results
|
||||
mkdir results
|
||||
|
||||
default_qps=1
|
||||
default_output_len=1
|
||||
benchmark $default_qps $default_output_len
|
||||
|
||||
}
|
||||
|
||||
|
||||
main "$@"
|
||||
@@ -0,0 +1,157 @@
|
||||
#!/bin/bash
|
||||
|
||||
# Requirement: 2x GPUs.
|
||||
|
||||
|
||||
# Model: meta-llama/Meta-Llama-3.1-8B-Instruct
|
||||
# Query: 1024 input tokens, 6 output tokens, QPS 2/4/6/8, 100 requests
|
||||
# Resource: 2x GPU
|
||||
# Approaches:
|
||||
# 2. Chunked prefill: 2 vllm instance with tp=4, equivalent to 1 tp=4 instance with QPS 4
|
||||
# 3. Disaggregated prefill: 1 prefilling instance and 1 decoding instance
|
||||
# Prefilling instance: max_output_token=1
|
||||
# Decoding instance: force the input tokens be the same across requests to bypass prefilling
|
||||
|
||||
set -ex
|
||||
|
||||
kill_gpu_processes() {
|
||||
# kill all processes on GPU.
|
||||
pgrep pt_main_thread | xargs -r kill -9
|
||||
pgrep python3 | xargs -r kill -9
|
||||
# vLLM now names the process with VLLM prefix after https://github.com/vllm-project/vllm/pull/21445
|
||||
pgrep VLLM | xargs -r kill -9
|
||||
for port in 8000 8100 8200; do lsof -t -i:$port | xargs -r kill -9; done
|
||||
sleep 1
|
||||
}
|
||||
|
||||
wait_for_server() {
|
||||
# wait for vllm server to start
|
||||
# return 1 if vllm server crashes
|
||||
local port=$1
|
||||
timeout 1200 bash -c "
|
||||
until curl -s localhost:${port}/v1/completions > /dev/null; do
|
||||
sleep 1
|
||||
done" && return 0 || return 1
|
||||
}
|
||||
|
||||
|
||||
launch_chunked_prefill() {
|
||||
model="meta-llama/Meta-Llama-3.1-8B-Instruct"
|
||||
# disagg prefill
|
||||
CUDA_VISIBLE_DEVICES=0 vllm serve $model \
|
||||
--port 8100 \
|
||||
--max-model-len 10000 \
|
||||
--enable-chunked-prefill \
|
||||
--gpu-memory-utilization 0.6 &
|
||||
CUDA_VISIBLE_DEVICES=1 vllm serve $model \
|
||||
--port 8200 \
|
||||
--max-model-len 10000 \
|
||||
--enable-chunked-prefill \
|
||||
--gpu-memory-utilization 0.6 &
|
||||
wait_for_server 8100
|
||||
wait_for_server 8200
|
||||
python3 round_robin_proxy.py &
|
||||
sleep 1
|
||||
}
|
||||
|
||||
|
||||
launch_disagg_prefill() {
|
||||
model="meta-llama/Meta-Llama-3.1-8B-Instruct"
|
||||
# disagg prefill
|
||||
CUDA_VISIBLE_DEVICES=0 vllm serve $model \
|
||||
--port 8100 \
|
||||
--max-model-len 10000 \
|
||||
--gpu-memory-utilization 0.6 \
|
||||
--kv-transfer-config \
|
||||
'{"kv_connector":"P2pNcclConnector","kv_role":"kv_producer","kv_rank":0,"kv_parallel_size":2,"kv_buffer_size":5e9}' &
|
||||
|
||||
CUDA_VISIBLE_DEVICES=1 vllm serve $model \
|
||||
--port 8200 \
|
||||
--max-model-len 10000 \
|
||||
--gpu-memory-utilization 0.6 \
|
||||
--kv-transfer-config \
|
||||
'{"kv_connector":"P2pNcclConnector","kv_role":"kv_consumer","kv_rank":1,"kv_parallel_size":2,"kv_buffer_size":5e9}' &
|
||||
|
||||
wait_for_server 8100
|
||||
wait_for_server 8200
|
||||
python3 disagg_prefill_proxy_server.py &
|
||||
sleep 1
|
||||
}
|
||||
|
||||
|
||||
benchmark() {
|
||||
results_folder="./results"
|
||||
model="meta-llama/Meta-Llama-3.1-8B-Instruct"
|
||||
dataset_name="sonnet"
|
||||
dataset_path="../sonnet_4x.txt"
|
||||
num_prompts=100
|
||||
qps=$1
|
||||
prefix_len=50
|
||||
input_len=1024
|
||||
output_len=$2
|
||||
tag=$3
|
||||
|
||||
vllm bench serve \
|
||||
--backend vllm \
|
||||
--model $model \
|
||||
--dataset-name $dataset_name \
|
||||
--dataset-path $dataset_path \
|
||||
--sonnet-input-len $input_len \
|
||||
--sonnet-output-len "$output_len" \
|
||||
--sonnet-prefix-len $prefix_len \
|
||||
--num-prompts $num_prompts \
|
||||
--port 8000 \
|
||||
--save-result \
|
||||
--result-dir $results_folder \
|
||||
--result-filename "$tag"-qps-"$qps".json \
|
||||
--request-rate "$qps"
|
||||
|
||||
sleep 2
|
||||
}
|
||||
|
||||
|
||||
main() {
|
||||
|
||||
(which wget && which curl) || (apt-get update && apt-get install -y wget curl)
|
||||
(which jq) || (apt-get -y install jq)
|
||||
(which socat) || (apt-get -y install socat)
|
||||
(which lsof) || (apt-get -y install lsof)
|
||||
|
||||
pip install quart httpx matplotlib aiohttp datasets
|
||||
|
||||
cd "$(dirname "$0")"
|
||||
|
||||
cd ..
|
||||
# create sonnet-4x.txt so that we can sample 2048 tokens for input
|
||||
echo "" > sonnet_4x.txt
|
||||
for _ in {1..4}
|
||||
do
|
||||
cat sonnet.txt >> sonnet_4x.txt
|
||||
done
|
||||
cd disagg_benchmarks
|
||||
|
||||
rm -rf results
|
||||
mkdir results
|
||||
|
||||
default_output_len=6
|
||||
|
||||
export VLLM_HOST_IP=$(hostname -I | awk '{print $1}')
|
||||
|
||||
launch_chunked_prefill
|
||||
for qps in 2 4 6 8; do
|
||||
benchmark $qps $default_output_len chunked_prefill
|
||||
done
|
||||
kill_gpu_processes
|
||||
|
||||
launch_disagg_prefill
|
||||
for qps in 2 4 6 8; do
|
||||
benchmark $qps $default_output_len disagg_prefill
|
||||
done
|
||||
kill_gpu_processes
|
||||
|
||||
python3 visualize_benchmark_results.py
|
||||
|
||||
}
|
||||
|
||||
|
||||
main "$@"
|
||||
@@ -0,0 +1,260 @@
|
||||
# SPDX-License-Identifier: Apache-2.0
|
||||
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
import argparse
|
||||
import asyncio
|
||||
import logging
|
||||
import os
|
||||
import time
|
||||
import uuid
|
||||
from urllib.parse import urlparse
|
||||
|
||||
import aiohttp
|
||||
from quart import Quart, Response, make_response, request
|
||||
|
||||
# Configure logging
|
||||
logging.basicConfig(level=logging.INFO)
|
||||
logger = logging.getLogger(__name__)
|
||||
|
||||
|
||||
def parse_args():
|
||||
"""parse command line arguments"""
|
||||
parser = argparse.ArgumentParser(description="vLLM P/D disaggregation proxy server")
|
||||
|
||||
# Add args
|
||||
parser.add_argument(
|
||||
"--timeout",
|
||||
type=float,
|
||||
default=6 * 60 * 60,
|
||||
help="Timeout for backend service requests in seconds (default: 21600)",
|
||||
)
|
||||
parser.add_argument(
|
||||
"--port",
|
||||
type=int,
|
||||
default=8000,
|
||||
help="Port to run the server on (default: 8000)",
|
||||
)
|
||||
parser.add_argument(
|
||||
"--prefill-url",
|
||||
type=str,
|
||||
default="http://localhost:8100",
|
||||
help="Prefill service base URL (protocol + host[:port])",
|
||||
)
|
||||
parser.add_argument(
|
||||
"--decode-url",
|
||||
type=str,
|
||||
default="http://localhost:8200",
|
||||
help="Decode service base URL (protocol + host[:port])",
|
||||
)
|
||||
parser.add_argument(
|
||||
"--kv-host",
|
||||
type=str,
|
||||
default="localhost",
|
||||
help="Hostname or IP used by KV transfer (default: localhost)",
|
||||
)
|
||||
parser.add_argument(
|
||||
"--prefill-kv-port",
|
||||
type=int,
|
||||
default=14579,
|
||||
help="Prefill KV port (default: 14579)",
|
||||
)
|
||||
parser.add_argument(
|
||||
"--decode-kv-port",
|
||||
type=int,
|
||||
default=14580,
|
||||
help="Decode KV port (default: 14580)",
|
||||
)
|
||||
|
||||
return parser.parse_args()
|
||||
|
||||
|
||||
def main():
|
||||
"""parse command line arguments"""
|
||||
args = parse_args()
|
||||
|
||||
# Initialize configuration using command line parameters
|
||||
AIOHTTP_TIMEOUT = aiohttp.ClientTimeout(total=args.timeout)
|
||||
PREFILL_SERVICE_URL = args.prefill_url
|
||||
DECODE_SERVICE_URL = args.decode_url
|
||||
PORT = args.port
|
||||
|
||||
PREFILL_KV_ADDR = f"{args.kv_host}:{args.prefill_kv_port}"
|
||||
DECODE_KV_ADDR = f"{args.kv_host}:{args.decode_kv_port}"
|
||||
|
||||
logger.info(
|
||||
"Proxy resolved KV addresses -> prefill: %s, decode: %s",
|
||||
PREFILL_KV_ADDR,
|
||||
DECODE_KV_ADDR,
|
||||
)
|
||||
|
||||
app = Quart(__name__)
|
||||
|
||||
# Attach the configuration object to the application instance so helper
|
||||
# coroutines can read the resolved backend URLs and timeouts without using
|
||||
# globals.
|
||||
app.config.update(
|
||||
{
|
||||
"AIOHTTP_TIMEOUT": AIOHTTP_TIMEOUT,
|
||||
"PREFILL_SERVICE_URL": PREFILL_SERVICE_URL,
|
||||
"DECODE_SERVICE_URL": DECODE_SERVICE_URL,
|
||||
"PREFILL_KV_ADDR": PREFILL_KV_ADDR,
|
||||
"DECODE_KV_ADDR": DECODE_KV_ADDR,
|
||||
}
|
||||
)
|
||||
|
||||
def _normalize_base_url(url: str) -> str:
|
||||
"""Remove any trailing slash so path joins behave predictably."""
|
||||
return url.rstrip("/")
|
||||
|
||||
def _get_host_port(url: str) -> str:
|
||||
"""Return the hostname:port portion for logging and KV headers."""
|
||||
parsed = urlparse(url)
|
||||
host = parsed.hostname or "localhost"
|
||||
port = parsed.port
|
||||
if port is None:
|
||||
port = 80 if parsed.scheme == "http" else 443
|
||||
return f"{host}:{port}"
|
||||
|
||||
PREFILL_BASE = _normalize_base_url(PREFILL_SERVICE_URL)
|
||||
DECODE_BASE = _normalize_base_url(DECODE_SERVICE_URL)
|
||||
KV_TARGET = _get_host_port(DECODE_SERVICE_URL)
|
||||
|
||||
def _build_headers(request_id: str) -> dict[str, str]:
|
||||
"""Construct the headers expected by vLLM's P2P disagg connector."""
|
||||
headers: dict[str, str] = {"X-Request-Id": request_id, "X-KV-Target": KV_TARGET}
|
||||
api_key = os.environ.get("OPENAI_API_KEY")
|
||||
if api_key:
|
||||
headers["Authorization"] = f"Bearer {api_key}"
|
||||
return headers
|
||||
|
||||
async def _run_prefill(
|
||||
request_path: str,
|
||||
payload: dict,
|
||||
headers: dict[str, str],
|
||||
request_id: str,
|
||||
):
|
||||
url = f"{PREFILL_BASE}{request_path}"
|
||||
start_ts = time.perf_counter()
|
||||
logger.info("[prefill] start request_id=%s url=%s", request_id, url)
|
||||
try:
|
||||
async with (
|
||||
aiohttp.ClientSession(timeout=AIOHTTP_TIMEOUT) as session,
|
||||
session.post(url=url, json=payload, headers=headers) as resp,
|
||||
):
|
||||
if resp.status != 200:
|
||||
error_text = await resp.text()
|
||||
raise RuntimeError(
|
||||
f"Prefill backend error {resp.status}: {error_text}"
|
||||
)
|
||||
await resp.read()
|
||||
logger.info(
|
||||
"[prefill] done request_id=%s status=%s elapsed=%.2fs",
|
||||
request_id,
|
||||
resp.status,
|
||||
time.perf_counter() - start_ts,
|
||||
)
|
||||
except asyncio.TimeoutError as exc:
|
||||
raise RuntimeError(f"Prefill service timeout at {url}") from exc
|
||||
except aiohttp.ClientError as exc:
|
||||
raise RuntimeError(f"Prefill service unavailable at {url}") from exc
|
||||
|
||||
async def _stream_decode(
|
||||
request_path: str,
|
||||
payload: dict,
|
||||
headers: dict[str, str],
|
||||
request_id: str,
|
||||
):
|
||||
url = f"{DECODE_BASE}{request_path}"
|
||||
# Stream tokens from the decode service once the prefill stage has
|
||||
# materialized KV caches on the target workers.
|
||||
logger.info("[decode] start request_id=%s url=%s", request_id, url)
|
||||
try:
|
||||
async with (
|
||||
aiohttp.ClientSession(timeout=AIOHTTP_TIMEOUT) as session,
|
||||
session.post(url=url, json=payload, headers=headers) as resp,
|
||||
):
|
||||
if resp.status != 200:
|
||||
error_text = await resp.text()
|
||||
logger.error(
|
||||
"Decode backend error %s - %s", resp.status, error_text
|
||||
)
|
||||
err_msg = (
|
||||
'{"error": "Decode backend error ' + str(resp.status) + '"}'
|
||||
)
|
||||
yield err_msg.encode()
|
||||
return
|
||||
logger.info(
|
||||
"[decode] streaming response request_id=%s status=%s",
|
||||
request_id,
|
||||
resp.status,
|
||||
)
|
||||
async for chunk_bytes in resp.content.iter_chunked(1024):
|
||||
yield chunk_bytes
|
||||
logger.info("[decode] finished streaming request_id=%s", request_id)
|
||||
except asyncio.TimeoutError:
|
||||
logger.error("Decode service timeout at %s", url)
|
||||
yield b'{"error": "Decode service timeout"}'
|
||||
except aiohttp.ClientError as exc:
|
||||
logger.error("Decode service error at %s: %s", url, exc)
|
||||
yield b'{"error": "Decode service unavailable"}'
|
||||
|
||||
async def process_request():
|
||||
"""Process a single request through prefill and decode stages"""
|
||||
try:
|
||||
original_request_data = await request.get_json()
|
||||
|
||||
# Create prefill request (max_tokens=1)
|
||||
prefill_request = original_request_data.copy()
|
||||
prefill_request["max_tokens"] = 1
|
||||
if "max_completion_tokens" in prefill_request:
|
||||
prefill_request["max_completion_tokens"] = 1
|
||||
|
||||
# Execute prefill stage
|
||||
# The request id encodes both KV socket addresses so the backend can
|
||||
# shuttle tensors directly via NCCL once the prefill response
|
||||
# completes.
|
||||
request_id = (
|
||||
f"___prefill_addr_{PREFILL_KV_ADDR}___decode_addr_"
|
||||
f"{DECODE_KV_ADDR}_{uuid.uuid4().hex}"
|
||||
)
|
||||
|
||||
headers = _build_headers(request_id)
|
||||
await _run_prefill(request.path, prefill_request, headers, request_id)
|
||||
|
||||
# Execute decode stage and stream response
|
||||
# Pass the unmodified user request so the decode phase can continue
|
||||
# sampling with the already-populated KV cache.
|
||||
generator = _stream_decode(
|
||||
request.path, original_request_data, headers, request_id
|
||||
)
|
||||
response = await make_response(generator)
|
||||
response.timeout = None # Disable timeout for streaming response
|
||||
return response
|
||||
|
||||
except Exception:
|
||||
logger.exception("Error processing request")
|
||||
return Response(
|
||||
response=b'{"error": "Internal server error"}',
|
||||
status=500,
|
||||
content_type="application/json",
|
||||
)
|
||||
|
||||
@app.route("/v1/completions", methods=["POST"])
|
||||
async def handle_request():
|
||||
"""Handle incoming API requests with concurrency and rate limiting"""
|
||||
try:
|
||||
return await process_request()
|
||||
except asyncio.CancelledError:
|
||||
logger.warning("Request cancelled")
|
||||
return Response(
|
||||
response=b'{"error": "Request cancelled"}',
|
||||
status=503,
|
||||
content_type="application/json",
|
||||
)
|
||||
|
||||
# Start the Quart server with host can be set to 0.0.0.0
|
||||
app.run(port=PORT)
|
||||
|
||||
|
||||
if __name__ == "__main__":
|
||||
main()
|
||||
@@ -0,0 +1,63 @@
|
||||
# SPDX-License-Identifier: Apache-2.0
|
||||
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
import asyncio
|
||||
import itertools
|
||||
|
||||
import aiohttp
|
||||
from aiohttp import web
|
||||
|
||||
|
||||
class RoundRobinProxy:
|
||||
def __init__(self, target_ports):
|
||||
self.target_ports = target_ports
|
||||
self.port_cycle = itertools.cycle(self.target_ports)
|
||||
|
||||
async def handle_request(self, request):
|
||||
target_port = next(self.port_cycle)
|
||||
target_url = f"http://localhost:{target_port}{request.path_qs}"
|
||||
|
||||
async with aiohttp.ClientSession() as session:
|
||||
try:
|
||||
# Forward the request
|
||||
async with session.request(
|
||||
method=request.method,
|
||||
url=target_url,
|
||||
headers=request.headers,
|
||||
data=request.content,
|
||||
) as response:
|
||||
# Start sending the response
|
||||
resp = web.StreamResponse(
|
||||
status=response.status, headers=response.headers
|
||||
)
|
||||
await resp.prepare(request)
|
||||
|
||||
# Stream the response content
|
||||
async for chunk in response.content.iter_any():
|
||||
await resp.write(chunk)
|
||||
|
||||
await resp.write_eof()
|
||||
return resp
|
||||
|
||||
except Exception as e:
|
||||
return web.Response(text=f"Error: {str(e)}", status=500)
|
||||
|
||||
|
||||
async def main():
|
||||
proxy = RoundRobinProxy([8100, 8200])
|
||||
app = web.Application()
|
||||
app.router.add_route("*", "/{path:.*}", proxy.handle_request)
|
||||
|
||||
runner = web.AppRunner(app)
|
||||
await runner.setup()
|
||||
site = web.TCPSite(runner, "localhost", 8000)
|
||||
await site.start()
|
||||
|
||||
print("Proxy server started on http://localhost:8000")
|
||||
|
||||
# Keep the server running
|
||||
await asyncio.Event().wait()
|
||||
|
||||
|
||||
if __name__ == "__main__":
|
||||
asyncio.run(main())
|
||||
@@ -0,0 +1,47 @@
|
||||
# SPDX-License-Identifier: Apache-2.0
|
||||
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
import json
|
||||
|
||||
import matplotlib.pyplot as plt
|
||||
import pandas as pd
|
||||
|
||||
if __name__ == "__main__":
|
||||
data = []
|
||||
for name in ["disagg_prefill", "chunked_prefill"]:
|
||||
for qps in [2, 4, 6, 8]:
|
||||
with open(f"results/{name}-qps-{qps}.json") as f:
|
||||
x = json.load(f)
|
||||
x["name"] = name
|
||||
x["qps"] = qps
|
||||
data.append(x)
|
||||
|
||||
df = pd.DataFrame.from_dict(data)
|
||||
dis_df = df[df["name"] == "disagg_prefill"]
|
||||
chu_df = df[df["name"] == "chunked_prefill"]
|
||||
|
||||
plt.style.use("bmh")
|
||||
plt.rcParams["font.size"] = 20
|
||||
|
||||
for key in [
|
||||
"mean_ttft_ms",
|
||||
"median_ttft_ms",
|
||||
"p99_ttft_ms",
|
||||
"mean_itl_ms",
|
||||
"median_itl_ms",
|
||||
"p99_itl_ms",
|
||||
]:
|
||||
fig, ax = plt.subplots(figsize=(11, 7))
|
||||
plt.plot(
|
||||
dis_df["qps"], dis_df[key], label="disagg_prefill", marker="o", linewidth=4
|
||||
)
|
||||
plt.plot(
|
||||
chu_df["qps"], chu_df[key], label="chunked_prefill", marker="o", linewidth=4
|
||||
)
|
||||
ax.legend()
|
||||
|
||||
ax.set_xlabel("QPS")
|
||||
ax.set_ylabel(key)
|
||||
ax.set_ylim(bottom=0)
|
||||
fig.savefig(f"results/{key}.png")
|
||||
plt.close(fig)
|
||||
@@ -33,7 +33,6 @@ from vllm.distributed.device_communicators.custom_all_reduce import CustomAllred
|
||||
from vllm.distributed.device_communicators.flashinfer_all_reduce import (
|
||||
FlashInferAllReduce,
|
||||
)
|
||||
from vllm.distributed.device_communicators.push_all_reduce import PushAllReduce
|
||||
from vllm.distributed.device_communicators.pynccl import (
|
||||
PyNcclCommunicator,
|
||||
register_nccl_symmetric_ops,
|
||||
@@ -81,7 +80,6 @@ class CommunicatorBenchmark:
|
||||
|
||||
# Initialize communicators
|
||||
self.custom_allreduce = None
|
||||
self.push_ar_comm = None
|
||||
self.pynccl_comm = None
|
||||
self.symm_mem_comm = None
|
||||
self.symm_mem_comm_multimem = None
|
||||
@@ -108,23 +106,6 @@ class CommunicatorBenchmark:
|
||||
)
|
||||
self.custom_allreduce = None
|
||||
|
||||
try:
|
||||
self.push_ar_comm = PushAllReduce(
|
||||
group=self.cpu_group,
|
||||
device=self.device,
|
||||
max_size=self.max_size_override,
|
||||
)
|
||||
if not self.push_ar_comm.disabled:
|
||||
logger.info("Rank %s: PushAllReduce initialized", self.rank)
|
||||
else:
|
||||
logger.info("Rank %s: PushAllReduce disabled", self.rank)
|
||||
self.push_ar_comm = None
|
||||
except Exception as e:
|
||||
logger.warning(
|
||||
"Rank %s: Failed to initialize PushAllReduce: %s", self.rank, e
|
||||
)
|
||||
self.push_ar_comm = None
|
||||
|
||||
try:
|
||||
self.pynccl_comm = PyNcclCommunicator(
|
||||
group=self.cpu_group, device=self.device
|
||||
@@ -235,19 +216,6 @@ class CommunicatorBenchmark:
|
||||
)
|
||||
)
|
||||
|
||||
if self.push_ar_comm is not None:
|
||||
comm = self.push_ar_comm
|
||||
communicators.append(
|
||||
(
|
||||
"push_ar",
|
||||
lambda t, c=comm: c.all_reduce(t),
|
||||
lambda t, c=comm: c.should_use(t),
|
||||
comm.capture(),
|
||||
{},
|
||||
None,
|
||||
)
|
||||
)
|
||||
|
||||
if self.pynccl_comm is not None:
|
||||
comm = self.pynccl_comm
|
||||
communicators.append(
|
||||
|
||||
@@ -1,277 +0,0 @@
|
||||
# SPDX-License-Identifier: Apache-2.0
|
||||
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
# Copyright (c) 2025 FlyDSL Project Contributors
|
||||
|
||||
import json
|
||||
import os
|
||||
|
||||
import torch
|
||||
from aiter.test_common import run_perftest
|
||||
|
||||
from vllm.model_executor.layers.fused_moe import fused_experts
|
||||
from vllm.model_executor.layers.fused_moe.activation import MoEActivation
|
||||
from vllm.model_executor.layers.fused_moe.config import (
|
||||
int4_w4a16_moe_quant_config,
|
||||
)
|
||||
from vllm.model_executor.layers.fused_moe.fused_flydsl_moe import fused_flydsl_moe
|
||||
from vllm.model_executor.layers.quantization.compressed_tensors.compressed_tensors_moe import ( # noqa: E501
|
||||
compressed_tensors_moe_w4a16_flydsl,
|
||||
)
|
||||
from vllm.platforms import current_platform
|
||||
|
||||
RoutingBuffers = tuple[
|
||||
torch.Tensor, # sorted_token_ids
|
||||
torch.Tensor, # sorted_weights
|
||||
torch.Tensor, # sorted_expert_ids
|
||||
torch.Tensor, # num_valid_ids (shape [1], i32)
|
||||
int, # sorted_size
|
||||
int, # blocks
|
||||
]
|
||||
|
||||
MODEL_PARAMS_TO_TUNE = [
|
||||
# (num_experts, inter_dim, hidden_size, topk)
|
||||
(384, 256, 7168, 8), # Kimi K2.5 TP=8
|
||||
(384, 512, 7168, 8), # Kimi K2.5 TP=4
|
||||
]
|
||||
|
||||
NUM_TOKENS_TO_TUNE = [
|
||||
1,
|
||||
2,
|
||||
4,
|
||||
8,
|
||||
16,
|
||||
24,
|
||||
32,
|
||||
48,
|
||||
64,
|
||||
128,
|
||||
256,
|
||||
512,
|
||||
1024,
|
||||
2048,
|
||||
4096,
|
||||
8192,
|
||||
]
|
||||
|
||||
TILE_M_SEARCH_SPACE = [16, 32, 64, 128, 256]
|
||||
TILE_N_SEARCH_SPACE = [16, 32, 64, 128, 256]
|
||||
TILE_K_SEARCH_SPACE = [16, 32, 64, 128, 256, 512]
|
||||
TILE_N2_SEARCH_SPACE = [16, 32, 64, 128, 256]
|
||||
TILE_K2_SEARCH_SPACE = [16, 32, 64, 128, 256, 512]
|
||||
|
||||
TILE_CONFIGS = []
|
||||
for tile_m in TILE_M_SEARCH_SPACE:
|
||||
for tile_n in TILE_N_SEARCH_SPACE:
|
||||
for tile_k in TILE_K_SEARCH_SPACE:
|
||||
for tile_n2 in TILE_N2_SEARCH_SPACE:
|
||||
for tile_k2 in TILE_K2_SEARCH_SPACE:
|
||||
TILE_CONFIGS.append(
|
||||
{
|
||||
"tile_m": tile_m,
|
||||
"tile_n": tile_n,
|
||||
"tile_k": tile_k,
|
||||
"tile_n2": tile_n2,
|
||||
"tile_k2": tile_k2,
|
||||
}
|
||||
)
|
||||
|
||||
|
||||
def tune_flydsl_moe_w4a16(
|
||||
device: str = "cuda", num_iters: int = 100, num_warmup: int = 10
|
||||
):
|
||||
packed_factor = 8
|
||||
w13_num_shards = 2
|
||||
params_dtype = torch.bfloat16
|
||||
group_size = 32
|
||||
scale_factor = 0.01
|
||||
|
||||
for model_params in MODEL_PARAMS_TO_TUNE:
|
||||
num_experts = model_params[0]
|
||||
inter_dim = model_params[1]
|
||||
hidden_size = model_params[2]
|
||||
topk = model_params[3]
|
||||
print(
|
||||
f"\nTuning: num_experts={num_experts}, inter_dim={inter_dim}, "
|
||||
f"hidden_size={hidden_size}, topk={topk}...\n"
|
||||
)
|
||||
|
||||
w2_scales_size = inter_dim
|
||||
num_groups_w2 = w2_scales_size // group_size
|
||||
num_groups_w13 = hidden_size // group_size
|
||||
|
||||
w13_weight = torch.randint(
|
||||
0,
|
||||
255,
|
||||
(num_experts, hidden_size // packed_factor, w13_num_shards * inter_dim),
|
||||
dtype=torch.int32,
|
||||
device=device,
|
||||
)
|
||||
|
||||
w2_weight = torch.randint(
|
||||
0,
|
||||
255,
|
||||
(num_experts, inter_dim // packed_factor, hidden_size),
|
||||
dtype=torch.int32,
|
||||
device=device,
|
||||
)
|
||||
w13_scale = scale_factor * torch.randn(
|
||||
num_experts,
|
||||
num_groups_w13,
|
||||
w13_num_shards * inter_dim,
|
||||
dtype=params_dtype,
|
||||
device=device,
|
||||
)
|
||||
w2_scale = scale_factor * torch.randn(
|
||||
num_experts, num_groups_w2, hidden_size, dtype=params_dtype, device=device
|
||||
)
|
||||
|
||||
w13 = w13_weight
|
||||
w13 = compressed_tensors_moe_w4a16_flydsl._gptq_int32_to_flydsl_packed(w13)
|
||||
w13 = w13.view(-1).contiguous()
|
||||
|
||||
w2 = w2_weight
|
||||
w2 = compressed_tensors_moe_w4a16_flydsl._gptq_int32_to_flydsl_packed(w2)
|
||||
w2 = w2.view(-1).contiguous()
|
||||
|
||||
w13_scale_flydsl = w13_scale
|
||||
w2_scale_flydsl = w2_scale
|
||||
|
||||
if group_size > 0 and w13_scale.dim() == 3 and w13_scale.shape[1] > 1:
|
||||
E, G, N = w13_scale.shape
|
||||
w13_scale_flydsl = (
|
||||
w13_scale_flydsl.view(E, G // 2, 2, N)
|
||||
.permute(0, 1, 3, 2)
|
||||
.contiguous()
|
||||
.view(-1)
|
||||
.contiguous()
|
||||
)
|
||||
elif w13_scale.dim() == 3 and w13_scale.shape[1] == 1:
|
||||
w13_scale_flydsl = w13_scale_flydsl.squeeze(1)
|
||||
|
||||
if group_size > 0 and w2_scale.dim() == 3 and w2_scale.shape[1] > 1:
|
||||
E, G, N = w2_scale.shape
|
||||
w2_scale_flydsl = (
|
||||
w2_scale_flydsl.view(E, G // 2, 2, N)
|
||||
.permute(0, 1, 3, 2)
|
||||
.contiguous()
|
||||
.view(-1)
|
||||
.contiguous()
|
||||
)
|
||||
elif w2_scale.dim() == 3 and w2_scale.shape[1] == 1:
|
||||
w2_scale_flydsl = w2_scale_flydsl.squeeze(1)
|
||||
|
||||
w13_scale_flydsl = w13_scale_flydsl.contiguous()
|
||||
w2_scale_flydsl = w2_scale_flydsl.contiguous()
|
||||
|
||||
w13.is_shuffled = True
|
||||
w2.is_shuffled = True
|
||||
|
||||
w13_weight_scale = w13_scale.transpose(1, 2).contiguous()
|
||||
w2_weight_scale = w2_scale.transpose(1, 2).contiguous()
|
||||
w13_weight_packed = w13_weight.transpose(1, 2).contiguous().view(torch.uint8)
|
||||
w2_weight_packed = w2_weight.transpose(1, 2).contiguous().view(torch.uint8)
|
||||
|
||||
moe_quant_config = int4_w4a16_moe_quant_config(
|
||||
w1_scale=w13_weight_scale,
|
||||
w2_scale=w2_weight_scale,
|
||||
w1_zp=None,
|
||||
w2_zp=None,
|
||||
block_shape=[0, group_size],
|
||||
)
|
||||
|
||||
tuned_config = {}
|
||||
|
||||
for num_tokens in NUM_TOKENS_TO_TUNE:
|
||||
score = torch.rand(
|
||||
(num_tokens, num_experts), device=device, dtype=torch.float32
|
||||
)
|
||||
topk_vals, topk_ids = torch.topk(score, k=topk, dim=1)
|
||||
topk_weights = torch.softmax(topk_vals, dim=1).to(torch.float32)
|
||||
x = torch.randn(
|
||||
(num_tokens, hidden_size), dtype=torch.bfloat16, device=device
|
||||
)
|
||||
us_best = float("inf")
|
||||
for tile_config in TILE_CONFIGS:
|
||||
try:
|
||||
tile_m = tile_config["tile_m"]
|
||||
tile_n = tile_config["tile_n"]
|
||||
tile_k = tile_config["tile_k"]
|
||||
tile_n2 = tile_config["tile_n2"]
|
||||
tile_k2 = tile_config["tile_k2"]
|
||||
|
||||
model_dim = x.shape[1]
|
||||
assert model_dim % 64 == 0
|
||||
assert model_dim % tile_k == 0
|
||||
assert inter_dim % tile_n == 0
|
||||
assert model_dim % tile_n2 == 0
|
||||
assert inter_dim % tile_k2 == 0
|
||||
assert ((tile_m * tile_k2) % 256) == 0
|
||||
bytes_per_thread_x = (tile_m * tile_k2) // 256
|
||||
assert (bytes_per_thread_x % 4) == 0
|
||||
|
||||
out, _us = run_perftest(
|
||||
fused_flydsl_moe,
|
||||
x,
|
||||
w13,
|
||||
w2,
|
||||
num_experts,
|
||||
inter_dim,
|
||||
topk_weights,
|
||||
topk_ids,
|
||||
num_iters=num_iters,
|
||||
num_warmup=num_warmup,
|
||||
w1_scale=w13_scale_flydsl,
|
||||
w2_scale=w2_scale_flydsl,
|
||||
topk=topk_weights.shape[-1],
|
||||
group_size=group_size,
|
||||
doweight_stage1=False,
|
||||
scale_is_bf16=True,
|
||||
config=tile_config,
|
||||
)
|
||||
torch.accelerator.synchronize()
|
||||
except Exception:
|
||||
torch.accelerator.synchronize()
|
||||
continue
|
||||
else:
|
||||
us = _us.item()
|
||||
if us < us_best:
|
||||
out_ref = fused_experts(
|
||||
x,
|
||||
w13_weight_packed,
|
||||
w2_weight_packed,
|
||||
topk_weights=topk_weights,
|
||||
topk_ids=topk_ids,
|
||||
activation=MoEActivation.SILU,
|
||||
apply_router_weight_on_input=False,
|
||||
global_num_experts=num_experts,
|
||||
expert_map=None,
|
||||
quant_config=moe_quant_config,
|
||||
)
|
||||
try:
|
||||
assert torch.allclose(out, out_ref, atol=0.5, rtol=0.1)
|
||||
except Exception:
|
||||
continue
|
||||
else:
|
||||
print(
|
||||
f"For [num_tokens={num_tokens}, num_experts={num_experts}, " # noqa: E501
|
||||
f"inter_dim={inter_dim}] found new best " # noqa: E501
|
||||
f"config={tile_config}, us={us:0.3f}"
|
||||
)
|
||||
us_best = us
|
||||
tuned_config[str(num_tokens)] = tile_config
|
||||
device_name = current_platform.get_device_name().replace(" ", "_")
|
||||
tuned_config_file_name = (
|
||||
f"E={num_experts},N={inter_dim},device_name={device_name},"
|
||||
f"dtype=int4_w4a16,backend=flydsl.json"
|
||||
)
|
||||
tuner_dir_path = os.path.dirname(os.path.realpath(__file__))
|
||||
store_path = os.path.join(tuner_dir_path, tuned_config_file_name)
|
||||
with open(store_path, "w") as f:
|
||||
json.dump(tuned_config, f, indent=4)
|
||||
print(
|
||||
f"\nTuned config for num_tokens={num_tokens} was stored at {store_path}\n" # noqa: E501
|
||||
)
|
||||
|
||||
|
||||
if __name__ == "__main__":
|
||||
tune_flydsl_moe_w4a16(device="cuda")
|
||||
@@ -1,465 +0,0 @@
|
||||
# SPDX-License-Identifier: Apache-2.0
|
||||
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
"""Benchmark the fused MoE-LoRA fast path (one-shot) vs two-kernel baseline.
|
||||
|
||||
The "one_shot" provider goes through `vllm.lora.ops.triton_ops.fused_moe_lora`
|
||||
which dispatches to the single-kernel one-shot implementation when
|
||||
fully_sharded=False (the prefill default).
|
||||
|
||||
The "two_kernel" provider drives `fused_moe_lora_shrink` + `fused_moe_lora_expand`
|
||||
directly, bypassing the dispatch and matching the legacy two-kernel path's
|
||||
work distribution. This isolates the win from kernel fusion.
|
||||
|
||||
Run:
|
||||
.venv/bin/python -m benchmarks.kernels.benchmark_fused_moe_lora_one_shot
|
||||
.venv/bin/python -m benchmarks.kernels.benchmark_fused_moe_lora_one_shot \\
|
||||
--model qwen3moe
|
||||
"""
|
||||
|
||||
from __future__ import annotations
|
||||
|
||||
import argparse
|
||||
import os
|
||||
import random
|
||||
|
||||
import torch
|
||||
|
||||
from vllm import _custom_ops as ops
|
||||
from vllm.lora.ops.triton_ops import (
|
||||
fused_moe_lora,
|
||||
fused_moe_lora_expand,
|
||||
fused_moe_lora_shrink,
|
||||
)
|
||||
from vllm.triton_utils import triton
|
||||
|
||||
DTYPE = torch.bfloat16
|
||||
DEVICE = "cuda"
|
||||
|
||||
|
||||
# ----- input fabrication -----------------------------------------------------
|
||||
|
||||
|
||||
def _round_up(x: int, base: int) -> int:
|
||||
return ((x + base - 1) // base) * base
|
||||
|
||||
|
||||
def _ceildiv(x: int, y: int) -> int:
|
||||
return (x + y - 1) // y
|
||||
|
||||
|
||||
def _assign_loras(num_tokens: int, num_sequences: int, max_loras: int) -> torch.Tensor:
|
||||
tokens_per_seq = num_tokens // num_sequences
|
||||
rem = num_tokens % num_sequences
|
||||
out = torch.empty(num_tokens, dtype=torch.int32)
|
||||
start = 0
|
||||
for i in range(num_sequences):
|
||||
end = start + tokens_per_seq + (1 if i < rem else 0)
|
||||
out[start:end] = random.randint(0, max_loras - 1)
|
||||
start = end
|
||||
return out
|
||||
|
||||
|
||||
def _assign_experts(num_tokens: int, num_experts: int, top_k: int):
|
||||
expert_indices = torch.empty((num_tokens, top_k), dtype=torch.int32)
|
||||
for i in range(num_tokens):
|
||||
expert_indices[i] = torch.randperm(num_experts)[:top_k]
|
||||
weights = torch.rand((num_tokens, top_k), dtype=torch.float32)
|
||||
weights = weights / weights.sum(dim=1, keepdim=True)
|
||||
return expert_indices, weights
|
||||
|
||||
|
||||
def _make_inputs(
|
||||
M: int,
|
||||
K: int,
|
||||
N_per_slice: int,
|
||||
rank: int,
|
||||
num_experts: int,
|
||||
top_k: int,
|
||||
max_loras: int,
|
||||
num_slices: int,
|
||||
block_size_m: int,
|
||||
):
|
||||
"""Mirrors the production caller's tensor layout."""
|
||||
torch.manual_seed(0)
|
||||
random.seed(0)
|
||||
|
||||
num_sequences = max(1, min(M, 8))
|
||||
topk_ids_cpu, topk_weights_cpu = _assign_experts(M, num_experts, top_k)
|
||||
token_lora_cpu = _assign_loras(M, num_sequences, max_loras)
|
||||
lora_ids_cpu = torch.full((max_loras + 1,), -1, dtype=torch.int32)
|
||||
uniq = torch.unique(token_lora_cpu, sorted=True)
|
||||
lora_ids_cpu[: uniq.size(0)].copy_(uniq)
|
||||
|
||||
topk_ids = topk_ids_cpu.to(DEVICE)
|
||||
topk_weights = topk_weights_cpu.to(device=DEVICE, dtype=DTYPE)
|
||||
token_lora_mapping = token_lora_cpu.to(DEVICE)
|
||||
lora_ids = lora_ids_cpu.to(DEVICE)
|
||||
adapter_enabled = torch.ones(max_loras + 1, dtype=torch.int32, device=DEVICE)
|
||||
|
||||
lora_a = [
|
||||
torch.randn((max_loras, num_experts, rank, K), dtype=DTYPE, device=DEVICE)
|
||||
/ max(K, 1) ** 0.5
|
||||
for _ in range(num_slices)
|
||||
]
|
||||
lora_b = [
|
||||
torch.randn(
|
||||
(max_loras, num_experts, N_per_slice, rank),
|
||||
dtype=DTYPE,
|
||||
device=DEVICE,
|
||||
)
|
||||
/ max(rank, 1) ** 0.5
|
||||
for _ in range(num_slices)
|
||||
]
|
||||
hidden = torch.randn((M, K), dtype=DTYPE, device=DEVICE)
|
||||
out_template = torch.zeros(
|
||||
(M, top_k, num_slices * N_per_slice), dtype=DTYPE, device=DEVICE
|
||||
)
|
||||
|
||||
# Sorted-path metadata (the prefill default).
|
||||
max_pad = topk_ids.numel() + num_experts * (block_size_m - 1)
|
||||
max_pad = _round_up(max_pad, block_size_m)
|
||||
max_blocks = _ceildiv(max_pad, block_size_m)
|
||||
sorted_token_ids = torch.empty(
|
||||
(max_loras * max_pad,), dtype=torch.int32, device=DEVICE
|
||||
)
|
||||
expert_ids = torch.empty(
|
||||
(max_loras * max_blocks,), dtype=torch.int32, device=DEVICE
|
||||
)
|
||||
num_post = torch.empty((max_loras,), dtype=torch.int32, device=DEVICE)
|
||||
ops.moe_lora_align_block_size(
|
||||
topk_ids,
|
||||
token_lora_mapping,
|
||||
num_experts,
|
||||
block_size_m,
|
||||
max_loras,
|
||||
max_pad,
|
||||
max_blocks,
|
||||
sorted_token_ids,
|
||||
expert_ids,
|
||||
num_post,
|
||||
adapter_enabled,
|
||||
lora_ids,
|
||||
)
|
||||
expert_ids = expert_ids.view(max_loras, -1).contiguous()
|
||||
sorted_token_ids = sorted_token_ids.view(max_loras, -1).contiguous()
|
||||
num_active = torch.tensor([max_loras + 1], dtype=torch.int32, device="cpu")
|
||||
|
||||
return dict(
|
||||
hidden=hidden,
|
||||
lora_a=lora_a,
|
||||
lora_b=lora_b,
|
||||
topk_weights=topk_weights,
|
||||
sorted_token_ids=sorted_token_ids,
|
||||
expert_ids=expert_ids,
|
||||
num_post=num_post,
|
||||
token_lora_mapping=token_lora_mapping,
|
||||
lora_ids=lora_ids,
|
||||
num_active=num_active,
|
||||
adapter_enabled=adapter_enabled,
|
||||
out_template=out_template,
|
||||
# bookkeeping
|
||||
M=M,
|
||||
K=K,
|
||||
N_per_slice=N_per_slice,
|
||||
rank=rank,
|
||||
num_experts=num_experts,
|
||||
top_k=top_k,
|
||||
max_loras=max_loras,
|
||||
num_slices=num_slices,
|
||||
block_size_m=block_size_m,
|
||||
)
|
||||
|
||||
|
||||
# ----- providers -------------------------------------------------------------
|
||||
|
||||
|
||||
def _run_one_shot(inp: dict):
|
||||
"""Drive `fused_moe_lora` with fully_sharded=False -> one-shot fast path."""
|
||||
out = inp["out_template"].clone()
|
||||
fused_moe_lora(
|
||||
out,
|
||||
inp["hidden"],
|
||||
inp["lora_a"],
|
||||
inp["lora_b"],
|
||||
inp["topk_weights"],
|
||||
inp["sorted_token_ids"],
|
||||
inp["expert_ids"],
|
||||
inp["num_post"],
|
||||
inp["token_lora_mapping"],
|
||||
inp["rank"],
|
||||
inp["top_k"],
|
||||
inp["lora_ids"],
|
||||
inp["num_active"],
|
||||
inp["adapter_enabled"],
|
||||
inp["block_size_m"],
|
||||
64,
|
||||
32,
|
||||
8,
|
||||
4,
|
||||
3,
|
||||
1,
|
||||
inp["block_size_m"],
|
||||
64,
|
||||
32,
|
||||
8,
|
||||
4,
|
||||
3,
|
||||
1,
|
||||
False,
|
||||
False,
|
||||
0,
|
||||
)
|
||||
return out
|
||||
|
||||
|
||||
def _run_two_kernel(inp: dict):
|
||||
"""Drive `fused_moe_lora_shrink` + `fused_moe_lora_expand` directly,
|
||||
bypassing the dispatch. Matches the legacy two-kernel work distribution.
|
||||
"""
|
||||
M = inp["M"]
|
||||
top_k = inp["top_k"]
|
||||
rank = inp["rank"]
|
||||
num_slices = inp["num_slices"]
|
||||
N_per_slice = inp["N_per_slice"]
|
||||
K = inp["K"]
|
||||
num_experts = inp["num_experts"]
|
||||
block_m = inp["block_size_m"]
|
||||
|
||||
intermediate = torch.zeros((num_slices, M, top_k, rank), dtype=DTYPE, device=DEVICE)
|
||||
out = inp["out_template"].clone()
|
||||
EM = inp["sorted_token_ids"].shape[1]
|
||||
num_tokens = M * top_k
|
||||
|
||||
fused_moe_lora_shrink(
|
||||
intermediate,
|
||||
inp["hidden"],
|
||||
inp["lora_a"],
|
||||
inp["topk_weights"],
|
||||
inp["sorted_token_ids"],
|
||||
inp["expert_ids"],
|
||||
inp["num_post"],
|
||||
inp["token_lora_mapping"],
|
||||
top_k,
|
||||
inp["lora_ids"],
|
||||
inp["adapter_enabled"],
|
||||
torch.device(DEVICE),
|
||||
rank,
|
||||
M,
|
||||
EM,
|
||||
K,
|
||||
num_tokens,
|
||||
num_experts,
|
||||
num_slices,
|
||||
block_m,
|
||||
64,
|
||||
32,
|
||||
8,
|
||||
4,
|
||||
3,
|
||||
1,
|
||||
inp["num_active"],
|
||||
False,
|
||||
)
|
||||
fused_moe_lora_expand(
|
||||
out,
|
||||
intermediate,
|
||||
inp["lora_b"],
|
||||
inp["topk_weights"],
|
||||
inp["sorted_token_ids"],
|
||||
inp["expert_ids"],
|
||||
inp["num_post"],
|
||||
inp["token_lora_mapping"],
|
||||
top_k,
|
||||
inp["lora_ids"],
|
||||
inp["adapter_enabled"],
|
||||
torch.device(DEVICE),
|
||||
rank,
|
||||
M,
|
||||
EM,
|
||||
K,
|
||||
num_tokens,
|
||||
num_experts,
|
||||
num_slices,
|
||||
rank,
|
||||
N_per_slice,
|
||||
block_m,
|
||||
64,
|
||||
32,
|
||||
8,
|
||||
4,
|
||||
3,
|
||||
1,
|
||||
inp["num_active"],
|
||||
False,
|
||||
0,
|
||||
)
|
||||
return out
|
||||
|
||||
|
||||
PROVIDER_FNS = {
|
||||
"one_shot": _run_one_shot,
|
||||
"two_kernel": _run_two_kernel,
|
||||
}
|
||||
|
||||
|
||||
# ----- model presets ---------------------------------------------------------
|
||||
|
||||
|
||||
MODEL_PRESETS: dict[str, dict] = {
|
||||
# Mixtral-8x7B style: E=8, top_k=2, hidden=4096, intermediate=14336
|
||||
"mixtral": dict(
|
||||
K=4096,
|
||||
N_per_slice=7168,
|
||||
num_experts=8,
|
||||
top_k=2,
|
||||
max_loras=4,
|
||||
num_slices=2,
|
||||
block_size_m=64,
|
||||
),
|
||||
# Qwen3-MoE / DeepSeek-V2 style: E=64, top_k=8, hidden=2048, inter=1408
|
||||
"qwen3moe": dict(
|
||||
K=2048,
|
||||
N_per_slice=1408,
|
||||
num_experts=64,
|
||||
top_k=8,
|
||||
max_loras=4,
|
||||
num_slices=2,
|
||||
block_size_m=64,
|
||||
),
|
||||
# GLM-5.1 (zai-org/GLM-5.1-FP8): E=256, top_k=8, hidden=6144,
|
||||
# moe_intermediate=2048
|
||||
"glm5_1": dict(
|
||||
K=6144,
|
||||
N_per_slice=2048,
|
||||
num_experts=256,
|
||||
top_k=8,
|
||||
max_loras=4,
|
||||
num_slices=2,
|
||||
block_size_m=64,
|
||||
),
|
||||
}
|
||||
|
||||
|
||||
M_RANGE = [16, 64, 256, 1024, 4096, 16384]
|
||||
RANK_RANGE = [8, 16, 32, 64]
|
||||
|
||||
|
||||
def get_benchmark(model: str, max_loras: int | None = None):
|
||||
preset = dict(MODEL_PRESETS[model])
|
||||
if max_loras is not None:
|
||||
preset["max_loras"] = max_loras
|
||||
|
||||
@triton.testing.perf_report(
|
||||
triton.testing.Benchmark(
|
||||
x_names=["M", "rank"],
|
||||
x_vals=[(M, R) for M in M_RANGE for R in RANK_RANGE],
|
||||
line_arg="provider",
|
||||
line_vals=list(PROVIDER_FNS.keys()),
|
||||
line_names=["one_shot (fused)", "two_kernel (legacy)"],
|
||||
styles=[("red", "-"), ("blue", "-")],
|
||||
ylabel="ms",
|
||||
plot_name=f"fused_moe_lora-{model}-loras{preset['max_loras']}",
|
||||
args={"preset": preset},
|
||||
)
|
||||
)
|
||||
def benchmark(M, rank, provider, preset):
|
||||
inp = _make_inputs(
|
||||
M=M,
|
||||
K=preset["K"],
|
||||
N_per_slice=preset["N_per_slice"],
|
||||
rank=rank,
|
||||
num_experts=preset["num_experts"],
|
||||
top_k=preset["top_k"],
|
||||
max_loras=preset["max_loras"],
|
||||
num_slices=preset["num_slices"],
|
||||
block_size_m=preset["block_size_m"],
|
||||
)
|
||||
fn = PROVIDER_FNS[provider]
|
||||
quantiles = [0.5, 0.2, 0.8]
|
||||
ms, min_ms, max_ms = triton.testing.do_bench(
|
||||
lambda: fn(inp), quantiles=quantiles
|
||||
)
|
||||
return ms, max_ms, min_ms
|
||||
|
||||
return benchmark
|
||||
|
||||
|
||||
# ----- correctness sanity ---------------------------------------------------
|
||||
|
||||
|
||||
def calculate_diff(model: str, M: int, rank: int, max_loras: int | None = None):
|
||||
preset = dict(MODEL_PRESETS[model])
|
||||
if max_loras is not None:
|
||||
preset["max_loras"] = max_loras
|
||||
inp = _make_inputs(
|
||||
M=M,
|
||||
K=preset["K"],
|
||||
N_per_slice=preset["N_per_slice"],
|
||||
rank=rank,
|
||||
num_experts=preset["num_experts"],
|
||||
top_k=preset["top_k"],
|
||||
max_loras=preset["max_loras"],
|
||||
num_slices=preset["num_slices"],
|
||||
block_size_m=preset["block_size_m"],
|
||||
)
|
||||
out_one = _run_one_shot(inp)
|
||||
out_two = _run_two_kernel(inp)
|
||||
max_abs = (out_one.float() - out_two.float()).abs().max().item()
|
||||
print(
|
||||
f" model={model:<9} M={M:<6} rank={rank:<3} "
|
||||
f"max|one_shot - two_kernel|={max_abs:.4g} "
|
||||
f"ref|max|={out_two.float().abs().max().item():.3g}"
|
||||
)
|
||||
if max_abs <= 5e-2:
|
||||
print(" ✅ outputs match within bf16 tolerance")
|
||||
else:
|
||||
print(" ❌ outputs differ beyond expected bf16 noise")
|
||||
|
||||
|
||||
# ----- main ------------------------------------------------------------------
|
||||
|
||||
|
||||
if __name__ == "__main__":
|
||||
parser = argparse.ArgumentParser()
|
||||
parser.add_argument(
|
||||
"--model",
|
||||
type=str,
|
||||
default="mixtral",
|
||||
choices=list(MODEL_PRESETS.keys()),
|
||||
help="Model preset to sweep",
|
||||
)
|
||||
parser.add_argument(
|
||||
"--save-path",
|
||||
type=str,
|
||||
default="./configs/fused_moe_lora_one_shot/",
|
||||
help="Directory to save benchmark results",
|
||||
)
|
||||
parser.add_argument(
|
||||
"--check-only",
|
||||
action="store_true",
|
||||
help="Run correctness sanity check only, no perf sweep",
|
||||
)
|
||||
parser.add_argument(
|
||||
"--max-loras",
|
||||
type=int,
|
||||
default=None,
|
||||
help="Override max_loras in the model preset (number of LoRA adapters "
|
||||
"active in the batch). Defaults to the preset's value.",
|
||||
)
|
||||
args = parser.parse_args()
|
||||
|
||||
print(f"Correctness check ({args.model}):")
|
||||
calculate_diff(args.model, M=256, rank=32, max_loras=args.max_loras)
|
||||
if args.check_only:
|
||||
raise SystemExit(0)
|
||||
|
||||
effective_max_loras = (
|
||||
args.max_loras
|
||||
if args.max_loras is not None
|
||||
else MODEL_PRESETS[args.model]["max_loras"]
|
||||
)
|
||||
print(f"\nGPU: {torch.cuda.get_device_name()}")
|
||||
print(f"Model preset: {args.model} max_loras={effective_max_loras}\n")
|
||||
benchmark = get_benchmark(args.model, max_loras=args.max_loras)
|
||||
os.makedirs(args.save_path, exist_ok=True)
|
||||
benchmark.run(print_data=True, save_path=args.save_path)
|
||||
Some files were not shown because too many files have changed in this diff Show More
Reference in New Issue
Block a user