Compare commits

..
Author SHA1 Message Date
dependabot[bot]andGitHub f36d57ec34 Bump actions/setup-python from 6.1.0 to 6.2.0
Bumps [actions/setup-python](https://github.com/actions/setup-python) from 6.1.0 to 6.2.0.
- [Release notes](https://github.com/actions/setup-python/releases)
- [Commits](https://github.com/actions/setup-python/compare/83679a892e2d95755f2dac6acb0bfd1e9ac5d548...a309ff8b426b58ec0e2a45f0f869d46889d02405)

---
updated-dependencies:
- dependency-name: actions/setup-python
  dependency-version: 6.2.0
  dependency-type: direct:production
  update-type: version-update:semver-minor
...

Signed-off-by: dependabot[bot] <support@github.com>
2026-06-02 18:31:03 +00:00
1908 changed files with 57303 additions and 180741 deletions
-23
View File
@@ -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"
+34 -65
View File
@@ -1,73 +1,42 @@
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
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
- >
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 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
- |
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
'
python3 - <<PY
import torch, vllm
print(torch.__version__)
print(vllm.__version__)
PY
echo AMD image smoke OK
'
- 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"
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
+14 -15
View File
@@ -28,19 +28,18 @@ steps:
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"
# 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: []
@@ -91,7 +90,7 @@ steps:
- 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"
@@ -136,7 +135,7 @@ steps:
- |
bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 45m "
pytest -x -v -s tests/models/multimodal/generation --ignore=tests/models/multimodal/generation/test_pixtral.py -m cpu_model --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB"
parallelism: 4
parallelism: 3
- label: "Arm CPU Test"
depends_on: []
@@ -1,80 +0,0 @@
group: Intel
steps:
- label: ":docker: Build XPU image"
soft_fail: true
optional: true
depends_on: []
key: image-build-xpu
commands:
- bash -lc '.buildkite/image_build/image_build_xpu.sh "public.ecr.aws/q9t5s3a7" "vllm-ci-test-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
- label: "XPU example Test"
depends_on:
- image-build-xpu
timeout_in_minutes: 30
optional: true
device: intel_gpu
agent_tags:
label: production
gpu: 2+
mem: 24+
no_plugin: true
env:
REGISTRY: "public.ecr.aws/q9t5s3a7"
REPO: "vllm-ci-test-repo"
source_file_dependencies:
- .buildkite/hardware_tests/intel_xpu_ci/test-intel.yaml
- .buildkite/scripts/hardware_ci/run-intel-ci-test.sh
commands:
- >-
bash .buildkite/scripts/hardware_ci/run-intel-test.sh
'bash .buildkite/scripts/hardware_ci/run-intel-ci-test.sh example'
- label: "XPU V1 test"
depends_on:
- image-build-xpu
timeout_in_minutes: 30
optional: true
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
env:
REGISTRY: "public.ecr.aws/q9t5s3a7"
REPO: "vllm-ci-test-repo"
source_file_dependencies:
- .buildkite/hardware_tests/intel_xpu_ci/test-intel.yaml
- .buildkite/scripts/hardware_ci/run-intel-ci-test.sh
commands:
- >-
bash .buildkite/scripts/hardware_ci/run-intel-test.sh
'bash .buildkite/scripts/hardware_ci/run-intel-ci-test.sh v1'
- label: "XPU server test"
depends_on:
- image-build-xpu
timeout_in_minutes: 30
optional: true
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
env:
REGISTRY: "public.ecr.aws/q9t5s3a7"
REPO: "vllm-ci-test-repo"
source_file_dependencies:
- .buildkite/hardware_tests/intel_xpu_ci/test-intel.yaml
- .buildkite/scripts/hardware_ci/run-intel-ci-test.sh
commands:
- >-
bash .buildkite/scripts/hardware_ci/run-intel-test.sh
'bash .buildkite/scripts/hardware_ci/run-intel-ci-test.sh server'
+2 -3
View File
@@ -21,13 +21,12 @@ else
exit 0
fi
# build for arm64 GPU targets: Grace/GH200 (sm_90) and DGX Spark/GB10
# (sm_121, family-covered by 12.0 under CUDA 13)
# 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 12.0" \
--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 \
@@ -1,26 +0,0 @@
group: Basic Correctness
depends_on:
- image-build-xpu
steps:
- label: XPU Sleep Mode
timeout_in_minutes: 30
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
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'
-4
View File
@@ -5,10 +5,6 @@ steps:
- label: Engine (1 GPU)
timeout_in_minutes: 30
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
working_dir: "."
env:
@@ -1,27 +0,0 @@
group: Expert Parallelism
depends_on:
- image-build-xpu
steps:
- label: EPLB Algorithm
key: eplb-algorithm
timeout_in_minutes: 45
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
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'
-4
View File
@@ -5,10 +5,6 @@ steps:
- label: vLLM IR Tests
timeout_in_minutes: 30
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
working_dir: "."
env:
-24
View File
@@ -5,10 +5,6 @@ steps:
- label: LoRA Runtime + Utils
timeout_in_minutes: 45
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 24+
no_plugin: true
working_dir: "."
env:
@@ -38,10 +34,6 @@ steps:
- label: LoRA Fused/MoE Kernels
timeout_in_minutes: 45
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
working_dir: "."
env:
@@ -62,10 +54,6 @@ steps:
- label: LoRA Punica Kernels
timeout_in_minutes: 45
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
working_dir: "."
env:
@@ -86,10 +74,6 @@ steps:
- label: LoRA Punica FP8/XPU Ops
timeout_in_minutes: 45
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
working_dir: "."
env:
@@ -110,10 +94,6 @@ steps:
- label: LoRA Models
timeout_in_minutes: 45
device: intel_gpu
agent_tags:
label: production
gpu: 2+
mem: 24+
no_plugin: true
working_dir: "."
env:
@@ -137,10 +117,6 @@ steps:
- label: LoRA Multimodal
timeout_in_minutes: 45
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
working_dir: "."
env:
+9 -168
View File
@@ -5,10 +5,6 @@ steps:
- label: V1 Core + KV + Metrics
timeout_in_minutes: 30
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
working_dir: "."
env:
@@ -35,54 +31,6 @@ steps:
- label: V1 Sample + Logits
timeout_in_minutes: 30
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
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/logger.py
- vllm/model_executor/
- vllm/platforms/
- vllm/sampling_params.py
- vllm/transformers_utils/
- vllm/utils/
- vllm/v1/
- tests/v1/sample
- tests/v1/logits_processors
- tests/v1/test_oracle.py
- tests/v1/test_request.py
- tests/v1/test_outputs.py
commands:
- >-
bash .buildkite/scripts/hardware_ci/run-intel-test.sh
'pip install lm_eval[api]>=0.4.12 &&
export VLLM_WORKER_MULTIPROC_METHOD=spawn &&
cd tests &&
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 &&
pytest -v -s v1/sample/test_logprobs.py &&
pytest -v -s v1/sample/test_logprobs_e2e.py'
- label: XPU CPU Offload
timeout_in_minutes: 60
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
working_dir: "."
env:
@@ -91,124 +39,17 @@ steps:
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
- tests/v1/sample
- tests/v1/logits_processors
- tests/v1/test_oracle.py
- tests/v1/test_request.py
- tests/v1/test_outputs.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
agent_tags:
label: production
gpu: 1+
mem: 16+
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
agent_tags:
label: production
gpu: 2+
mem: 16+
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
agent_tags:
label: production
gpu: 1+
mem: 24+
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/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'
@@ -1,62 +0,0 @@
group: Model Runner V2 Intel
depends_on:
- image-build-xpu
steps:
- label: Model Runner V2 Core Tests (Intel)
timeout_in_minutes: 45
device: intel_gpu
agent_tags:
label: production
gpu: 2+
mem: 16+
no_plugin: true
working_dir: "."
env:
REGISTRY: "public.ecr.aws/q9t5s3a7"
REPO: "vllm-ci-test-repo"
VLLM_TEST_DEVICE: "xpu"
source_file_dependencies:
- vllm/v1/worker/gpu/
- vllm/v1/worker/gpu_worker.py
- vllm/v1/core/sched/
- vllm/v1/attention/
- tests/v1/engine/test_llm_engine.py
- tests/v1/e2e/
commands:
- >-
bash .buildkite/scripts/hardware_ci/run-intel-test.sh
'export VLLM_USE_V2_MODEL_RUNNER=1 &&
cd tests &&
pytest -v -s v1/engine/test_llm_engine.py -k "not test_engine_metrics" &&
ENFORCE_EAGER=1 pytest -v -s v1/e2e/general/test_async_scheduling.py -k "not ngram" &&
pytest -v -s v1/e2e/general/test_min_tokens.py'
- label: Model Runner V2 Examples (Intel)
timeout_in_minutes: 45
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 24+
no_plugin: true
working_dir: "."
env:
REGISTRY: "public.ecr.aws/q9t5s3a7"
REPO: "vllm-ci-test-repo"
VLLM_TEST_DEVICE: "xpu"
source_file_dependencies:
- vllm/v1/worker/gpu/
- vllm/v1/core/sched/
- vllm/v1/worker/gpu_worker.py
- examples/basic/offline_inference/
- examples/generate/multimodal/
- examples/features/
commands:
- >-
bash .buildkite/scripts/hardware_ci/run-intel-test.sh
'export VLLM_USE_V2_MODEL_RUNNER=1 &&
cd examples &&
python3 basic/offline_inference/chat.py &&
python3 basic/offline_inference/generate.py --model facebook/opt-125m &&
python3 generate/multimodal/vision_language_offline.py --seed 0 &&
python3 features/automatic_prefix_caching/prefix_caching_offline.py'
@@ -1,131 +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
agent_tags:
label: production
gpu: 1+
mem: 16+
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
agent_tags:
label: production
gpu: 1+
mem: 16+
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
agent_tags:
label: production
gpu: 1+
mem: 24+
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
agent_tags:
label: production
gpu: 1+
mem: 16+
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
agent_tags:
label: production
gpu: 1+
mem: 16+
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
+3 -39
View File
@@ -19,10 +19,6 @@ steps:
- image-build-xpu
timeout_in_minutes: 30
device: intel_gpu
agent_tags:
label: production
gpu: 2+
mem: 24+
no_plugin: true
env:
REGISTRY: "public.ecr.aws/q9t5s3a7"
@@ -44,19 +40,13 @@ steps:
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 superjob/Qwen3-4B-Instruct-2507-GPTQ-Int4 --max-model-len 8192
'
- label: "XPU V1 test"
depends_on:
- image-build-xpu
timeout_in_minutes: 30
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
env:
REGISTRY: "public.ecr.aws/q9t5s3a7"
@@ -68,24 +58,19 @@ steps:
- >-
bash .buildkite/scripts/hardware_ci/run-intel-test.sh
'cd tests &&
bash v1/kv_connector/nixl_integration/run_xpu_disagg_accuracy_test.sh &&
pytest -v -s v1/core --ignore=v1/core/test_reset_prefix_cache_e2e.py --ignore=v1/core/test_scheduler_e2e.py &&
pytest -v -s v1/engine --ignore=v1/engine/test_output_processor.py &&
pytest -v -s v1/sample --ignore=v1/sample/test_logprobs.py --ignore=v1/sample/test_logprobs_e2e.py -k "not test_topk_only and not test_topp_only and not test_topk_and_topp" &&
pytest -v -s v1/worker --ignore=v1/worker/test_gpu_model_runner.py --ignore=v1/worker/test_worker_memory_snapshot.py &&
pytest -v -s v1/structured_output &&
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 --ignore=v1/spec_decode/test_speculators_correctness.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
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
env:
REGISTRY: "public.ecr.aws/q9t5s3a7"
@@ -98,26 +83,5 @@ steps:
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 entrypoints/openai/chat_completion/test_audio_in_video.py &&
pytest -v -s benchmarks/test_serve_cli.py'
- label: "XPU quantization test"
depends_on:
- image-build-xpu
timeout_in_minutes: 30
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
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
'cd tests &&
pytest -v -s quantization/test_auto_round.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,
@@ -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,
+4 -16
View File
@@ -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"
@@ -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
-3
View File
@@ -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"
+47 -151
View File
@@ -1,178 +1,74 @@
#!/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>.
#
# Default output path: ci-<build>-<uuid_first_13_chars>.log (e.g.
# ci-68478-019e6b07-daae.log). Jobs in the same build share the UUID's
# first 8 chars, so the second segment is needed for uniqueness when
# fetching multiple jobs in parallel. The script refuses to overwrite an
# existing output file; pass an explicit path or set CI_FETCH_LOG_FORCE=1
# to override.
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)
JOB=$(echo "$1" | grep -oE '[0-9a-f]{8}-[0-9a-f-]+' | head -n 1)
OUT="${2:-}"
[ -n "$BUILD" ] || die "Could not parse build number from: $1"
;;
[0-9]*)
[ $# -ge 2 ] || usage
else
if [ $# -lt 2 ]; then usage; fi
BUILD="$1"
JOB="$2"
OUT="${3:-}"
;;
*)
fi
if [ -z "$BUILD" ] || [ -z "$JOB" ]; then
echo "Could not parse build number or job UUID from: $1" >&2
usage
;;
esac
fi
# Jobs in the same build share the UUID's first segment, so include the
# second segment (chars 9-13, e.g. "019e6b07-daae") to keep default filenames
# unique when fetching multiple jobs from one build in parallel.
if [ -z "$OUT" ]; then
OUT="ci-${BUILD}-${JOB:0:13}.log"
fi
if [ -e "$OUT" ] && [ -z "${CI_FETCH_LOG_FORCE:-}" ]; then
echo "Refusing to overwrite existing $OUT (set CI_FETCH_LOG_FORCE=1 or pass an explicit output path)." >&2
exit 1
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"
+15 -149
View File
@@ -28,10 +28,8 @@
###############################################################################
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
@@ -54,108 +52,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
@@ -347,30 +243,22 @@ report_docker_usage
# --- 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 run uses --rm, so the container is normally already gone when the
# EXIT trap runs. Cleanup is best-effort and must not affect the test result.
docker rm -f "${container_name}" >/dev/null 2>&1 || 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
on_exit() {
local exit_code=$?
remove_docker_container
exit "$exit_code"
}
trap on_exit EXIT
# --- Prepare commands ---
echo "--- Running container"
@@ -379,14 +267,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 +306,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 +383,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 +390,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"
@@ -1,51 +0,0 @@
#!/bin/bash
set -euo pipefail
test_suite="${1:-}"
if [[ -z "${test_suite}" ]]; then
echo "Usage: $0 <example|v1|server>" >&2
exit 1
fi
case "${test_suite}" in
example)
pip install tblib==3.1.0
python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --block-size 64 --enforce-eager
python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --block-size 64 -O3 -cc.cudagraph_mode=NONE
python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --block-size 64 --enforce-eager -tp 2 --distributed-executor-backend mp
python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --block-size 64 --enforce-eager --attention-backend=TRITON_ATTN
python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --block-size 64 --enforce-eager --quantization fp8
python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --block-size 64 --enforce-eager --kv-cache-dtype fp8
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
;;
v1)
cd tests
pytest -v -s v1/core --ignore=v1/core/test_reset_prefix_cache_e2e.py --ignore=v1/core/test_scheduler_e2e.py
pytest -v -s v1/engine --ignore=v1/engine/test_output_processor.py
pytest -v -s v1/sample --ignore=v1/sample/test_logprobs.py --ignore=v1/sample/test_logprobs_e2e.py -k "not test_topk_only and not test_topp_only and not test_topk_and_topp"
pytest -v -s v1/worker --ignore=v1/worker/test_gpu_model_runner.py --ignore=v1/worker/test_worker_memory_snapshot.py
pytest -v -s v1/structured_output
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 --ignore=v1/spec_decode/test_speculators_correctness.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
;;
server)
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
;;
*)
echo "Unknown Intel test suite: ${test_suite}" >&2
exit 1
;;
esac
@@ -243,10 +243,8 @@ container_name="xpu_${BUILDKITE_COMMIT}_$(tr -dc A-Za-z0-9 < /dev/urandom | head
# ---- Command source selection ----
commands=""
commands_source=""
if [[ -n "${VLLM_TEST_COMMANDS:-}" ]]; then
commands="${VLLM_TEST_COMMANDS}"
commands_source="env"
echo "Commands sourced from VLLM_TEST_COMMANDS (quoting preserved)"
elif [[ $# -gt 0 ]]; then
all_yaml=true
@@ -305,12 +303,8 @@ if [[ -z "$commands" ]]; then
fi
echo "Raw commands: $commands"
if [[ "$commands_source" != "env" ]]; then
commands=$(re_quote_pytest_markers "$commands")
echo "After re-quoting: $commands"
else
echo "Skipping re-quoting for VLLM_TEST_COMMANDS input"
fi
commands=$(re_quote_pytest_markers "$commands")
echo "After re-quoting: $commands"
commands=$(apply_intel_test_overrides "$commands")
echo "Final commands: $commands"
@@ -330,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
}
@@ -346,12 +357,9 @@ 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..."
if ! docker image inspect "${IMAGE}" >/dev/null 2>&1; then
echo 'Image missing before container creation, pulling again...'
timeout 900 docker pull "${IMAGE}"
echo "Pull step completed"
fi
docker create \
@@ -364,8 +372,6 @@ export HF_TOKEN ZE_AFFINITY_MASK
--entrypoint='' \
-e HF_TOKEN \
-e ZE_AFFINITY_MASK \
-e BUILDKITE_PARALLEL_JOB \
-e BUILDKITE_PARALLEL_JOB_COUNT \
-e CMDS \
--name "${container_name}" \
"${IMAGE}" \
@@ -4,11 +4,6 @@
set -euo pipefail
if python3 -c "import torch; raise SystemExit(0 if torch.version.hip is not None else 1)"; then
uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
exit 0
fi
REQUIREMENTS_FILE="${KV_CONNECTORS_REQUIREMENTS:-/vllm-workspace/requirements/kv_connectors.txt}"
uv pip install --system -r "${REQUIREMENTS_FILE}"
@@ -110,36 +110,6 @@ install_uv() {
| 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
@@ -162,7 +132,6 @@ run_style_clippy() {
run_tests() {
install_uv
setup_pyo3_python
install_cargo_nextest
log_section "Running cargo nextest"
+462 -376
View File
File diff suppressed because it is too large Load Diff
+2 -16
View File
@@ -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
+2 -2
View File
@@ -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
+1 -1
View File
@@ -23,4 +23,4 @@ steps:
- benchmarks/attention_benchmarks/
- vllm/v1/attention/
commands:
- python3 benchmarks/attention_benchmarks/benchmark.py --backends flash flashinfer --batch-specs "8q1s1k"
- python3 benchmarks/attention_benchmarks/benchmark.py --backends flash flashinfer --batch-specs "8q1s1k" --repeats 1 --warmup-iters 1
-14
View File
@@ -61,20 +61,6 @@ steps:
- bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
- 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
-16
View File
@@ -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
@@ -74,16 +68,6 @@ steps:
- tests/v1/e2e/general/
commands:
- pytest -v -s v1/e2e/general --ignore v1/e2e/general/test_async_scheduling.py
mirror:
amd:
device: mi250_1
timeout_in_minutes: 35
depends_on:
- image-build-amd
source_file_dependencies:
- vllm/v1/
- tests/v1/e2e/general/
- vllm/platforms/rocm.py
- label: V1 e2e (2 GPUs)
key: v1-e2e-2-gpus
+82 -106
View File
@@ -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/offline_mode --ignore=entrypoints/openai --ignore=entrypoints/serve --ignore=entrypoints/test_chat_utils.py --ignore=entrypoints/pooling --ignore=entrypoints/speech_to_text --ignore=tests/entrypoints/generate
- label: Entrypoints Integration (LLM)
key: entrypoints-integration-llm
@@ -21,20 +20,84 @@ 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
soft_fail: true
depends_on:
- image-build-amd
- label: Entrypoints Integration (API Server)
key: entrypoints-integration-api-server
- 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"
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_oot_registration.py
mirror:
amd:
device: mi325_1
soft_fail: true
timeout_in_minutes: 80
depends_on:
- image-build-amd
- 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"
source_file_dependencies:
- vllm/
- tests/entrypoints/openai
- tests/entrypoints/test_chat_utils
- tests/entrypoints/generate
- tests/tool_use
commands:
- pytest -v -s entrypoints/openai/completion --ignore=entrypoints/openai/completion/test_tensorizer_entrypoint.py
- pytest -v -s entrypoints/test_chat_utils.py
- pytest -v -s entrypoints/generate
- pytest -v -s tool_use
mirror:
amd:
device: mi325_1
soft_fail: true
timeout_in_minutes: 60
depends_on:
- image-build-amd
- 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
- 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
mirror:
amd:
device: mi325_1
soft_fail: true
timeout_in_minutes: 60
depends_on:
- image-build-amd
- label: Entrypoints Integration (API Server 2)
device: h200_35gb
key: entrypoints-integration-api-server-2
timeout_in_minutes: 130
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
@@ -47,78 +110,10 @@ steps:
mirror:
amd:
device: mi325_1
soft_fail: true
depends_on:
- image-build-amd
- 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"
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
mirror:
amd:
device: mi325_1
timeout_in_minutes: 80
depends_on:
- image-build-amd
- 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"
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
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
timeout_in_minutes: 60
depends_on:
- image-build-amd
- 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: Entrypoints Integration (Speech to Text)
device: h200_35gb
key: entrypoints-integration-speech_to_text
@@ -131,18 +126,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 +137,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
@@ -163,20 +156,3 @@ steps:
- vllm/entrypoints/openai/
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/
+2 -50
View File
@@ -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
@@ -47,10 +46,8 @@ steps:
- 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
- tests/kernels/test_top_k_per_row.py # it runs on Blackwell too - some kernels have arch-specific optimizations
commands:
- pytest -v -s kernels/test_fused_deepseek_v4_*.py
- pytest -v -s kernels/test_top_k_per_row.py
- label: Deepseek V4 Kernel Test (B200)
key: deepseek-v4-kernel-test-b200
@@ -76,33 +73,6 @@ steps:
commands:
- pytest -v -s kernels/attention --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
parallelism: 2
mirror:
amd:
device: mi325_1
timeout_in_minutes: 55
depends_on:
- image-build-amd
source_file_dependencies:
- csrc/attention/
- vllm/v1/attention
- vllm/model_executor/layers/attention
- tests/kernels/attention
- vllm/_aiter_ops.py
- vllm/envs.py
- vllm/platforms/rocm.py
- 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
@@ -120,7 +90,6 @@ steps:
source_file_dependencies:
- csrc/quantization/
- vllm/model_executor/layers/quantization
- vllm/config/
- tests/kernels/quantization
- tests/kernels/quantization/test_rocm_skinny_gemms.py
- vllm/_aiter_ops.py
@@ -144,22 +113,6 @@ steps:
- pytest -v -s kernels/moe --ignore=kernels/moe/test_modular_oai_triton_moe.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
- pytest -v -s kernels/moe/test_modular_oai_triton_moe.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
parallelism: 5
mirror:
amd:
device: mi325_1
timeout_in_minutes: 50
source_file_dependencies:
- csrc/quantization/cutlass_w8a8/moe/
- csrc/moe/
- tests/kernels/moe
- vllm/model_executor/layers/fused_moe/
- vllm/distributed/device_communicators/
- vllm/envs.py
- vllm/config
- vllm/_aiter_ops.py
- vllm/platforms/rocm.py
depends_on:
- image-build-amd
- label: Kernels Mamba Test
key: kernels-mamba-test
@@ -270,7 +223,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/
@@ -346,4 +299,3 @@ steps:
- vllm/config
commands:
- pytest -v -s kernels/moe/test_moe_layer.py
- pytest -v -s kernels/moe/test_deepep_v2_moe.py
-82
View File
@@ -12,21 +12,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
@@ -101,16 +86,6 @@ steps:
num_devices: 8
commands:
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-h200.txt
mirror:
amd:
device: mi300_8
timeout_in_minutes: 180
depends_on:
- image-build-amd
commands:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- export PYTORCH_ROCM_ARCH=gfx942 # Limit Quark compilation to save time
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-mi3xx.txt
- label: MoE Refactor Integration Test (H100 - TEMPORARY)
key: moe-refactor-integration-test-h100-temporary
@@ -136,49 +111,6 @@ steps:
commands:
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/moe-refactor-dp-ep/config-b200.txt
- label: LM Eval Humming (A100 - TEMPORARY)
key: lm-eval-humming-a100
timeout_in_minutes: 30
device: a100
optional: true
num_devices: 1
source_file_dependencies:
- vllm/model_executor/layers/quantization/humming.py
- vllm/model_executor/layers/quantization/utils/humming_utils.py
- vllm/model_executor/layers/fused_moe/experts/fused_humming_moe.py
- vllm/model_executor/layers/fused_moe/oracle/mxfp4.py
commands:
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config.txt
- label: LM Eval Humming (H100 - TEMPORARY)
key: lm-eval-humming-h100
timeout_in_minutes: 30
device: h100
optional: true
num_devices: 1
source_file_dependencies:
- vllm/model_executor/layers/quantization/humming.py
- vllm/model_executor/layers/quantization/utils/humming_utils.py
- vllm/model_executor/layers/fused_moe/experts/fused_humming_moe.py
- vllm/model_executor/layers/fused_moe/oracle/mxfp4.py
commands:
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config.txt
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config-act-fp8.txt
- label: LM Eval Humming (B200 - TEMPORARY)
key: lm-eval-humming-b200
timeout_in_minutes: 30
device: b200-k8s
optional: true
num_devices: 1
source_file_dependencies:
- vllm/model_executor/layers/quantization/humming.py
- vllm/model_executor/layers/quantization/utils/humming_utils.py
- vllm/model_executor/layers/fused_moe/experts/fused_humming_moe.py
- vllm/model_executor/layers/fused_moe/oracle/mxfp4.py
commands:
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config.txt
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config-act-fp8.txt
- label: LM Eval TurboQuant KV Cache
key: lm-eval-turboquant-kv-cache
@@ -220,20 +152,6 @@ steps:
- uv pip install --system 'gpt-oss[eval]==0.0.5'
- pytest -s -v evals/gpt_oss/test_gpqa_correctness.py --config-list-file=configs/models-b200.txt
- label: GPQA Eval (GPT-OSS) (DGX Spark)
key: gpqa-eval-gpt-oss-spark
timeout_in_minutes: 120
device: dgx-spark
optional: true
num_devices: 1
source_file_dependencies:
- csrc/
- vllm/model_executor/layers/quantization
- tests/evals/gpt_oss/
commands:
- uv pip install --system 'gpt-oss[eval]==0.0.5'
- pytest -s -v evals/gpt_oss/test_gpqa_correctness.py --config-list-file=configs/models-spark.txt
- label: MRCR Eval Small Models
device: h200_35gb
timeout_in_minutes: 30
-11
View File
@@ -12,17 +12,6 @@ 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
parallelism: 4
mirror:
amd:
device: mi325_1
working_dir: "/vllm-workspace/tests"
timeout_in_minutes: 60
source_file_dependencies:
- vllm/lora
- tests/lora
- vllm/platforms/rocm.py
depends_on:
- image-build-amd
- label: LoRA TP (Distributed)
+2 -32
View File
@@ -21,12 +21,6 @@ steps:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
# TODO: create another `optional` test group for slow tests
- pytest -v -s -m 'not slow_test' v1/spec_decode
mirror:
amd:
device: mi300_1
timeout_in_minutes: 65
depends_on:
- image-build-amd
- label: V1 Sample + Logits
key: v1-sample-logits
@@ -105,12 +99,6 @@ steps:
# Integration test for streaming correctness (requires special branch).
- pip install -U git+https://github.com/robertgshaw2-redhat/lm-evaluation-harness.git@streaming-api
- pytest -v -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine
mirror:
amd:
device: mi325_1
timeout_in_minutes: 60
depends_on:
- image-build-amd
- label: V1 Others (CPU)
key: v1-others-cpu
@@ -150,26 +138,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
@@ -308,7 +281,6 @@ steps:
- vllm/model_executor/layers/quantization/quark/
- vllm/multimodal/
- vllm/outputs.py
- vllm/parser/
- vllm/platforms/
- vllm/pooling_params.py
- vllm/ray/
@@ -320,7 +292,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
@@ -328,25 +299,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
@@ -94,13 +94,11 @@ 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
@@ -51,7 +51,6 @@ steps:
torch_nightly: {}
amd:
device: mi325_1
timeout_in_minutes: 90
depends_on:
- image-build-amd
commands:
+2 -16
View File
@@ -30,6 +30,7 @@ steps:
- pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
- pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "qwen3 or gemma"
- pytest -v -s models/multimodal/generation/test_qwen2_5_vl.py -m core_model
- pytest -v -s models/multimodal/generation/test_vit_cudagraph.py -m core_model
mirror:
amd:
device: mi325_1
@@ -62,15 +63,9 @@ steps:
- tests/models/multimodal
commands:
- pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
- 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/generation/test_vit_cudagraph.py --ignore models/multimodal/processing
- pytest -v -s models/multimodal/generation/test_vit_cudagraph.py -m core_model
- 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
- pytest 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 # Otherwise, mp_method="spawn" doesn't work
mirror:
amd:
device: mi325_1
depends_on:
- image-build-amd
- label: Multi-Modal Processor (CPU)
key: multi-modal-processor-cpu
@@ -158,12 +153,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
+3 -21
View File
@@ -27,10 +27,6 @@ steps:
- pip install -e ./plugins/bge_m3_sparse_plugin
- pytest -v -s plugins_tests/test_bge_m3_sparse_io_processor_plugins.py
- pip uninstall bge_m3_sparse_plugin -y
# test colbert_query io_processor plugin
- pip install -e ./plugins/colbert_query_plugin
- pytest -v -s plugins_tests/test_colbert_query_io_processor_plugins.py
- pip uninstall colbert_query_plugin -y
# end io_processor plugins test
# begin stat_logger plugins test
- pip install -e ./plugins/vllm_add_dummy_stat_logger
@@ -41,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
-6
View File
@@ -107,12 +107,6 @@ steps:
- tests/compile/passes
commands:
- pytest -s -v compile/passes --ignore compile/passes/distributed
mirror:
amd:
device: mi300_1
timeout_in_minutes: 180
depends_on:
- image-build-amd
- label: PyTorch Fullgraph Smoke Test
key: pytorch-fullgraph-smoke-test
-12
View File
@@ -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
-4
View File
@@ -99,13 +99,9 @@ steps:
- vllm/v1/engine/
- vllm/v1/worker/
- tests/utils.py
- tests/v1/distributed/test_external_lb_dp.py
- tests/v1/distributed/test_hybrid_lb_dp.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"
- TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_external_lb_dp.py -k "not 4 and not server_info"
- TP_SIZE=1 DP_SIZE=4 pytest -v -s v1/distributed/test_hybrid_lb_dp.py -k "not 4 and not server_info"
-57
View File
@@ -12,20 +12,6 @@ steps:
- tests/v1/e2e/spec_decode/
commands:
- pytest -v -s v1/e2e/spec_decode -k "eagle_correctness"
mirror:
amd:
device: mi325_1
timeout_in_minutes: 45
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 Eagle Nightly B200
key: spec-decode-eagle-nightly-b200
@@ -51,21 +37,6 @@ steps:
- 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
@@ -90,20 +61,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
@@ -115,20 +72,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
@@ -1,35 +0,0 @@
---
name: ci-fails-buildkite
description: Fetch and diagnose vLLM Buildkite CI failure logs. Use when investigating failing CI jobs on a PR or build, when the user pastes a buildkite.com URL, or asks to fetch/diagnose CI logs.
---
# Diagnosing vLLM Buildkite CI Failures
Buildkite logs are public; no login needed.
`.buildkite/scripts/ci-fetch-log.sh` saves each log as `ci-<build>-<job-name>.log`, stripped of timestamps and ANSI codes. Existing files are kept; set `CI_FETCH_LOG_FORCE=1` to refetch.
## Fetching logs
```bash
# All failed jobs in a PR's latest build (current branch's PR if omitted):
.buildkite/scripts/ci-fetch-log.sh --pr <PR>
# All failed jobs in a build (--soft also includes soft-failed jobs;
# --all fetches every finished job):
.buildkite/scripts/ci-fetch-log.sh "https://buildkite.com/vllm/ci/builds/<N>"
# One job — `gh pr checks` URLs (#<job_uuid>) and web UI URLs (?sid=) both
# work; pass "-" as a second argument to stream to stdout:
.buildkite/scripts/ci-fetch-log.sh "https://buildkite.com/vllm/ci/builds/<N>#<job_uuid>"
```
To clean an already-downloaded log with `.buildkite/scripts/ci-clean-log.sh`:
```bash
./ci-clean-log.sh ci.log
```
## Reference
See [docs/contributing/ci/failures.md](../../../docs/contributing/ci/failures.md) for the full guide: filing CI failure issues, investigating/bisecting, reproducing flaky tests, and daily triage.
+16 -11
View File
@@ -2,14 +2,15 @@
# for more info about CODEOWNERS file
# This lists cover the "core" components of vLLM that require careful review
/vllm/compilation @zou3519 @youkaichao @ProExpertProg @BoyuanFeng
/vllm/compilation @zou3519 @youkaichao @ProExpertProg @BoyuanFeng @vadiklyutiy
/vllm/distributed/kv_transfer @NickLucche @ApostaC @orozery @xuechendi
/vllm/lora @jeejeelee
/vllm/model_executor/layers/attention @LucasWilkinson @MatthewBonanni
/vllm/model_executor/layers/fused_moe @mgoin @pavanimajety @zyongye
/vllm/model_executor/layers/quantization @mgoin @robertgshaw2-redhat @tlrmchlsmth @yewentao256 @pavanimajety @zyongye
/vllm/model_executor/layers/mamba @tdoublep @tomeras91
/vllm/model_executor/layers/mamba/gdn/qwen_gdn_linear_attn.py @tdoublep @ZJY0516 @vadiklyutiy
/vllm/model_executor/layers/mamba/gdn_linear_attn.py @tdoublep @ZJY0516 @vadiklyutiy
/vllm/model_executor/layers/rotary_embedding.py @vadiklyutiy
/vllm/model_executor/model_loader @22quinn
/vllm/model_executor/layers/batch_invariant.py @yewentao256
/vllm/ir @ProExpertProg
@@ -22,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
@@ -38,11 +34,10 @@
/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
@@ -79,7 +74,7 @@
/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
@@ -119,6 +114,16 @@
/vllm/model_executor/models/transformers @hmellor
/tests/models/test_transformers.py @hmellor
# Observability
/vllm/config/observability.py @markmc
/vllm/v1/metrics @markmc
/tests/v1/metrics @markmc
/vllm/tracing.py @markmc
/tests/v1/tracing/test_tracing.py @markmc
/vllm/config/kv_events.py @markmc
/vllm/distributed/kv_events.py @markmc
/tests/distributed/test_events.py @markmc
# Docs
/docs/mkdocs @hmellor
/docs/**/*.yml @hmellor
-5
View File
@@ -1,5 +0,0 @@
# Custom self-hosted runner labels (e.g. the autoscaling vllm-runners pool) so
# actionlint doesn't flag them as unknown in `runs-on`.
self-hosted-runner:
labels:
- vllm-runners
+1
View File
@@ -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 -11
View File
@@ -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:
@@ -144,12 +153,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:
@@ -388,13 +397,9 @@ pull_request_rules:
- or:
- files~=^tests/tool_use/
- files~=^tests/tool_parsers/
- files~=^tests/parser/
- files~=^tests/reasoning/
- files~=^tests/entrypoints/openai/.*tool.*
- files~=^tests/entrypoints/anthropic/.*tool.*
- files~=^vllm/tool_parsers/
- files~=^vllm/parser/
- files~=^vllm/reasoning/
- files=docs/features/tool_calling.md
- files~=^examples/tool_calling/
actions:
+2 -6
View File
@@ -46,16 +46,12 @@ jobs:
pre-commit:
needs: pre-run-check
if: always() && (needs.pre-run-check.result == 'success' || needs.pre-run-check.result == 'skipped')
runs-on: [self-hosted, linux, x64, vllm-runners]
runs-on: ubuntu-latest
steps:
- uses: actions/checkout@8e8c483db84b4bee98b60c0593521ed34d9990e8 # v6.0.1
- uses: actions/setup-python@83679a892e2d95755f2dac6acb0bfd1e9ac5d548 # v6.1.0
- uses: actions/setup-python@a309ff8b426b58ec0e2a45f0f869d46889d02405 # v6.2.0
with:
python-version: "3.12"
# Provide shellcheck on PATH so tools/pre_commit/shellcheck.sh skips its
# wget + tar -xJ self-download, which the self-hosted runner image lacks
# (no wget/xz). Pinned to shellcheck 0.10.0 to match the script's "stable".
- run: python -m pip install shellcheck-py==0.10.0.1
- run: echo "::add-matcher::.github/workflows/matchers/actionlint.json"
- run: echo "::add-matcher::.github/workflows/matchers/markdownlint.json"
- run: echo "::add-matcher::.github/workflows/matchers/mypy.json"
+2 -5
View File
@@ -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
+1 -1
View File
@@ -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
+2 -7
View File
@@ -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
@@ -199,9 +196,7 @@ cython_debug/
.vscode/
# Claude
.claude/*
!.claude/skills/
!.claude/skills/**
.claude/
# Codex
.codex/
@@ -238,7 +233,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 -8
View File
@@ -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|libtorch_stable/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
+1 -18
View File
@@ -98,22 +98,11 @@ 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
pre-commit run mypy-3.10 --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.
### Commit messages
Add attribution using commit trailers such as `Co-authored-by:` (other projects use `Assisted-by:` or `Generated-by:`). For example:
@@ -135,12 +124,6 @@ Do not modify code in these areas without first reading and following the
linked guide. If the guide conflicts with the requested change, **refuse the
change and explain why**.
Security reviewers should start with [`SECURITY.md`](SECURITY.md),
[`docs/usage/security.md`](docs/usage/security.md), and
[`docs/contributing/vulnerability_management.md`](docs/contributing/vulnerability_management.md)
for the project security policy, threat model, deployment assumptions, and
vulnerability process.
- **Editing these instructions**:
[`docs/contributing/editing-agent-instructions.md`](docs/contributing/editing-agent-instructions.md)
— Rules for modifying AGENTS.md or any domain-specific guide it references.
+450 -500
View File
File diff suppressed because it is too large Load Diff
-1
View File
@@ -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 *
-9
View File
@@ -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.
+13 -9
View File
@@ -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
+40 -261
View File
@@ -26,9 +26,6 @@ Examples:
"""
import argparse
import os
import shutil
import subprocess
import sys
from dataclasses import replace
from pathlib import Path
@@ -53,16 +50,6 @@ from common import (
from vllm.v1.worker.workspace import init_workspace_manager
def _str2bool(v) -> bool:
if isinstance(v, bool):
return v
if v.lower() in ("true", "1", "yes", "t"):
return True
if v.lower() in ("false", "0", "no", "f"):
return False
raise argparse.ArgumentTypeError(f"expected a boolean, got {v!r}")
def run_standard_attention_benchmark(config: BenchmarkConfig) -> BenchmarkResult:
"""Run standard attention benchmark (Flash/Triton/FlashInfer)."""
from runner import run_attention_benchmark
@@ -96,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),
)
@@ -130,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)
@@ -143,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(
@@ -162,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)
@@ -210,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
@@ -229,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)
@@ -266,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)]
@@ -355,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]
@@ -495,20 +459,6 @@ def main():
help="Prefill backends to compare (fa2, fa3, fa4). "
"Uses the first decode backend for impl construction.",
)
parser.add_argument(
"--fp8-output-scale",
type=float,
help="Static per-tensor scale enabling the MLA prefill FP8-output "
"comparison on FA4 (fused write vs standalone post-quant).",
)
parser.add_argument(
"--fuse-quant-op",
nargs="+",
type=_str2bool,
help="FP8-output write path(s) to run: false = bf16 attention + "
"standalone static-FP8 quant, true = FA4 writes FP8 directly. "
"Default: both.",
)
# Batch specifications
parser.add_argument(
@@ -524,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",
@@ -565,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(
@@ -642,12 +545,6 @@ def main():
# Prefill backends (e.g., ["fa3", "fa4"])
args.prefill_backends = yaml_config.get("prefill_backends", None)
# FP8 output benchmark knobs; CLI wins.
if args.fp8_output_scale is None:
args.fp8_output_scale = yaml_config.get("fp8_output_scale", None)
if args.fuse_quant_op is None:
args.fuse_quant_op = yaml_config.get("fuse_quant_op", None)
# Check for special modes
args.mode = yaml_config.get("mode", None)
@@ -679,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:
@@ -720,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}"
@@ -739,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)
@@ -789,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)
@@ -808,68 +662,8 @@ 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).[/]"
)
# FA4 fused FP8 output vs standalone post-quant, on the same fa4 kernel:
# the delta is the post-quant kernel the fused path removes.
fp8_output_scale = getattr(args, "fp8_output_scale", None)
if fp8_output_scale is not None:
decode_backend = backends[0]
fuse_variants = args.fuse_quant_op or [False, True]
label_of = {False: "post_quant", True: "fused"}
console.print(
f"[yellow]FP8 output comparison @ scale={fp8_output_scale} "
f"(prefill=fa4, decode impl={decode_backend})[/]"
)
fp8_results = []
total = len(fuse_variants) * len(args.batch_specs)
with tqdm(total=total, desc="FP8 output benchmarking") as pbar:
for spec in args.batch_specs:
for fuse in fuse_variants:
config = BenchmarkConfig(
backend=decode_backend,
batch_spec=spec,
num_layers=args.num_layers,
head_dim=args.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,
prefill_backend="fa4",
)
result = run_benchmark(
config, output_scale=fp8_output_scale, fuse_quant_op=fuse
)
label = label_of[fuse]
labeled_config = replace(result.config, backend=label)
result = replace(result, config=labeled_config)
fp8_results.append(result)
if not result.success:
console.print(f"[red]Error {label} {spec}: {result.error}[/]")
pbar.update(1)
console.print("\n[bold green]FP8 Output Results:[/]")
formatter = ResultsFormatter(console)
labels = [label_of[f] for f in fuse_variants]
formatter.print_table(fp8_results, labels, compare_to_fastest=True)
all_results = fp8_results
# Handle special mode: decode_vs_prefill comparison
elif hasattr(args, "mode") and args.mode == "decode_vs_prefill":
if hasattr(args, "mode") and args.mode == "decode_vs_prefill":
console.print("[yellow]Mode: Decode vs Prefill pipeline comparison[/]")
console.print(
"[dim]For each query length, testing both decode and prefill pipelines[/]"
@@ -914,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
@@ -955,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"],
@@ -977,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"),
@@ -987,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:[/]")
@@ -1069,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,
@@ -1098,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
@@ -1132,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)
@@ -1155,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:
@@ -1183,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,
)
@@ -1200,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)
+6 -54
View File
@@ -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
@@ -1,44 +0,0 @@
# MLA prefill FP8-output microbenchmark (FA4).
# Compares the fused FP8 write against bf16 attention + a standalone static-FP8
# quant; the delta is the post-quant kernel the fused path removes.
# DeepSeek-Coder-V2-Lite dims; FA4 needs SM100/110.
#
# Usage:
# python benchmark.py --config configs/mla_fa4_fp8_output.yaml
description: "MLA prefill FA4 fused-FP8 output vs post-quant"
model:
name: "deepseek-v2-lite"
num_layers: 27
num_q_heads: 16
num_kv_heads: 1
head_dim: 576
kv_lora_rank: 512
qk_nope_head_dim: 128
qk_rope_head_dim: 64
v_head_dim: 128
block_size: 128
# Pure prefill (q_len == kv_len) so every token goes through forward_mha.
batch_specs:
- "q512"
- "q1k"
- "q2k"
- "q4k"
- "q8k"
- "2q4k"
- "4q4k"
- "8q4k"
# Only used to construct the MLA impl; the pure-prefill specs skip decode.
decode_backends:
- CUTLASS_MLA
# Sweep the two FP8 write paths (prefill backend is fixed to fa4).
fp8_output_scale: 0.1
fuse_quant_op: [false, true]
device: "cuda:0"
repeats: 50
warmup_iters: 10
@@ -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
+42 -90
View File
@@ -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,
)
@@ -708,8 +704,6 @@ def _run_single_benchmark(
device: torch.device,
indexer=None,
kv_cache_dtype: str | None = None,
output_scale: float | None = None,
fuse_quant_op: bool = False,
) -> BenchmarkResult:
"""
Run a single benchmark iteration.
@@ -723,11 +717,6 @@ def _run_single_benchmark(
mla_dims: MLA dimension configuration
device: Target device
indexer: Optional MockIndexer for sparse backends
output_scale: Static per-tensor FP8 scale for prefill output. None
keeps the plain bf16 output (no quantization).
fuse_quant_op: With output_scale set, True lets the prefill kernel write
FP8 directly; False runs bf16 attention then a standalone static-FP8
quant. The delta isolates the saved post-quant kernel.
Returns:
BenchmarkResult with timing statistics
@@ -831,86 +820,63 @@ def _run_single_benchmark(
num_prefill, mla_dims, query_fmt, device, torch.bfloat16
)
# Prefill FP8 output: fused (kernel writes e4m3) vs separate post-quant.
prefill_fp8_output = None
prefill_output_scale = None
prefill_quant_op = None
if has_prefill and output_scale is not None:
from vllm.platforms import current_platform
prefill_output_scale = torch.tensor(
[output_scale], device=device, dtype=torch.float32
)
if fuse_quant_op:
prefill_fp8_output = torch.empty_like(
prefill_inputs["output"], dtype=current_platform.fp8_dtype()
)
else:
from vllm.model_executor.layers.quantization.input_quant_fp8 import (
QuantFP8,
)
from vllm.model_executor.layers.quantization.utils.quant_utils import (
GroupShape,
)
prefill_quant_op = QuantFP8(static=True, group_shape=GroupShape.PER_TENSOR)
fused_output = output_scale is not None and fuse_quant_op
# Build forward function (runs a single decode/prefill pass)
# Build forward function
def forward_fn():
results = []
if has_decode:
results.append(impl.forward_mqa(decode_inputs, kv_cache, metadata, layer))
if has_prefill:
out = impl.forward_mha(
prefill_inputs["q"],
prefill_inputs["k_c_normed"],
prefill_inputs["k_pe"],
kv_cache,
metadata,
prefill_inputs["k_scale"],
prefill_fp8_output if fused_output else prefill_inputs["output"],
prefill_output_scale if fused_output else None,
)
if fused_output:
out = prefill_fp8_output
elif prefill_quant_op is not None:
out, _ = prefill_quant_op(
prefill_inputs["output"], prefill_output_scale
results.append(
impl.forward_mha(
prefill_inputs["q"],
prefill_inputs["k_c_normed"],
prefill_inputs["k_pe"],
kv_cache,
metadata,
prefill_inputs["k_scale"],
prefill_inputs["output"],
)
results.append(out)
)
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,
)
@@ -920,8 +886,6 @@ def _run_mla_benchmark_batched(
configs_with_params: list[tuple], # [(config, threshold, num_splits), ...]
index_topk: int = 2048,
prefill_backend: str | None = None,
output_scale: float | None = None,
fuse_quant_op: bool = False,
) -> list[BenchmarkResult]:
"""
Unified batched MLA benchmark runner for all backends.
@@ -1061,8 +1025,6 @@ def _run_mla_benchmark_batched(
device,
indexer=indexer,
kv_cache_dtype=kv_cache_dtype,
output_scale=output_scale,
fuse_quant_op=fuse_quant_op,
)
results.append(result)
@@ -1090,8 +1052,6 @@ def run_mla_benchmark(
num_kv_splits: int | None = None,
index_topk: int = 2048,
prefill_backend: str | None = None,
output_scale: float | None = None,
fuse_quant_op: bool = False,
) -> BenchmarkResult | list[BenchmarkResult]:
"""
Unified MLA benchmark runner for all backends.
@@ -1111,9 +1071,6 @@ def run_mla_benchmark(
index_topk: Topk value for sparse MLA backends (default 2048)
prefill_backend: Prefill backend name (e.g., "fa3", "fa4").
When set, forces the specified FlashAttention version for prefill.
output_scale: Static per-tensor FP8 scale for prefill output (None = bf16).
fuse_quant_op: With output_scale set, fuse the FP8 write into the prefill
kernel vs a standalone post-quant kernel. See _run_single_benchmark.
Returns:
BenchmarkResult (single mode) or list of BenchmarkResult (batched mode)
@@ -1138,12 +1095,7 @@ def run_mla_benchmark(
# Use unified batched execution
results = _run_mla_benchmark_batched(
backend,
configs_with_params,
index_topk,
prefill_backend=prefill_backend,
output_scale=output_scale,
fuse_quant_op=fuse_quant_op,
backend, configs_with_params, index_topk, prefill_backend=prefill_backend
)
# Return single result or list based on input
+60 -53
View File
@@ -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"),
@@ -92,6 +92,7 @@ def run_baseline(
llm = LLM(
model=model,
enable_prefix_caching=False,
enable_chunked_prefill=False,
**extra_args,
)
sampling_params = SamplingParams(max_tokens=1)
@@ -193,6 +194,7 @@ async def _run_extraction_async(
engine_args = AsyncEngineArgs(
model=model,
enable_prefix_caching=False,
enable_chunked_prefill=False,
max_num_batched_tokens=40960,
max_model_len=40960,
speculative_config={
-358
View File
@@ -1,358 +0,0 @@
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
"""Benchmark and regression-test pinned (page-locked) CPU memory for vLLM.
Verifies that enabling pinned memory does not regress throughput or latency
compared to unpinned memory. Each condition runs in an isolated ``spawn``
subprocess so both start from a cold CUDA context, giving an unbiased
comparison.
Usage
-----
Run all tests with the default model::
python benchmarks/benchmark_pin_memory.py -v
Override the model and optional max-model-len::
python benchmarks/benchmark_pin_memory.py --model unsloth/Qwen3-1.7B -v
python benchmarks/benchmark_pin_memory.py --model unsloth/Qwen3-1.7B \
--max-model-len 8192 -v
Run only throughput or latency tests::
python benchmarks/benchmark_pin_memory.py -v -k test_throughput
python benchmarks/benchmark_pin_memory.py -v -k test_latency
Run only the v1 or v2 runner variant::
python benchmarks/benchmark_pin_memory.py -v -k v1
python benchmarks/benchmark_pin_memory.py -v -k v2
Note: on WSL2, v1 runner tests are skipped because pin memory is not available
for the v1 runner without cpu_offload_gb. Run on other platforms to exercise v1.
"""
import argparse
import json
import multiprocessing
import sys
import tempfile
import pytest
# Allow up to 2% degradation. Both benchmark runs start from an identical
# cold CUDA context (separate spawn subprocesses), so the measured difference
# reflects the genuine pin_memory overhead rather than cold/warm ordering bias.
_THROUGHPUT_TOLERANCE = 0.98
_THROUGHPUT_NUM_REQUESTS = 200
_THROUGHPUT_INPUT_LEN = 128
_THROUGHPUT_OUTPUT_LEN = 512
_THROUGHPUT_MAX_NUM_SEQS = 128
# Latency benchmark constants — match latency.py defaults.
_LATENCY_TOLERANCE = 1.02 # Allow up to 2% latency regression.
_LATENCY_BATCH_SIZE = 64
_LATENCY_INPUT_LEN = 32
_LATENCY_OUTPUT_LEN = 128
_LATENCY_WARMUP_ITERS = 5
_LATENCY_BENCH_ITERS = 15
_DEFAULT_MODEL = "unsloth/Qwen3-1.7B"
_DEFAULT_MAX_MODEL_LEN = 16384
def _benchmark_args() -> argparse.Namespace:
parser = argparse.ArgumentParser(add_help=False)
parser.add_argument("--model", default=_DEFAULT_MODEL)
parser.add_argument("--max-model-len", type=int, default=_DEFAULT_MAX_MODEL_LEN)
args, _ = parser.parse_known_args()
return args
@pytest.fixture
def model() -> str:
return _benchmark_args().model
@pytest.fixture
def max_model_len() -> int:
return _benchmark_args().max_model_len
def _skip_if_pin_memory_not_available(engine_args_kwargs: dict) -> None:
"""Skip the current pytest test if pin_memory is unavailable for this config."""
import vllm.utils.platform_utils as pu
from vllm.config import set_current_vllm_config
from vllm.engine.arg_utils import EngineArgs
vllm_config = EngineArgs(**engine_args_kwargs).create_engine_config()
with set_current_vllm_config(vllm_config):
pu.is_pin_memory_available.cache_clear()
if not pu.is_pin_memory_available():
import os
runner = "v2" if os.environ.get("VLLM_USE_V2_MODEL_RUNNER") == "1" else "v1"
model = engine_args_kwargs.get("model", "unknown")
print(
f"\033[33mSKIP: pin_memory not available for "
f"{runner} runner, model={model}\033[0m"
)
pytest.skip("pin_memory not available for this configuration")
def _throughput_worker(
pin: bool,
engine_args_kwargs: dict,
q: "multiprocessing.Queue[float]",
v2_mode: bool = False,
) -> None:
"""Run throughput benchmark in a fresh spawn subprocess.
Delegates to vllm/benchmarks/throughput.py main() using the random dataset,
so the methodology matches the official benchmark. Results are written to a
temp JSON file and forwarded through the queue as tokens/s.
v2_mode: when True, monkeypatches is_uva_available() to always return True
so the v2 model runner's UVA buffers remain functional even when pin=False.
This isolates the non-UVA pin_memory paths in v2.
"""
import vllm.utils.platform_utils as pu
from vllm.platforms import current_platform
pu.is_pin_memory_available.cache_clear()
pu.is_uva_available.cache_clear()
type(current_platform).is_pin_memory_available = classmethod(lambda cls: pin)
if v2_mode:
pu.is_uva_available = lambda: True
from vllm.benchmarks.throughput import add_cli_args
from vllm.benchmarks.throughput import main as throughput_main
parser = argparse.ArgumentParser()
add_cli_args(parser)
args = parser.parse_args([])
for key, val in engine_args_kwargs.items():
setattr(args, key, val)
args.max_num_seqs = _THROUGHPUT_MAX_NUM_SEQS
args.dataset_name = "random"
args.input_len = _THROUGHPUT_INPUT_LEN
args.output_len = _THROUGHPUT_OUTPUT_LEN
# Nullify defaults that conflict with explicit input/output_len.
args.random_input_len = None
args.random_output_len = None
args.random_prefix_len = None
args.num_prompts = _THROUGHPUT_NUM_REQUESTS
args.seed = 0
args.disable_detokenize = True
with tempfile.NamedTemporaryFile(mode="w", suffix=".json", delete=False) as f:
tmp_path = f.name
args.output_json = tmp_path
throughput_main(args)
with open(tmp_path) as f:
results = json.load(f)
q.put(results["tokens_per_second"])
def _run_throughput_benchmark(
pin: bool,
engine_args_kwargs: dict,
v2_mode: bool = False,
) -> float:
ctx = multiprocessing.get_context("spawn")
q = ctx.Queue()
p = ctx.Process(
target=_throughput_worker,
args=(pin, engine_args_kwargs, q, v2_mode),
)
p.start()
p.join()
if p.exitcode != 0:
raise RuntimeError(
f"Throughput benchmark subprocess (pin={pin}) exited with code {p.exitcode}"
)
return q.get()
def _latency_worker(
pin: bool,
engine_args_kwargs: dict,
q: "multiprocessing.Queue[dict]",
v2_mode: bool = False,
) -> None:
"""Run latency benchmark in a fresh spawn subprocess.
Follows latency.py methodology: fixed batch of dummy token IDs, warmup
iterations to reach steady state, then timed iterations reduced to avg
and percentiles. Results are written to a temp JSON file by latency_main
and forwarded through the queue.
"""
import vllm.utils.platform_utils as pu
from vllm.platforms import current_platform
pu.is_pin_memory_available.cache_clear()
pu.is_uva_available.cache_clear()
type(current_platform).is_pin_memory_available = classmethod(lambda cls: pin)
if v2_mode:
pu.is_uva_available = lambda: True
from vllm.benchmarks.latency import add_cli_args
from vllm.benchmarks.latency import main as latency_main
parser = argparse.ArgumentParser()
add_cli_args(parser)
args = parser.parse_args([])
for key, val in engine_args_kwargs.items():
setattr(args, key, val)
args.input_len = _LATENCY_INPUT_LEN
args.output_len = _LATENCY_OUTPUT_LEN
args.batch_size = _LATENCY_BATCH_SIZE
args.num_iters_warmup = _LATENCY_WARMUP_ITERS
args.num_iters = _LATENCY_BENCH_ITERS
args.profile = False
args.disable_detokenize = True
with tempfile.NamedTemporaryFile(mode="w", suffix=".json", delete=False) as f:
tmp_path = f.name
args.output_json = tmp_path
latency_main(args)
with open(tmp_path) as f:
results = json.load(f)
q.put(results)
def _run_latency_benchmark(
pin: bool,
engine_args_kwargs: dict,
v2_mode: bool = False,
) -> dict:
ctx = multiprocessing.get_context("spawn")
q = ctx.Queue()
p = ctx.Process(
target=_latency_worker,
args=(pin, engine_args_kwargs, q, v2_mode),
)
p.start()
p.join()
if p.exitcode != 0:
raise RuntimeError(
f"Latency benchmark subprocess (pin={pin}) exited with code {p.exitcode}"
)
return q.get()
@pytest.mark.parametrize(
"test_v2_runner",
[
pytest.param(False, id="v1"),
pytest.param(True, id="v2"),
],
)
class TestPinnedMemory:
"""Verify pinned memory yields >= throughput vs unpinned via real vLLM inference."""
def test_throughput(self, monkeypatch, test_v2_runner, model, max_model_len):
"""Benchmark throughput with pin_memory forced on then off.
Delegates to vllm/benchmarks/throughput.py main() with the random
dataset. Each condition runs in an isolated spawn subprocess so both
start from a cold CUDA context, giving an unbiased comparison.
"""
monkeypatch.setenv("VLLM_ENABLE_V1_MULTIPROCESSING", "0")
monkeypatch.setenv("VLLM_USE_V2_MODEL_RUNNER", "1" if test_v2_runner else "0")
engine_args_kwargs = dict(
model=model,
gpu_memory_utilization=0.88,
max_model_len=max_model_len,
enable_prefix_caching=False,
)
_skip_if_pin_memory_not_available(engine_args_kwargs)
unpinned_tps = _run_throughput_benchmark(
False, engine_args_kwargs, v2_mode=test_v2_runner
)
pinned_tps = _run_throughput_benchmark(
True, engine_args_kwargs, v2_mode=test_v2_runner
)
pct_diff = (pinned_tps - unpinned_tps) / unpinned_tps * 100
runner = "v2" if test_v2_runner else "v1"
print(
f"\n=== Throughput results ({runner} runner, {model}) ==="
f"\npin_memory=True: {pinned_tps:.1f} tok/s"
f"\npin_memory=False: {unpinned_tps:.1f} tok/s"
f"\nDifference: {pct_diff:+.1f}% (pinned vs unpinned)"
)
assert pinned_tps >= unpinned_tps * _THROUGHPUT_TOLERANCE, (
f"Pinned throughput ({pinned_tps:.1f} tok/s) fell more than "
f"{(1.0 - _THROUGHPUT_TOLERANCE) * 100:.1f}% below "
f"unpinned ({unpinned_tps:.1f} tok/s)."
)
def test_latency(self, monkeypatch, test_v2_runner, model, max_model_len):
"""Benchmark per-batch latency with pin_memory forced on then off.
Follows vllm/benchmarks/latency.py: fixed dummy-token batch, warmup
iterations to reach steady state, then timed iterations reduced to avg
and percentiles. Subprocesses run serially so each gets a cold CUDA
context without GPU memory pressure from the other run.
"""
monkeypatch.setenv("VLLM_ENABLE_V1_MULTIPROCESSING", "0")
monkeypatch.setenv("VLLM_USE_V2_MODEL_RUNNER", "1" if test_v2_runner else "0")
engine_args_kwargs = dict(
model=model,
gpu_memory_utilization=0.88,
max_model_len=max_model_len,
enable_prefix_caching=False,
)
_skip_if_pin_memory_not_available(engine_args_kwargs)
unpinned = _run_latency_benchmark(
False, engine_args_kwargs, v2_mode=test_v2_runner
)
pinned = _run_latency_benchmark(
True, engine_args_kwargs, v2_mode=test_v2_runner
)
pct_diff = (
(pinned["avg_latency"] - unpinned["avg_latency"])
/ unpinned["avg_latency"]
* 100
)
runner = "v2" if test_v2_runner else "v1"
print(
f"\n=== Latency results ({runner} runner, {model}) ==="
f"\npin_memory=True: avg={pinned['avg_latency']:.3f}s"
f" p50={pinned['percentiles']['50']:.3f}s"
f" p99={pinned['percentiles']['99']:.3f}s"
f"\npin_memory=False: avg={unpinned['avg_latency']:.3f}s"
f" p50={unpinned['percentiles']['50']:.3f}s"
f" p99={unpinned['percentiles']['99']:.3f}s"
f"\nDifference: {pct_diff:+.1f}% (pinned vs unpinned)"
)
assert pinned["avg_latency"] <= unpinned["avg_latency"] * _LATENCY_TOLERANCE, (
f"Pinned avg latency ({pinned['avg_latency']:.3f}s) exceeded "
f"unpinned ({unpinned['avg_latency']:.3f}s) by more than "
f"{(_LATENCY_TOLERANCE - 1.0) * 100:.1f}%."
)
if __name__ == "__main__":
_parser = argparse.ArgumentParser(add_help=False)
_parser.add_argument("--model", default=_DEFAULT_MODEL)
_parser.add_argument("--max-model-len", type=int, default=_DEFAULT_MAX_MODEL_LEN)
_, _remaining = _parser.parse_known_args()
sys.exit(pytest.main([__file__] + _remaining))
@@ -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)
@@ -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")
+4 -7
View File
@@ -250,7 +250,7 @@ def benchmark_config(
num_experts=num_experts,
experts_per_token=topk,
hidden_dim=hidden_size,
intermediate_size=shard_intermediate_size,
intermediate_size_per_partition=shard_intermediate_size,
num_local_experts=num_experts,
num_logical_experts=num_experts,
activation=MoEActivation.SILU,
@@ -271,6 +271,7 @@ def benchmark_config(
moe_config=moe_config,
quant_config=quant_config,
),
inplace=not disable_inplace(),
)
with override_config(config):
@@ -278,6 +279,7 @@ def benchmark_config(
x, input_gating, topk, renormalize=not use_deep_gemm
)
inplace = not disable_inplace()
if use_deep_gemm:
return deep_gemm_experts.apply(
x,
@@ -296,6 +298,7 @@ def benchmark_config(
w2,
topk_weights,
topk_ids,
inplace=inplace,
quant_config=quant_config,
)
@@ -792,12 +795,6 @@ def get_model_params(config):
topk = text_config.num_experts_per_tok
intermediate_size = text_config.moe_intermediate_size
hidden_size = text_config.hidden_size
elif architecture == "DiffusionGemmaForBlockDiffusion":
text_config = config.get_text_config()
E = text_config.num_experts
topk = text_config.top_k_experts
intermediate_size = text_config.moe_intermediate_size
hidden_size = text_config.hidden_size
elif architecture == "HunYuanMoEV1ForCausalLM":
E = config.num_experts
topk = config.moe_topk[0]
-248
View File
@@ -1,248 +0,0 @@
#!/bin/bash
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
#
# Reproducible demonstration of the KV cache watermark (`--watermark`) for
# reducing preemption thrashing.
#
# The watermark is the fraction of total KV cache blocks the scheduler keeps
# free when admitting a waiting/preempted request into the running queue.
#
# Why this workload triggers thrashing:
# Requests are admitted based on the KV cache they need *at admission time*.
# With `--scheduler-reserve-full-isl` (default) the input length is reserved up
# front, but the *output* length is unknown and unreserved. A decode-heavy
# workload (output >> input) at high concurrency therefore over-admits while
# requests are short, then runs out of KV cache as they all grow during decode
# -> the scheduler preempts (recompute) recently-admitted requests, re-prefills
# them later, and repeats. The watermark keeps a block of KV cache free so
# running requests can grow into it instead of triggering this churn.
#
# This script launches `vllm serve` under a deliberately KV-constrained config
# and a decode-heavy workload, sweeping the watermark across several values, and
# reports the preemption count (scraped from /metrics), throughput, and latency
# percentiles for each. It then plots the results.
#
# Default workload: concurrency 200, input ~300 tokens, output ~4000 tokens
# (+/- 20% variance), sized to run each config for ~5 minutes.
#
# Usage:
# benchmarks/kv_cache_watermark.sh
# MODEL=Qwen/Qwen2.5-14B-Instruct TP=2 benchmarks/kv_cache_watermark.sh
#
# Run inside the vLLM virtualenv (so `vllm` and `python` resolve to it).
set -euo pipefail
# ---- Config (override via environment) -------------------------------------
MODEL=${MODEL:-Qwen/Qwen2.5-7B-Instruct}
TP=${TP:-1}
PORT=${PORT:-8000}
URL="http://127.0.0.1:${PORT}"
# Constrain the KV cache to a *near-critical* size: large enough that the engine
# can run stably, but small enough that greedy over-admission tips it into
# preemption thrashing. (Independent of GPU size, so the demo is reproducible.)
# At the default workload this fits ~1.5x the mean concurrent KV demand.
KV_CACHE_MEMORY_GB=${KV_CACHE_MEMORY_GB:-16}
MAX_MODEL_LEN=${MAX_MODEL_LEN:-8192}
MAX_NUM_SEQS=${MAX_NUM_SEQS:-256}
# Optional weight loader (e.g. fastsafetensors on the GCP cluster).
LOAD_FORMAT=${LOAD_FORMAT:-auto}
# Decode-heavy workload: moderate input, long output, with length variance. The
# long output means preempted requests have generated a lot before eviction, so
# resuming them re-prefills a long sequence (high recomputation cost).
INPUT_LEN=${INPUT_LEN:-1000}
OUTPUT_LEN=${OUTPUT_LEN:-5000}
RANGE_RATIO=${RANGE_RATIO:-0.2}
CONCURRENCY=${CONCURRENCY:-128}
# Enough prompts to keep each config saturated for ~5+ minutes.
NUM_PROMPTS=${NUM_PROMPTS:-450}
OUTDIR=${OUTDIR:-./watermark_bench_results}
# Watermark fractions compared. "label value" per line; value=0 disables it.
CONFIGS=${CONFIGS:-"off 0
w0.02 0.02
w0.05 0.05
w0.10 0.10
w0.15 0.15"}
KV_CACHE_MEMORY_BYTES=$((KV_CACHE_MEMORY_GB * 1024 * 1024 * 1024))
mkdir -p "$OUTDIR"
SERVER_PID=""
cleanup() { [[ -n "$SERVER_PID" ]] && kill "$SERVER_PID" 2>/dev/null || true; }
trap cleanup EXIT
scrape_preemptions() {
# Sum the vllm:num_preemptions_total counter across engines.
python - "${URL}/metrics" <<'PY'
import sys, urllib.request
total = 0.0
try:
body = urllib.request.urlopen(sys.argv[1], timeout=10).read().decode("utf-8", "replace")
for line in body.splitlines():
if line.startswith("vllm:num_preemptions_total"):
total += float(line.rsplit(" ", 1)[-1])
except Exception as e: # noqa: BLE001
print(f"scrape error: {e}", file=sys.stderr)
print(int(total))
PY
}
wait_for_server() {
for _ in $(seq 1 300); do
if curl -s "${URL}/health" >/dev/null 2>&1; then return 0; fi
if ! kill -0 "$SERVER_PID" 2>/dev/null; then
echo "ERROR: server process exited during startup" >&2; return 1
fi
sleep 5
done
echo "ERROR: server did not become ready" >&2; return 1
}
run_one() {
local label=$1 watermark=$2
echo
echo "==================== watermark: ${label} (${watermark}) ===================="
vllm serve "$MODEL" \
--tensor-parallel-size "$TP" \
--load-format "$LOAD_FORMAT" \
--kv-cache-memory-bytes "$KV_CACHE_MEMORY_BYTES" \
--max-model-len "$MAX_MODEL_LEN" \
--max-num-seqs "$MAX_NUM_SEQS" \
--no-enable-prefix-caching \
--watermark "$watermark" \
--port "$PORT" >"${OUTDIR}/serve_${label}.log" 2>&1 &
SERVER_PID=$!
wait_for_server
sleep 5
local pre post
pre=$(scrape_preemptions)
vllm bench serve \
--backend vllm \
--base-url "$URL" \
--model "$MODEL" \
--dataset-name random \
--random-input-len "$INPUT_LEN" \
--random-output-len "$OUTPUT_LEN" \
--random-range-ratio "$RANGE_RATIO" \
--ignore-eos \
--num-prompts "$NUM_PROMPTS" \
--max-concurrency "$CONCURRENCY" \
--percentile-metrics "ttft,tpot,itl,e2el" \
--metric-percentiles "50,90,99" \
--save-result \
--result-dir "$OUTDIR" \
--result-filename "bench_${label}.json"
post=$(scrape_preemptions)
echo "${label} ${watermark} $((post - pre))" >>"${OUTDIR}/preemptions.txt"
kill "$SERVER_PID" 2>/dev/null || true
for _ in $(seq 1 60); do curl -s "${URL}/health" >/dev/null 2>&1 || break; sleep 2; done
SERVER_PID=""
sleep 10
}
: >"${OUTDIR}/preemptions.txt"
while read -r label watermark; do
[[ -z "${label:-}" ]] && continue
run_one "$label" "$watermark"
done <<<"$CONFIGS"
echo
echo "==================== summary ===================="
python - "$OUTDIR" <<'PY'
import json, os, sys
outdir = sys.argv[1]
pre = {}
order = []
for line in open(os.path.join(outdir, "preemptions.txt")):
label, watermark, n = line.split()
pre[label] = (float(watermark), int(n))
order.append(label)
def g(d, *names):
for n in names:
if d.get(n) is not None:
return d[n]
return float("nan")
cols = ["watermark", "frac", "preempt", "out_tok/s", "req/s",
"TTFT_p50", "TTFT_p99", "ITL_p99", "E2EL_p50"]
print(" ".join(f"{c:>10}" for c in cols))
rows = []
for label in order:
watermark, n = pre[label]
d = json.load(open(os.path.join(outdir, f"bench_{label}.json")))
rows.append(dict(
label=label, watermark=watermark, preempt=n,
out_tok_s=g(d, "output_throughput"),
req_s=g(d, "request_throughput"),
ttft_p50=g(d, "p50_ttft_ms", "median_ttft_ms"),
ttft_p99=g(d, "p99_ttft_ms"),
itl_p99=g(d, "p99_itl_ms"),
e2el_p50=g(d, "p50_e2el_ms", "median_e2el_ms"),
))
print(" ".join(f"{str(v):>10}" for v in [
label, watermark, n,
f"{rows[-1]['out_tok_s']:.0f}",
f"{rows[-1]['req_s']:.3f}",
f"{rows[-1]['ttft_p50']/1000:.2f}",
f"{rows[-1]['ttft_p99']/1000:.2f}",
f"{rows[-1]['itl_p99']:.2f}",
f"{rows[-1]['e2el_p50']/1000:.1f}",
]))
print("\n(TTFT/E2EL in seconds; ITL in ms. Lower preempt is better.)")
# ---- Plot -------------------------------------------------------------------
try:
import matplotlib
matplotlib.use("Agg")
import matplotlib.pyplot as plt
except Exception as e: # noqa: BLE001
print(f"\n(skip plot: matplotlib unavailable: {e})")
sys.exit(0)
x = [r["watermark"] for r in rows]
xt = [f"{r['watermark']:g}\n({r['label']})" for r in rows]
idx = list(range(len(rows)))
fig, axes = plt.subplots(2, 2, figsize=(12, 8))
fig.suptitle(
f"KV cache watermark sweep — {os.path.basename(os.path.abspath(outdir))}",
fontsize=12,
)
ax = axes[0][0]
ax.bar(idx, [r["preempt"] for r in rows], color="tab:red")
ax.set_title("Preemptions (lower is better)")
ax.set_ylabel("preemptions")
ax.set_xticks(idx); ax.set_xticklabels(xt)
ax = axes[0][1]
ax.plot(idx, [r["out_tok_s"] for r in rows], "o-", color="tab:green")
ax.set_title("Output throughput (higher is better)")
ax.set_ylabel("tokens/s")
ax.set_xticks(idx); ax.set_xticklabels(xt)
ax = axes[1][0]
ax.plot(idx, [r["itl_p99"] for r in rows], "o-", color="tab:blue")
ax.set_title("Inter-token latency p99 (lower is better)")
ax.set_ylabel("ITL p99 (ms)")
ax.set_xlabel("watermark fraction")
ax.set_xticks(idx); ax.set_xticklabels(xt)
ax = axes[1][1]
ax.plot(idx, [r["ttft_p50"] / 1000 for r in rows], "o-", label="TTFT p50")
ax.plot(idx, [r["ttft_p99"] / 1000 for r in rows], "o-", label="TTFT p99")
ax.plot(idx, [r["e2el_p50"] / 1000 for r in rows], "o-", label="E2EL p50")
ax.set_title("Latency (lower is better)")
ax.set_ylabel("seconds")
ax.set_xlabel("watermark fraction")
ax.set_xticks(idx); ax.set_xticklabels(xt)
ax.legend()
fig.tight_layout(rect=(0, 0, 1, 0.95))
out_png = os.path.join(outdir, "watermark_results.png")
fig.savefig(out_png, dpi=120)
print(f"\nWrote plot: {out_png}")
PY
@@ -65,32 +65,6 @@ class RequestArgs(NamedTuple):
limit_min_tokens: int # Use negative value for no limit
limit_max_tokens: int # Use negative value for no limit
timeout_sec: int
send_conversation_id: bool
headers: dict[str, str]
def parse_custom_header(header: str) -> tuple[str, str]:
separators = (":", "=")
for separator in separators:
if separator in header:
key, value = header.split(separator, 1)
key = key.strip()
value = value.strip()
if key:
return key, value
break
raise argparse.ArgumentTypeError(
"Headers must be provided as 'Header-Name: value' or 'Header-Name=value'"
)
def build_request_headers(
api_key: str | None, custom_headers: list[tuple[str, str]] | None
) -> dict[str, str]:
headers = dict(custom_headers or [])
if api_key:
headers["Authorization"] = f"Bearer {api_key}"
return headers
class BenchmarkArgs(NamedTuple):
@@ -244,11 +218,12 @@ async def send_request(
max_tokens: int | None = None,
timeout_sec: int = 120,
conversation_id: str | None = None,
headers: dict[str, str] | None = None,
) -> ServerResponse:
payload = {
"model": model,
"messages": messages,
"seed": 0,
"temperature": 0.0,
}
if conversation_id is not None:
@@ -258,17 +233,13 @@ async def send_request(
payload["stream"] = True
payload["stream_options"] = {"include_usage": False}
# if min_tokens is not None:
# payload["min_tokens"] = min_tokens
if min_tokens is not None:
payload["min_tokens"] = min_tokens
if max_tokens is not None:
payload["max_tokens"] = max_tokens
request_headers = {"Content-Type": "application/json"}
if conversation_id is not None:
request_headers["X-Session-ID"] = str(conversation_id)
if headers is not None:
request_headers.update(headers)
headers = {"Content-Type": "application/json"}
# Calculate the timeout for the request
if max_tokens is not None:
@@ -294,7 +265,7 @@ async def send_request(
most_recent_timestamp: int = start_time
async with session.post(
url=chat_url, json=payload, headers=request_headers, timeout=timeout
url=chat_url, json=payload, headers=headers, timeout=timeout
) as response:
http_status = HTTPStatus(response.status)
if http_status == HTTPStatus.OK:
@@ -346,8 +317,6 @@ async def send_request(
latency = time.perf_counter_ns() - start_time
if ttft is None:
if stream:
valid_response = False
# The response was a single chunk
ttft = latency
@@ -454,8 +423,7 @@ async def send_turn(
min_tokens,
max_tokens,
req_args.timeout_sec,
conversation_id=conv_id if req_args.send_conversation_id else None,
headers=req_args.headers,
conversation_id=conv_id,
)
if response.valid is False:
@@ -904,7 +872,6 @@ def get_client_config(
# Arguments for API requests
chat_url = f"{args.url}/v1/chat/completions"
model_name = args.served_model_name if args.served_model_name else args.model
headers = build_request_headers(args.api_key, args.header)
req_args = RequestArgs(
chat_url=chat_url,
@@ -913,8 +880,6 @@ def get_client_config(
limit_min_tokens=args.limit_min_tokens,
limit_max_tokens=args.limit_max_tokens,
timeout_sec=args.request_timeout_sec,
send_conversation_id=args.send_conversation_id,
headers=headers,
)
return client_args, req_args
@@ -1280,19 +1245,19 @@ def process_statistics(
)
async def get_server_info(url: str, headers: dict[str, str] | None = None) -> None:
async def get_server_info(url: str) -> None:
logger.info(f"{Color.BLUE}Collecting information from server: {url}{Color.RESET}")
async with aiohttp.ClientSession() as session:
# Get server version (not mandatory, "version" endpoint may not exist)
url_version = f"{url}/version"
async with session.get(url_version, headers=headers) as response:
async with session.get(url_version) as response:
if HTTPStatus(response.status) == HTTPStatus.OK:
text = await response.text()
logger.info(f"{Color.BLUE}Server version: {text}{Color.RESET}")
# Get available models
url_models = f"{url}/v1/models"
async with session.get(url_models, headers=headers) as response:
async with session.get(url_models) as response:
if HTTPStatus(response.status) == HTTPStatus.OK:
text = await response.text()
logger.info(f"{Color.BLUE}Models:{Color.RESET}")
@@ -1358,22 +1323,6 @@ async def main() -> None:
help="Base URL for the LLM API server",
)
parser.add_argument(
"--api-key",
type=str,
default=None,
help="API key to send as an Authorization bearer token",
)
parser.add_argument(
"--header",
action="append",
type=parse_custom_header,
default=None,
metavar="KEY=VALUE",
help="Custom request header. Can be specified multiple times. "
"Accepts 'Header-Name: value' or 'Header-Name=value'.",
)
parser.add_argument(
"-p",
"--num-clients",
@@ -1488,22 +1437,6 @@ async def main() -> None:
help="Disable stream/streaming mode (set 'stream' to False in the API request)",
)
parser.add_argument(
"--send-conversation-id",
default=False,
action="store_true",
help=(
"Inject a `conversation_id` field into each Chat Completions "
"payload. This is a non-standard OpenAI extension consumed by "
"vLLM's disaggregated multi-turn proxy "
"(examples/disaggregated/disaggregated_serving/"
"disagg_proxy_multiturn.py) to key cross-turn KV cache reuse. "
"Leave disabled (default) when targeting strict "
"OpenAI-compatible endpoints; enable when benchmarking the "
"disaggregated proxy."
),
)
parser.add_argument(
"-e",
"--excel-output",
@@ -1592,8 +1525,7 @@ async def main() -> None:
args.model, trust_remote_code=args.trust_remote_code
)
headers = build_request_headers(args.api_key, args.header)
await get_server_info(args.url, headers=headers)
await get_server_info(args.url)
# Load the input file (either conversations of configuration file)
logger.info(f"Reading input file: {args.input_file}")
+15 -4
View File
@@ -1,5 +1,5 @@
#!/bin/bash
# Build vLLM Rust artifacts and install them into the vllm package.
# Build the vllm-rs Rust frontend binary and install it into the vllm package.
# Usage: ./build_rust.sh [--debug]
#
# By default builds in release mode. Pass --debug for faster compile times
@@ -8,6 +8,8 @@
set -euo pipefail
REPO_ROOT="$(cd "$(dirname "$0")" && pwd)"
RUST_DIR="$REPO_ROOT/rust"
TARGET_PATH="${VLLM_RS_TARGET_PATH:-$REPO_ROOT/vllm/vllm-rs}"
# Read the required toolchain from rust-toolchain.toml.
TOOLCHAIN=$(grep '^channel' "$REPO_ROOT/rust-toolchain.toml" | sed 's/.*= *"\(.*\)"/\1/')
@@ -25,9 +27,18 @@ if ! rustup run "$TOOLCHAIN" rustc --version &>/dev/null; then
fi
if [[ "${1:-}" == "--debug" ]]; then
PROFILE_ARG="--debug"
PROFILE_ARGS=()
PROFILE_DIR="debug"
else
PROFILE_ARG="--release"
PROFILE_ARGS=(--release)
PROFILE_DIR="release"
fi
python3 "$REPO_ROOT/tools/build_rust.py" "$PROFILE_ARG"
cargo +"$TOOLCHAIN" build "${PROFILE_ARGS[@]}" \
--manifest-path "$RUST_DIR/Cargo.toml" \
--bin vllm-rs \
--features native-tls-vendored
mkdir -p "$(dirname "$TARGET_PATH")"
cp "$RUST_DIR/target/$PROFILE_DIR/vllm-rs" "$TARGET_PATH"
echo "Installed vllm-rs to $TARGET_PATH"
+3 -12
View File
@@ -166,10 +166,6 @@ elseif (S390_FOUND)
"-mtune=native")
elseif (CMAKE_SYSTEM_PROCESSOR MATCHES "riscv64")
message(STATUS "RISC-V detected")
if(DEFINED VLLM_RVV_VLEN AND NOT VLLM_RVV_VLEN GREATER 0)
message(FATAL_ERROR
"VLLM_RVV_VLEN must be a positive integer; got '${VLLM_RVV_VLEN}'")
endif()
# VLLM_RVV_VLEN selects the target VLEN. Auto-detected from /proc/cpuinfo
# by default; override with -DVLLM_RVV_VLEN=128 or -DVLLM_RVV_VLEN=256.
if(NOT DEFINED VLLM_RVV_VLEN)
@@ -193,7 +189,8 @@ elseif (CMAKE_SYSTEM_PROCESSOR MATCHES "riscv64")
"RISC-V RVV is available but VLEN could not be auto-detected. "
"Please specify VLEN explicitly:\n"
" -DVLLM_RVV_VLEN=128 (for VLEN=128 hardware)\n"
" -DVLLM_RVV_VLEN=256 (for VLEN=256 hardware, e.g. Spacemit X100)")
" -DVLLM_RVV_VLEN=256 (for VLEN=256 hardware, e.g. Spacemit X100)\n"
" -DVLLM_RVV_VLEN=0 (force scalar, no RVV)")
endif()
endif()
if(VLLM_RVV_VLEN AND VLLM_RVV_VLEN GREATER 0)
@@ -222,7 +219,7 @@ endif()
# Build oneDNN for GEMM kernels
if (ENABLE_X86_ISA OR (ASIMD_FOUND AND NOT APPLE_SILICON_FOUND) OR POWER9_FOUND OR POWER10_FOUND OR POWER11_FOUND OR RVV_FP16_FOUND OR RVV_BF16_FOUND)
if (ENABLE_X86_ISA OR (ASIMD_FOUND AND NOT APPLE_SILICON_FOUND) OR POWER9_FOUND OR POWER10_FOUND OR POWER11_FOUND)
# Fetch and build Arm Compute Library (ACL) as oneDNN's backend for AArch64
# TODO [fadara01]: remove this once ACL can be fetched and built automatically as a dependency of oneDNN
set(ONEDNN_AARCH64_USE_ACL OFF CACHE BOOL "")
@@ -438,12 +435,6 @@ if(USE_ONEDNN)
${VLLM_EXT_SRC})
endif()
if (CMAKE_SYSTEM_PROCESSOR MATCHES "riscv64")
set(VLLM_EXT_SRC
"csrc/cpu/sgl-kernels/gemm_int4.cpp"
${VLLM_EXT_SRC})
endif()
if (ENABLE_X86_ISA)
set(VLLM_EXT_SRC_SGL
"csrc/cpu/sgl-kernels/conv.cpp"
+29 -59
View File
@@ -8,73 +8,43 @@ if (DEFINED ENV{DEEPGEMM_SRC_DIR})
set(DEEPGEMM_SRC_DIR $ENV{DEEPGEMM_SRC_DIR})
endif()
# Local tree: set deepgemm_SOURCE_DIR directly (no FetchContent download).
# Upstream git: use FetchContent_Populate with explicit options (CMP0169 NEW
# disallows one-argument Populate(dep) after Declare; MakeAvailable would run
# DeepGEMM's top-level CMakeLists.txt, which vLLM must not load).
if(DEEPGEMM_SRC_DIR)
# cmake_path(ABSOLUTE_PATH <var> ...) reads the path from <var>; NORMALIZE is a
# flag (no trailing path argument). Resolve relative paths against vLLM root.
set(_deepgemm_user_src "${DEEPGEMM_SRC_DIR}")
cmake_path(ABSOLUTE_PATH _deepgemm_user_src
BASE_DIRECTORY "${CMAKE_SOURCE_DIR}"
NORMALIZE)
set(DEEPGEMM_SRC_DIR "${_deepgemm_user_src}")
if(NOT IS_DIRECTORY "${DEEPGEMM_SRC_DIR}")
message(FATAL_ERROR
"DEEPGEMM_SRC_DIR is not an existing directory: '${DEEPGEMM_SRC_DIR}'")
endif()
set(deepgemm_SOURCE_DIR "${DEEPGEMM_SRC_DIR}")
message(STATUS "DeepGEMM using local DEEPGEMM_SRC_DIR: ${deepgemm_SOURCE_DIR}")
FetchContent_Declare(
deepgemm
SOURCE_DIR ${DEEPGEMM_SRC_DIR}
CONFIGURE_COMMAND ""
BUILD_COMMAND ""
)
else()
# Keep in sync with tools/install_deepgemm.sh
set(_DEEPGEMM_UPSTREAM_REPO "https://github.com/deepseek-ai/DeepGEMM.git")
set(_DEEPGEMM_UPSTREAM_TAG "891d57b4db1071624b5c8fa0d1e51cb317fa709f")
set(_deepgemm_fc_root "${FETCHCONTENT_BASE_DIR}")
if(NOT _deepgemm_fc_root)
set(_deepgemm_fc_root "${CMAKE_BINARY_DIR}/_deps")
endif()
set(_deepgemm_src "${_deepgemm_fc_root}/deepgemm-src")
set(_deepgemm_bin "${_deepgemm_fc_root}/deepgemm-build")
set(_deepgemm_sub "${_deepgemm_fc_root}/deepgemm-subbuild")
if(EXISTS "${_deepgemm_src}/csrc/python_api.cpp")
set(deepgemm_SOURCE_DIR "${_deepgemm_src}")
set(deepgemm_BINARY_DIR "${_deepgemm_bin}")
else()
FetchContent_Populate(
deepgemm
SUBBUILD_DIR "${_deepgemm_sub}"
SOURCE_DIR "${_deepgemm_src}"
BINARY_DIR "${_deepgemm_bin}"
GIT_REPOSITORY "${_DEEPGEMM_UPSTREAM_REPO}"
GIT_TAG "${_DEEPGEMM_UPSTREAM_TAG}"
GIT_SUBMODULES "third-party/cutlass" "third-party/fmt"
GIT_PROGRESS TRUE
)
endif()
message(STATUS "DeepGEMM is available at ${deepgemm_SOURCE_DIR}")
# This ref should be kept in sync with tools/install_deepgemm.sh
FetchContent_Declare(
deepgemm
GIT_REPOSITORY https://github.com/deepseek-ai/DeepGEMM.git
GIT_TAG 891d57b4db1071624b5c8fa0d1e51cb317fa709f
GIT_SUBMODULES "third-party/cutlass" "third-party/fmt"
GIT_PROGRESS TRUE
CONFIGURE_COMMAND ""
BUILD_COMMAND ""
)
endif()
# DeepGEMM requires CUDA 12.3+ for SM90, 12.9+ for SM100 (official upstream),
# and 12.8+ for SM120 / SM12x. CUDA 13+ can use the family-specific SM12x
# arch; CUDA 12.x builds the arch-specific SM120/SM121 variants.
# Use FetchContent_Populate (not MakeAvailable) to avoid processing
# DeepGEMM's own CMakeLists.txt which has incompatible find_package calls.
FetchContent_GetProperties(deepgemm)
if(NOT deepgemm_POPULATED)
FetchContent_Populate(deepgemm)
endif()
message(STATUS "DeepGEMM is available at ${deepgemm_SOURCE_DIR}")
# DeepGEMM requires CUDA 12.3+ for SM90, 12.9+ for SM100
set(DEEPGEMM_SUPPORT_ARCHS)
if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.3)
list(APPEND DEEPGEMM_SUPPORT_ARCHS "9.0a")
endif()
if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8)
if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.9)
list(APPEND DEEPGEMM_SUPPORT_ARCHS "10.0f")
else()
list(APPEND DEEPGEMM_SUPPORT_ARCHS "10.0a")
endif()
if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
list(APPEND DEEPGEMM_SUPPORT_ARCHS "12.0f")
else()
list(APPEND DEEPGEMM_SUPPORT_ARCHS "12.0a" "12.1a")
endif()
if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.9)
list(APPEND DEEPGEMM_SUPPORT_ARCHS "10.0f")
elseif(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8)
list(APPEND DEEPGEMM_SUPPORT_ARCHS "10.0a")
endif()
cuda_archs_loose_intersection(DEEPGEMM_ARCHS
-73
View File
@@ -1,73 +0,0 @@
include(FetchContent)
# If FMHA_SM100_SRC_DIR is set, fmha_sm100 is installed from that directory
# instead of downloading. This is useful for local MSA development.
if(DEFINED ENV{FMHA_SM100_SRC_DIR})
set(FMHA_SM100_SRC_DIR $ENV{FMHA_SM100_SRC_DIR})
endif()
if(FMHA_SM100_SRC_DIR)
FetchContent_Declare(
fmha_sm100
SOURCE_DIR ${FMHA_SM100_SRC_DIR}
CONFIGURE_COMMAND ""
BUILD_COMMAND ""
)
else()
FetchContent_Declare(
fmha_sm100
GIT_REPOSITORY https://github.com/vllm-project/MSA.git
GIT_TAG fee783153f3efe57e3e933c5cb7e267a7cebcfb5
GIT_PROGRESS TRUE
CONFIGURE_COMMAND ""
BUILD_COMMAND ""
)
endif()
FetchContent_GetProperties(fmha_sm100)
if(NOT fmha_sm100_POPULATED)
FetchContent_Populate(fmha_sm100)
endif()
message(STATUS "fmha_sm100 is available at ${fmha_sm100_SOURCE_DIR}")
add_custom_target(fmha_sm100)
set(FMHA_SM100_PY_ROOT "${fmha_sm100_SOURCE_DIR}/python/fmha_sm100")
install(FILES
"${FMHA_SM100_PY_ROOT}/__init__.py"
"${FMHA_SM100_PY_ROOT}/api.py"
"${FMHA_SM100_PY_ROOT}/bench_utils.py"
"${FMHA_SM100_PY_ROOT}/jit.py"
"${FMHA_SM100_PY_ROOT}/sparse.py"
"${FMHA_SM100_PY_ROOT}/sparse_fmha_adapter.py"
DESTINATION vllm/third_party/fmha_sm100
COMPONENT fmha_sm100)
install(DIRECTORY "${FMHA_SM100_PY_ROOT}/csrc/"
DESTINATION vllm/third_party/fmha_sm100/csrc
COMPONENT fmha_sm100
PATTERN "__pycache__" EXCLUDE
PATTERN "*.pyc" EXCLUDE
PATTERN ".git*" EXCLUDE)
install(DIRECTORY "${FMHA_SM100_PY_ROOT}/cute/"
DESTINATION vllm/third_party/fmha_sm100/cute
COMPONENT fmha_sm100
PATTERN "__pycache__" EXCLUDE
PATTERN "*.pyc" EXCLUDE
PATTERN ".git*" EXCLUDE)
install(DIRECTORY "${FMHA_SM100_PY_ROOT}/cutlass/include/"
DESTINATION vllm/third_party/fmha_sm100/cutlass/include
COMPONENT fmha_sm100
PATTERN "__pycache__" EXCLUDE
PATTERN "*.pyc" EXCLUDE
PATTERN ".git*" EXCLUDE)
install(DIRECTORY "${FMHA_SM100_PY_ROOT}/cutlass/tools/util/include/"
DESTINATION vllm/third_party/fmha_sm100/cutlass/tools/util/include
COMPONENT fmha_sm100
PATTERN "__pycache__" EXCLUDE
PATTERN "*.pyc" EXCLUDE
PATTERN ".git*" EXCLUDE)
+32 -90
View File
@@ -6,83 +6,48 @@ if(DEFINED ENV{QUTLASS_SRC_DIR})
set(QUTLASS_SRC_DIR $ENV{QUTLASS_SRC_DIR})
endif()
# CMP0169 NEW: one-argument FetchContent_Populate(name) after Declare is invalid.
# Use explicit Populate(...) for git, or set SOURCE_DIR for local trees.
if(QUTLASS_SRC_DIR)
set(_qutlass_user_src "${QUTLASS_SRC_DIR}")
cmake_path(ABSOLUTE_PATH _qutlass_user_src
BASE_DIRECTORY "${CMAKE_SOURCE_DIR}"
NORMALIZE)
set(QUTLASS_SRC_DIR "${_qutlass_user_src}")
if(NOT IS_DIRECTORY "${QUTLASS_SRC_DIR}")
message(FATAL_ERROR
"[QUTLASS] QUTLASS_SRC_DIR is not an existing directory: '${QUTLASS_SRC_DIR}'")
endif()
set(qutlass_SOURCE_DIR "${QUTLASS_SRC_DIR}")
set(qutlass_BINARY_DIR "${CMAKE_BINARY_DIR}/qutlass-binary-dir-unused")
FetchContent_Declare(
qutlass
SOURCE_DIR ${QUTLASS_SRC_DIR}
CONFIGURE_COMMAND ""
BUILD_COMMAND ""
)
else()
set(_QUTLASS_UPSTREAM_REPO "https://github.com/IST-DASLab/qutlass.git")
set(_QUTLASS_UPSTREAM_TAG "830d2c4537c7396e14a02a46fbddd18b5d107c65")
set(_qutlass_fc_root "${FETCHCONTENT_BASE_DIR}")
if(NOT _qutlass_fc_root)
set(_qutlass_fc_root "${CMAKE_BINARY_DIR}/_deps")
endif()
set(_qutlass_src "${_qutlass_fc_root}/qutlass-src")
set(_qutlass_bin "${_qutlass_fc_root}/qutlass-build")
set(_qutlass_sub "${_qutlass_fc_root}/qutlass-subbuild")
if(EXISTS "${_qutlass_src}/qutlass/csrc/bindings.cpp")
set(qutlass_SOURCE_DIR "${_qutlass_src}")
set(qutlass_BINARY_DIR "${_qutlass_bin}")
else()
FetchContent_Populate(
qutlass
SUBBUILD_DIR "${_qutlass_sub}"
SOURCE_DIR "${_qutlass_src}"
BINARY_DIR "${_qutlass_bin}"
GIT_REPOSITORY "${_QUTLASS_UPSTREAM_REPO}"
GIT_TAG "${_QUTLASS_UPSTREAM_TAG}"
GIT_PROGRESS TRUE
)
endif()
FetchContent_Declare(
qutlass
GIT_REPOSITORY https://github.com/IST-DASLab/qutlass.git
GIT_TAG 830d2c4537c7396e14a02a46fbddd18b5d107c65
GIT_PROGRESS TRUE
CONFIGURE_COMMAND ""
BUILD_COMMAND ""
)
endif()
FetchContent_Populate(qutlass)
if(NOT qutlass_SOURCE_DIR)
message(FATAL_ERROR "[QUTLASS] source directory could not be resolved.")
endif()
message(STATUS "[QUTLASS] QuTLASS is available at ${qutlass_SOURCE_DIR}")
if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
cuda_archs_loose_intersection(QUTLASS_SM120_ARCHS "12.0f" "${CUDA_ARCHS}")
cuda_archs_loose_intersection(QUTLASS_SM100_ARCHS "10.0f" "${CUDA_ARCHS}")
cuda_archs_loose_intersection(QUTLASS_ARCHS "10.0f;12.0f" "${CUDA_ARCHS}")
else()
cuda_archs_loose_intersection(QUTLASS_SM120_ARCHS "12.0a;12.1a" "${CUDA_ARCHS}")
cuda_archs_loose_intersection(QUTLASS_SM100_ARCHS "10.0a;10.3a" "${CUDA_ARCHS}")
endif()
# QUTLASS uses TARGET_CUDA_ARCH as a single preprocessor selector for all its
# sources. Do not compile a mixed SM100/SM120 arch list with one selector; prefer
# SM100 when both families are requested because that is the primary deployed
# target for this extension today.
if(QUTLASS_SM100_ARCHS)
set(QUTLASS_ARCHS "${QUTLASS_SM100_ARCHS}")
set(QUTLASS_TARGET_CC 100)
if(QUTLASS_SM120_ARCHS)
message(WARNING
"[QUTLASS] Both SM100 and SM120 archs were requested; selecting SM100 "
"because TARGET_CUDA_ARCH is a single compile-time selector.")
endif()
elseif(QUTLASS_SM120_ARCHS)
set(QUTLASS_ARCHS "${QUTLASS_SM120_ARCHS}")
set(QUTLASS_TARGET_CC 120)
else()
set(QUTLASS_ARCHS)
cuda_archs_loose_intersection(QUTLASS_ARCHS "12.0a;12.1a;10.0a;10.3a" "${CUDA_ARCHS}")
endif()
if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND QUTLASS_ARCHS)
if(QUTLASS_ARCHS MATCHES "10\\.(0a|3a|0f)")
set(QUTLASS_TARGET_CC 100)
elseif(QUTLASS_ARCHS MATCHES "12\\.[01][af]?")
set(QUTLASS_TARGET_CC 120)
else()
message(FATAL_ERROR "[QUTLASS] internal error parsing CUDA_ARCHS='${QUTLASS_ARCHS}'.")
endif()
set(QUTLASS_SOURCES
csrc/qutlass_registration.cpp
${qutlass_SOURCE_DIR}/qutlass/csrc/bindings.cpp
${qutlass_SOURCE_DIR}/qutlass/csrc/gemm.cu
${qutlass_SOURCE_DIR}/qutlass/csrc/gemm_ada.cu
@@ -101,19 +66,8 @@ if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND QUTLASS_ARCHS)
if(CUTLASS_INCLUDE_DIR AND EXISTS "${CUTLASS_INCLUDE_DIR}/cutlass/cutlass.h")
list(APPEND QUTLASS_INCLUDES "${CUTLASS_INCLUDE_DIR}")
if(CUTLASS_TOOLS_UTIL_INCLUDE_DIR AND
EXISTS "${CUTLASS_TOOLS_UTIL_INCLUDE_DIR}/cutlass/util/packed_stride.hpp")
list(APPEND QUTLASS_INCLUDES "${CUTLASS_TOOLS_UTIL_INCLUDE_DIR}")
else()
get_filename_component(_qutlass_cutlass_root "${CUTLASS_INCLUDE_DIR}" DIRECTORY)
if(EXISTS "${_qutlass_cutlass_root}/tools/util/include/cutlass/util/packed_stride.hpp")
list(APPEND QUTLASS_INCLUDES "${_qutlass_cutlass_root}/tools/util/include")
endif()
endif()
elseif(EXISTS "${qutlass_SOURCE_DIR}/qutlass/third_party/cutlass/include/cutlass/cutlass.h")
list(APPEND QUTLASS_INCLUDES
"${qutlass_SOURCE_DIR}/qutlass/third_party/cutlass/include"
"${qutlass_SOURCE_DIR}/qutlass/third_party/cutlass/tools/util/include")
list(APPEND QUTLASS_INCLUDES "${qutlass_SOURCE_DIR}/qutlass/third_party/cutlass/include")
message(STATUS "[QUTLASS] Using QuTLASS vendored CUTLASS headers (no vLLM CUTLASS detected).")
else()
message(FATAL_ERROR "[QUTLASS] CUTLASS headers not found. "
@@ -125,23 +79,12 @@ if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND QUTLASS_ARCHS)
CUDA_ARCHS "${QUTLASS_ARCHS}"
)
# QuTLASS uses legacy ATen headers and cannot be built with TORCH_TARGET_VERSION.
# Keep it as its own extension (registers torch.ops._qutlass_C).
define_extension_target(
_qutlass_C
DESTINATION vllm
LANGUAGE ${VLLM_GPU_LANG}
SOURCES ${QUTLASS_SOURCES}
COMPILE_FLAGS ${VLLM_GPU_FLAGS}
ARCHITECTURES ${VLLM_GPU_ARCHES}
INCLUDE_DIRECTORIES ${QUTLASS_INCLUDES}
USE_SABI 3
WITH_SOABI)
target_compile_definitions(_qutlass_C PRIVATE
target_sources(_C PRIVATE ${QUTLASS_SOURCES})
target_include_directories(_C PRIVATE ${QUTLASS_INCLUDES})
target_compile_definitions(_C PRIVATE
QUTLASS_DISABLE_PYBIND=1
TARGET_CUDA_ARCH=${QUTLASS_TARGET_CC}
CUTLASS_ENABLE_DIRECT_CUDA_DRIVER_CALL=1)
)
set_property(SOURCE ${QUTLASS_SOURCES} APPEND PROPERTY COMPILE_OPTIONS
$<$<COMPILE_LANGUAGE:CUDA>:--expt-relaxed-constexpr --use_fast_math -O3>
@@ -156,5 +99,4 @@ else()
"[QUTLASS] Skipping build: no supported arch (12.0f / 10.0f) found in "
"CUDA_ARCHS='${CUDA_ARCHS}'.")
endif()
add_custom_target(_qutlass_C)
endif()

Some files were not shown because too many files have changed in this diff Show More