Compare commits

..
Author SHA1 Message Date
Zhuohan Li 0116e1cedc [Core] Enable extensible KV cache for all attention backends and Mamba
Port of internal D110967544. The extensible KV cache flow (reserve KV
virtual address space up front, capture CUDA graphs first, then size and
commit the KV cache from post-capture free memory) previously required a
block-major attention backend and rejected Mamba models. This enables it
for every backend layout and for Mamba / linear attention:

- ExtensibleTensor gains num_segments: the reservation is divided into
  equal segments that grow in lockstep, with committed bytes forming a
  prefix of each segment. Physical pages are mapped at
  allocation-granularity granules and deduped across overlapping ranges,
  so a granule straddling a segment boundary is mapped exactly once.
  resize_per_segment_(bytes, zero_new=True) zeroes only the newly
  committed logical range of each segment.
- Each KV cache buffer keeps its layers' physical layout and is committed
  as one prefix per layout segment. The segment count is derived from the
  backend's get_kv_cache_shape / get_kv_cache_block_dim / stride order:
  K/V-split layouts (e.g. FlashAttention) get one prefix per half,
  block-major layouts (e.g. FlashInfer, MLA) a single prefix. Mamba state
  pages are block-major per layer, and hybrid-model attention caches are
  re-strided to block-major, so both use a single segment.
- Removed the supports_extensible_kv_cache gate plumbing from EngineCore,
  Executor, Worker, WorkerBase and GPUModelRunner; a CUDA platform check
  remains in EngineCore.
- enable_extensible_kv_cache is reported as unsupported by the V2 model
  runner so V2-default models fall back to the V1 runner (which implements
  the flow); also fixed initialize_kv_cache being called with the
  extensible kwarg on runners that do not accept it, which broke every
  default V2-runner boot on this branch.

Tested on H100:
- tests/utils_/test_extensible_tensor.py (5 passed, incl. new segmented
  lockstep-grow/zero, granule-dedup and invalid-usage tests)
- tests/v1/worker/test_extensible_kv_cache.py (new, 6 passed: segment
  derivation, split grows both halves, block-major, legacy full commit,
  Mamba per-layer growth, hybrid attention+Mamba)
- E2E Qwen3-0.6B greedy with VLLM_ATTENTION_BACKEND=FLASH_ATTN (a K/V-split
  backend the old gate rejected): extensible generations byte-identical to
  the legacy path; log shows reserve then "Extended KV cache to 34663
  blocks". V2->V1 auto-fallback path verified as well.
2026-07-10 17:08:34 -07:00
Zhuohan Li 80f66afd81 add new files 2026-07-01 17:32:19 -07:00
Zhuohan Li 1fa4c3adb7 [Core] Demo implementation of extensible kv cache memory 2026-07-01 17:30:53 -07:00
1789 changed files with 24091 additions and 129430 deletions
-1
View File
@@ -3,7 +3,6 @@ job_dirs:
- ".buildkite/intel_jobs"
run_all_patterns:
- ".buildkite/ci_config_intel.yaml"
- ".buildkite/scripts/hardware_ci/run-intel-test.sh"
- "docker/Dockerfile"
- "docker/Dockerfile.xpu"
- "CMakeLists.txt"
-1
View File
@@ -8,7 +8,6 @@ run_all_patterns:
- "docker/docker-bake-rocm.hcl"
- ".buildkite/hardware_tests/amd.yaml"
- ".buildkite/scripts/ci-bake-rocm.sh"
- ".buildkite/scripts/rocm/"
- ".buildkite/scripts/hardware_ci/run-amd-test.py"
- ".buildkite/scripts/hardware_ci/run-amd-test.sh"
- "CMakeLists.txt"
+29 -38
View File
@@ -1,30 +1,5 @@
group: Hardware - AMD Build
# ROCm image flow:
# 1. Refresh the long-lived ROCm base image only when Dockerfile.rocm_base changes.
# 2. Build ci_base from either the stable base or the freshly refreshed base.
# 3. Build the per-commit ROCm CI image and smoke-test it before GPU jobs run.
steps:
- label: "AMD: :docker: refresh ROCm base"
key: refresh-rocm-base-amd
depends_on: []
device: amd_cpu
no_plugin: true
commands:
- bash .buildkite/scripts/rocm/refresh-base-image.sh
env:
DOCKER_BUILDKIT: "1"
BUILDKIT_PROGRESS: "tty"
TERM: "xterm-256color"
retry:
automatic:
- exit_status: 1 # Transient Docker/BuildKit failure
limit: 1
- exit_status: -1 # Agent was lost
limit: 1
- exit_status: -10 # Agent was lost
limit: 1
# 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
@@ -32,24 +7,19 @@ steps:
- label: "AMD: :docker: ensure ci_base"
key: ensure-ci-base-amd
soft_fail: false
depends_on:
- refresh-rocm-base-amd
depends_on: []
device: amd_cpu
no_plugin: true
commands:
- bash .buildkite/scripts/rocm/build-ci-base.sh
- bash .buildkite/scripts/ci-bake-rocm.sh ci-base-rocm-ci-with-deps
env:
DOCKER_BUILDKIT: "1"
BUILDKIT_PROGRESS: "tty"
TERM: "xterm-256color"
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 # Transient Docker/BuildKit failure
limit: 1
- exit_status: -1 # Agent was lost
limit: 1
- exit_status: -10 # Agent was lost
@@ -63,12 +33,35 @@ steps:
device: amd_cpu
no_plugin: true
commands:
- bash .buildkite/scripts/rocm/build-test-image.sh
- bash .buildkite/scripts/rocm/smoke-test-image.sh
- |
if [[ "${ROCM_CI_ARTIFACT_ONLY:-0}" == "1" ]]; then
echo "ROCM_CI_ARTIFACT_ONLY=1; building ROCm wheel artifact only"
IMAGE_TAG="" bash .buildkite/scripts/ci-bake-rocm.sh test-rocm-ci-with-artifacts
else
bash .buildkite/scripts/ci-bake-rocm.sh test-rocm-ci-with-wheel
fi
- |
docker run --rm --network=none --entrypoint /bin/bash "rocm/vllm-ci:${BUILDKITE_COMMIT}" -ec '
if [ ! -d /vllm-workspace ]; then echo Missing directory: /vllm-workspace >&2; exit 1; fi
if [ ! -d /vllm-workspace/tests ]; then echo Missing directory: /vllm-workspace/tests >&2; exit 1; fi
if [ ! -d /vllm-workspace/src/vllm ]; then echo Missing directory: /vllm-workspace/src/vllm >&2; exit 1; fi
if [ ! -x /vllm-workspace/src/vllm/vllm-rs ]; then echo Missing executable: /vllm-workspace/src/vllm/vllm-rs >&2; exit 1; fi
command -v python3
command -v uv
command -v pytest
if ! command -v amd-smi >/dev/null 2>&1 && ! command -v rocminfo >/dev/null 2>&1; then
echo No ROCm CLI found in image >&2
exit 1
fi
python3 - <<PY
import torch, vllm
print(torch.__version__)
print(vllm.__version__)
PY
echo AMD image smoke OK
'
env:
DOCKER_BUILDKIT: "1"
BUILDKIT_PROGRESS: "tty"
TERM: "xterm-256color"
VLLM_BAKE_FILE: "docker/docker-bake-rocm.hcl"
PYTORCH_ROCM_ARCH: "gfx90a;gfx942;gfx950"
IMAGE_TAG: "rocm/vllm-ci:$BUILDKITE_COMMIT"
@@ -76,8 +69,6 @@ steps:
VLLM_BRANCH: "$BUILDKITE_COMMIT"
retry:
automatic:
- exit_status: 1 # Transient Docker/BuildKit failure
limit: 1
- exit_status: -1 # Agent was lost
limit: 1
- exit_status: -10 # Agent was lost
+2 -21
View File
@@ -17,22 +17,16 @@ steps:
- tests/kernels/test_awq_int4_to_int8.py
- tests/kernels/quantization/test_cpu_fp8_scaled_mm.py
- tests/kernels/mamba/cpu/test_cpu_gdn_ops.py
- tests/kernels/mamba/test_cpu_short_conv.py
- tests/kernels/mamba/test_causal_conv1d.py
- tests/kernels/mamba/test_mamba_ssm.py
commands:
- |
bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 30m "
pytest -x -v -s tests/kernels/attention/test_cpu_attn.py
pytest -x -v -s tests/kernels/moe/test_cpu_fused_moe.py
pytest -x -v -s tests/kernels/moe/test_cpu_quant_fused_moe.py
pytest -x -v -s tests/kernels/mamba/test_cpu_short_conv.py
pytest -x -v -s tests/kernels/test_onednn.py
pytest -x -v -s tests/kernels/test_awq_int4_to_int8.py
pytest -x -v -s tests/kernels/quantization/test_cpu_fp8_scaled_mm.py
pytest -x -v -s tests/kernels/mamba/cpu/test_cpu_gdn_ops.py
pytest -x -v -s tests/kernels/mamba/test_causal_conv1d.py
pytest -x -v -s tests/kernels/mamba/test_mamba_ssm.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
@@ -145,22 +139,9 @@ steps:
commands:
- |
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 --ignore=tests/models/multimodal/generation/test_qwen2_5_vl.py -m cpu_model --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB"
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
- label: CPU-Qwen2.5-VL Multimodal Tests
depends_on: []
device: intel_cpu
no_plugin: true
source_file_dependencies:
# - vllm/
- vllm/model_executor/layers/rotary_embedding
- tests/models/multimodal/generation/
commands:
- |
bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 40m "
VLLM_CI_ENV=0 pytest -x -v -s tests/models/multimodal/generation/test_qwen2_5_vl.py"
- label: "Arm CPU Test"
depends_on: []
soft_fail: false
@@ -18,7 +18,7 @@ steps:
- label: "XPU example Test"
depends_on:
- image-build-xpu
timeout_in_minutes: 50
timeout_in_minutes: 30
optional: true
device: intel_gpu
agent_tags:
@@ -39,13 +39,13 @@ steps:
- label: "XPU V1 test"
depends_on:
- image-build-xpu
timeout_in_minutes: 70
timeout_in_minutes: 30
optional: true
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 24+
mem: 16+
no_plugin: true
env:
REGISTRY: "public.ecr.aws/q9t5s3a7"
@@ -60,7 +60,7 @@ steps:
- label: "XPU server test"
depends_on:
- image-build-xpu
timeout_in_minutes: 45
timeout_in_minutes: 30
optional: true
device: intel_gpu
agent_tags:
-20
View File
@@ -79,18 +79,12 @@ setup_buildx_builder() {
docker buildx ls | grep -E '^\*|^NAME' || docker buildx ls
}
annotate_image_tags() {
.buildkite/scripts/annotate-image-build.sh \
"${IMAGE_TAG:-}" "${IMAGE_TAG_LATEST:-}"
}
check_and_skip_if_image_exists() {
if [[ -n "${IMAGE_TAG:-}" ]]; then
echo "--- :mag: Checking if image exists"
if docker manifest inspect "${IMAGE_TAG}" >/dev/null 2>&1; then
echo "Image already exists: ${IMAGE_TAG}"
echo "Skipping build"
annotate_image_tags
exit 0
fi
echo "Image not found, proceeding with build"
@@ -177,18 +171,6 @@ BRANCH=$4
IMAGE_TAG=$5
IMAGE_TAG_LATEST=${6:-} # only used for main branch, optional
# When TORCH_NIGHTLY=1, build the base CI image against PyTorch nightly so the
# entire existing pipeline runs on nightly torch (CUDA/GPU lane only). Delegate
# to the dedicated nightly build (PYTORCH_NIGHTLY=1, CUDA 13.0) and tag it at the
# normal IMAGE_TAG that every test step already pulls -- no separate image tag,
# no duplicate "vLLM Against PyTorch Nightly" pipeline section.
if [[ "${TORCH_NIGHTLY:-0}" == "1" ]]; then
echo "--- :warning: TORCH_NIGHTLY=1 -- building base image on PyTorch nightly"
SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
exec "${SCRIPT_DIR}/image_build_torch_nightly.sh" \
"${REGISTRY}" "${REPO}" "${BUILDKITE_COMMIT}" "${BRANCH}" "${IMAGE_TAG}"
fi
# build config
TARGET="test-ci"
VLLM_BAKE_FILE_PATH="${VLLM_BAKE_FILE_PATH:-docker/docker-bake.hcl}"
@@ -272,5 +254,3 @@ echo "--- :docker: Building ${TARGET}"
docker --debug buildx bake -f "${VLLM_BAKE_FILE_PATH}" -f "${CI_HCL_PATH}" --progress plain "${TARGET}"
echo "--- :white_check_mark: Build complete"
annotate_image_tags
+19 -20
View File
@@ -9,31 +9,30 @@ fi
REGISTRY=$1
REPO=$2
BUILDKITE_COMMIT=$3
IMAGE="$REGISTRY/$REPO:$BUILDKITE_COMMIT-arm64"
# authenticate with AWS ECR
aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY" || true
# skip build if image already exists
if docker manifest inspect "$IMAGE" >/dev/null 2>&1; then
echo "Image found"
else
if [[ -z $(docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-arm64) ]]; then
echo "Image not found, proceeding with build..."
# build for arm64 GPU targets: Grace/GH200 (sm_90),
# Blackwell/Thor (sm_100/sm_103/sm_110), and DGX Spark/GB10
# (sm_121, family-covered by 12.0 under CUDA 13)
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 10.0 11.0 12.0" \
--build-arg USE_SCCACHE=1 \
--build-arg buildkite_commit="$BUILDKITE_COMMIT" \
--tag "$IMAGE" \
--target test \
--progress plain .
# push
docker push "$IMAGE"
else
echo "Image found"
exit 0
fi
.buildkite/scripts/annotate-image-build.sh "$IMAGE"
# build for arm64 GPU targets: Grace/GH200 (sm_90) and DGX Spark/GB10
# (sm_121, family-covered by 12.0 under CUDA 13)
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 USE_SCCACHE=1 \
--build-arg buildkite_commit="$BUILDKITE_COMMIT" \
--tag "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-arm64 \
--target test \
--progress plain .
# push
docker push "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-arm64
+15 -15
View File
@@ -9,26 +9,26 @@ fi
REGISTRY=$1
REPO=$2
BUILDKITE_COMMIT=$3
IMAGE="$REGISTRY/$REPO:$BUILDKITE_COMMIT-cpu"
# authenticate with AWS ECR
aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY" || true
# skip build if image already exists
if docker manifest inspect "$IMAGE" >/dev/null 2>&1; then
echo "Image found"
else
if [[ -z $(docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-cpu) ]]; then
echo "Image not found, proceeding with build..."
# build
docker build --file docker/Dockerfile.cpu \
--build-arg max_jobs=16 \
--build-arg buildkite_commit="$BUILDKITE_COMMIT" \
--build-arg VLLM_CPU_X86=true \
--tag "$IMAGE" \
--target vllm-test \
--progress plain .
# push
docker push "$IMAGE"
else
echo "Image found"
exit 0
fi
.buildkite/scripts/annotate-image-build.sh "$IMAGE"
# build
docker build --file docker/Dockerfile.cpu \
--build-arg max_jobs=16 \
--build-arg buildkite_commit="$BUILDKITE_COMMIT" \
--build-arg VLLM_CPU_X86=true \
--tag "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-cpu \
--target vllm-test \
--progress plain .
# push
docker push "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-cpu
+14 -14
View File
@@ -9,25 +9,25 @@ fi
REGISTRY=$1
REPO=$2
BUILDKITE_COMMIT=$3
IMAGE="$REGISTRY/$REPO:$BUILDKITE_COMMIT-arm64-cpu"
# authenticate with AWS ECR
aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY" || true
# skip build if image already exists
if docker manifest inspect "$IMAGE" >/dev/null 2>&1; then
echo "Image found"
else
if [[ -z $(docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-arm64-cpu) ]]; then
echo "Image not found, proceeding with build..."
# build
docker build --file docker/Dockerfile.cpu \
--build-arg max_jobs=16 \
--build-arg buildkite_commit="$BUILDKITE_COMMIT" \
--tag "$IMAGE" \
--target vllm-test \
--progress plain .
# push
docker push "$IMAGE"
else
echo "Image found"
exit 0
fi
.buildkite/scripts/annotate-image-build.sh "$IMAGE"
# build
docker build --file docker/Dockerfile.cpu \
--build-arg max_jobs=16 \
--build-arg buildkite_commit="$BUILDKITE_COMMIT" \
--tag "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-arm64-cpu \
--target vllm-test \
--progress plain .
# push
docker push "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-arm64-cpu
+15 -15
View File
@@ -9,26 +9,26 @@ fi
REGISTRY=$1
REPO=$2
BUILDKITE_COMMIT=$3
IMAGE="$REGISTRY/$REPO:$BUILDKITE_COMMIT-hpu"
# authenticate with AWS ECR
aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY" || true
# skip build if image already exists
if docker manifest inspect "$IMAGE" >/dev/null 2>&1; then
echo "Image found"
else
if [[ -z $(docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-hpu) ]]; then
echo "Image not found, proceeding with build..."
# build
docker build \
--file tests/pytorch_ci_hud_benchmark/Dockerfile.hpu \
--build-arg max_jobs=16 \
--build-arg buildkite_commit="$BUILDKITE_COMMIT" \
--tag "$IMAGE" \
--progress plain \
https://github.com/vllm-project/vllm-gaudi.git
# push
docker push "$IMAGE"
else
echo "Image found"
exit 0
fi
.buildkite/scripts/annotate-image-build.sh "$IMAGE"
# build
docker build \
--file tests/pytorch_ci_hud_benchmark/Dockerfile.hpu \
--build-arg max_jobs=16 \
--build-arg buildkite_commit="$BUILDKITE_COMMIT" \
--tag "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-hpu \
--progress plain \
https://github.com/vllm-project/vllm-gaudi.git
# push
docker push "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-hpu
@@ -40,7 +40,6 @@ docker buildx ls
echo "--- :mag: Checking if image already exists"
if docker manifest inspect "$IMAGE_TAG" >/dev/null 2>&1; then
echo "Image found: $IMAGE_TAG — skipping build"
.buildkite/scripts/annotate-image-build.sh "$IMAGE_TAG"
exit 0
fi
echo "Image not found, proceeding with build..."
@@ -67,5 +66,3 @@ docker buildx build --file docker/Dockerfile \
--progress plain .
echo "--- :white_check_mark: Torch nightly image build complete: $IMAGE_TAG"
.buildkite/scripts/annotate-image-build.sh "$IMAGE_TAG"
+14 -14
View File
@@ -9,26 +9,26 @@ fi
REGISTRY=$1
REPO=$2
BUILDKITE_COMMIT=$3
IMAGE="$REGISTRY/$REPO:$BUILDKITE_COMMIT-xpu"
# authenticate with AWS ECR
aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY" || true
aws ecr get-login-password --region us-east-1 | docker login --username AWS --password-stdin 936637512419.dkr.ecr.us-east-1.amazonaws.com || true
# skip build if image already exists
if docker manifest inspect "$IMAGE" &> /dev/null; then
echo "Image found"
else
if ! docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-xpu &> /dev/null; then
echo "Image not found, proceeding with build..."
# build
docker build \
--file docker/Dockerfile.xpu \
--build-arg max_jobs=16 \
--build-arg buildkite_commit="$BUILDKITE_COMMIT" \
--tag "$IMAGE" \
--progress plain .
# push
docker push "$IMAGE"
else
echo "Image found"
exit 0
fi
.buildkite/scripts/annotate-image-build.sh "$IMAGE"
# build
docker build \
--file docker/Dockerfile.xpu \
--build-arg max_jobs=16 \
--build-arg buildkite_commit="$BUILDKITE_COMMIT" \
--tag "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-xpu \
--progress plain .
# push
docker push "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-xpu
+1 -1
View File
@@ -3,7 +3,7 @@ depends_on:
- image-build-xpu
steps:
- label: XPU Sleep Mode
timeout_in_minutes: 45
timeout_in_minutes: 30
device: intel_gpu
agent_tags:
label: production
+2 -4
View File
@@ -81,12 +81,10 @@ steps:
'cd tests &&
export VLLM_WORKER_MULTIPROC_METHOD=spawn &&
set -o pipefail &&
pytest -v -s lora/test_punica_ops.py::test_kernels &&
pytest -v -s lora/test_punica_ops.py::test_kernels_hidden_size &&
pytest -v -s lora/test_punica_ops.py::test_add_lora_fused_moe_early_exit'
pytest -v -s lora/test_punica_ops.py --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[expand-0-xpu:0-dtype0-3-43264-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype1-1-2049-64-128-16]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-128-1-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-256-1-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-256-8-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels[expand-0-xpu:0-dtype0-3-2049-128-8-16]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-128-8-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels[expand-0-xpu:0-dtype1-1-2049-256-128-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype0-3-64256-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype1-2-29696-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype1-3-49408-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype0-2-16384-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[expand-0-xpu:0-dtype0-2-51328-32-4-4]" --deselect="tests/lora/test_kernels_hidden_size[shrink-0-xpu:0-dtype0-3-32000-32-4-4]" --deselect="tests/lora/test_kernels_hidden_size[shrink-0-xpu:0-dtype0-3-32000-32-4-4]"'
- label: LoRA Punica FP8/XPU Ops
timeout_in_minutes: 60
timeout_in_minutes: 45
device: intel_gpu
agent_tags:
label: production
+11 -33
View File
@@ -3,7 +3,7 @@ depends_on:
- image-build-xpu
steps:
- label: V1 Core + KV + Metrics
timeout_in_minutes: 45
timeout_in_minutes: 30
device: intel_gpu
agent_tags:
label: production
@@ -33,12 +33,12 @@ steps:
pytest -v -s v1/executor'
- label: V1 Sample + Logits
timeout_in_minutes: 90
timeout_in_minutes: 30
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 24+
mem: 16+
no_plugin: true
working_dir: "."
env:
@@ -72,31 +72,9 @@ steps:
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'
- label: Basic Models Tests (Initialization)
timeout_in_minutes: 60
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/test_initialization.py
- tests/models/registry.py
commands:
- >-
bash .buildkite/scripts/hardware_ci/run-intel-test.sh
'export VLLM_XPU_FUSED_MOE_USE_REF=1 &&
cd tests &&
pytest -v -s models/test_initialization.py::test_can_initialize_large_subset[Eagle3MiniMaxM2ForCausalLM]'
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
@@ -132,7 +110,7 @@ steps:
agent_tags:
label: production
gpu: 2+
mem: 16+
mem: 24+
no_plugin: true
working_dir: "."
env:
@@ -152,7 +130,7 @@ steps:
- label: Regression
key: regression
timeout_in_minutes: 50
timeout_in_minutes: 30
device: intel_gpu
agent_tags:
label: production
@@ -180,13 +158,13 @@ steps:
commands:
- >-
bash .buildkite/scripts/hardware_ci/run-intel-test.sh
'pip install modelscope\<1.38 &&
'pip install modelscope &&
cd tests &&
pytest -v -s test_regression.py'
- label: Metrics, Tracing (2 GPUs)
key: metrics-tracing-2-gpus
timeout_in_minutes: 45
timeout_in_minutes: 30
num_devices: 2
device: intel_gpu
agent_tags:
@@ -222,7 +200,7 @@ steps:
- label: Async Engine, Inputs, Utils, Worker
key: async-engine-inputs-utils-worker
timeout_in_minutes: 55
timeout_in_minutes: 30
device: intel_gpu
agent_tags:
label: production
@@ -4,12 +4,12 @@ depends_on:
steps:
- label: Distributed Model Tests (2 GPUs)
key: distributed-model-tests-2-gpus
timeout_in_minutes: 65
timeout_in_minutes: 50
device: intel_gpu
agent_tags:
label: production
gpu: 2+
mem: 16+
mem: 24+
no_plugin: true
working_dir: "."
env:
@@ -4,12 +4,12 @@ depends_on:
steps:
- label: "Multi-Modal Models (Standard) 1: qwen2"
key: multi-modal-models-standard-1-qwen2
timeout_in_minutes: 70
timeout_in_minutes: 45
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 24+
mem: 16+
no_plugin: true
working_dir: "."
env:
@@ -29,7 +29,7 @@ steps:
- label: "Multi-Modal Models (Standard) 2: qwen3 + gemma"
key: multi-modal-models-standard-2-qwen3-gemma
timeout_in_minutes: 70
timeout_in_minutes: 45
device: intel_gpu
agent_tags:
label: production
@@ -52,7 +52,7 @@ steps:
- label: "Multi-Modal Models (Standard) 3: llava + qwen2_vl"
key: multi-modal-models-standard-3-llava-qwen2-vl
timeout_in_minutes: 65
timeout_in_minutes: 45
device: intel_gpu
agent_tags:
label: production
@@ -81,7 +81,7 @@ steps:
agent_tags:
label: production
gpu: 1+
mem: 24+
mem: 16+
no_plugin: true
working_dir: "."
env:
@@ -100,12 +100,12 @@ steps:
- label: Multi-Modal Processor # 44min
key: multi-modal-processor
timeout_in_minutes: 60
timeout_in_minutes: 45
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 24+
mem: 16+
no_plugin: true
working_dir: "."
env:
-28
View File
@@ -1,28 +0,0 @@
group: Quantization
depends_on:
- image-build-xpu
steps:
- label: Quantization
key: quantization
timeout_in_minutes: 30
env:
REGISTRY: "public.ecr.aws/q9t5s3a7"
REPO: "vllm-ci-test-repo"
VLLM_TEST_DEVICE: "xpu"
no_plugin: true
working_dir: "."
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
source_file_dependencies:
- csrc/
- vllm/model_executor/layers/quantization
- tests/quantization
commands:
# - VLLM_TEST_FORCE_LOAD_FORMAT=auto pytest -v -s quantization/ --ignore quantization/test_blackwell_moe.py
- >-
bash .buildkite/scripts/hardware_ci/run-intel-test.sh
'VLLM_TEST_FORCE_LOAD_FORMAT=auto pytest -v -s tests/quantization/test_per_token_kv_cache.py --deselect="tests/quantization/test_per_token_kv_cache.py::test_triton_unified_attention_per_token_head_scale[int4-16-128-num_heads0-seq_lens1]"'
+5 -54
View File
@@ -17,7 +17,7 @@ steps:
- label: "XPU example Test"
depends_on:
- image-build-xpu
timeout_in_minutes: 50
timeout_in_minutes: 30
device: intel_gpu
agent_tags:
label: production
@@ -42,46 +42,21 @@ steps:
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 TheBloke/TinyLlama-1.1B-Chat-v0.3-AWQ --block-size 64 --enforce-eager &&
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
'
- label: "XPU W8A8 FP8 Linear Examples"
depends_on:
- image-build-xpu
timeout_in_minutes: 60
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 24+
no_plugin: true
env:
REGISTRY: "public.ecr.aws/q9t5s3a7"
REPO: "vllm-ci-test-repo"
VLLM_TEST_DEVICE: "xpu"
source_file_dependencies:
- vllm/
- .buildkite/intel_jobs/test-intel.yaml
commands:
- >-
bash .buildkite/scripts/hardware_ci/run-intel-test.sh
'python3 examples/basic/offline_inference/generate.py --linear-backend xpu --model RedHatAI/Meta-Llama-3.1-8B-Instruct-FP8 --enforce-eager --max-model-len 4096 &&
python3 examples/basic/offline_inference/generate.py --linear-backend xpu --model neuralmagic/Llama-3.2-1B-Instruct-FP8-dynamic --enforce-eager --max-model-len 4096 &&
python3 examples/basic/offline_inference/generate.py --linear-backend xpu --model meta-llama/Llama-3.2-1B-Instruct --quantization fp8 --enforce-eager --max-model-len 4096
'
- label: "XPU V1 test"
depends_on:
- image-build-xpu
timeout_in_minutes: 70
timeout_in_minutes: 30
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 24+
mem: 16+
no_plugin: true
env:
REGISTRY: "public.ecr.aws/q9t5s3a7"
@@ -99,13 +74,12 @@ steps:
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/e2e/general/test_correctness_sliding_window.py --deselect="tests/v1/e2e/general/test_correctness_sliding_window.py::test_sliding_window_retrieval[True-1-5-google/gemma-3-1b-it]" &&
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'
- label: "XPU server test"
depends_on:
- image-build-xpu
timeout_in_minutes: 45
timeout_in_minutes: 30
device: intel_gpu
agent_tags:
label: production
@@ -145,27 +119,4 @@ steps:
- >-
bash .buildkite/scripts/hardware_ci/run-intel-test.sh
'cd tests &&
pytest -v -s quantization/test_auto_round.py'
- label: "XPU compressed tensors FP8 test"
depends_on:
- image-build-xpu
timeout_in_minutes: 60
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"
VLLM_TEST_DEVICE: "xpu"
source_file_dependencies:
- vllm/
- tests/quantization/test_compressed_tensors.py
- .buildkite/intel_jobs/test-intel.yaml
commands:
- >-
bash .buildkite/scripts/hardware_ci/run-intel-test.sh
'cd tests &&
pytest -v -s quantization/test_compressed_tensors.py::test_compressed_tensors_fp8'
pytest -v -s quantization/test_auto_round.py'
@@ -1,7 +1,6 @@
# For hf script, without -t option (tensor parallel size).
# bash .buildkite/lm-eval-harness/run-lm-eval-mmlupro-vllm-baseline.sh -m meta-llama/Llama-4-Maverick-17B-128E-Instruct-FP8 -l 250 -t 8 -f 5
model_name: "meta-llama/Llama-4-Maverick-17B-128E-Instruct-FP8"
rocm_safetensors_load_strategy: lazy
required_gpu_arch:
- gfx942
- gfx950
@@ -72,11 +72,6 @@ def launch_lm_eval(eval_config, tp_size):
if moe_backend is not None:
model_args += f"moe_backend={moe_backend},"
if current_platform.is_rocm():
rocm_load_strategy = eval_config.get("rocm_safetensors_load_strategy")
if rocm_load_strategy is not None:
model_args += f"safetensors_load_strategy={rocm_load_strategy},"
env_vars = eval_config.get("env_vars", None)
with scoped_env_vars(env_vars):
results = lm_eval.simple_evaluate(
File diff suppressed because it is too large Load Diff
@@ -3,8 +3,7 @@
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
#
# Append a build artifact line to the Buildkite annotation.
# Usage: annotate-build-artifact.sh <label> <value> <context>
# Usage: annotate-build-artifact.sh <label> <value>
set -e
echo "- **${1}**: \`${2}\`" | \
buildkite-agent annotate --append --style 'info' \
--context "${3:?context is required}"
buildkite-agent annotate --append --style 'info' --context 'release-artifacts'
@@ -1,36 +0,0 @@
#!/bin/bash
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
#
# Append the Docker image tag(s) an image-build step pushed to a Buildkite
# annotation, so the built image tags show up on the build page instead of
# being buried in the job logs.
#
# Usage: annotate-image-build.sh <image_tag> [<image_tag> ...]
set -euo pipefail
# buildkite-agent only exists on Buildkite agents; no-op elsewhere so the
# image build scripts stay runnable locally.
if ! command -v buildkite-agent >/dev/null 2>&1; then
echo "buildkite-agent not found; skipping image tag annotation"
exit 0
fi
label="${BUILDKITE_LABEL:-Image build}"
content=""
for image in "$@"; do
[[ -n "$image" ]] || continue
content+="- **${label}**: \`${image}\`"$'\n'
done
if [[ -z "$content" ]]; then
echo "No image tags provided; nothing to annotate"
exit 0
fi
# Best-effort: a flaky annotation must never fail an otherwise successful
# (and expensive) image build.
if ! printf '%s' "$content" | \
buildkite-agent annotate --append --style 'info' --context 'docker-images'; then
echo "warning: failed to annotate build with image tags"
fi
-32
View File
@@ -1,32 +0,0 @@
#!/usr/bin/env bash
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
#
# Build the macOS arm64 CPU wheel natively on a macOS agent (the `macmini`
# queue) into artifacts/dist/ for upload-nightly-wheels.sh.
set -euo pipefail
# The Rust frontend build needs protoc.
if ! command -v protoc >/dev/null 2>&1; then
brew install protobuf
fi
# upload-nightly-wheels.sh expects exactly one wheel.
rm -rf artifacts/dist
mkdir -p artifacts/dist
export VLLM_TARGET_DEVICE=cpu
export VLLM_REQUIRE_RUST_FRONTEND=1
export MACOSX_DEPLOYMENT_TARGET=11.0
# uv's CPython is universal2; force an arm64-only build and tag so the wheel
# isn't mislabelled universal2 and installed on Intel Macs where import fails.
export ARCHFLAGS="-arch arm64"
export _PYTHON_HOST_PLATFORM="macosx-11.0-arm64"
export CMAKE_BUILD_PARALLEL_LEVEL="${CMAKE_BUILD_PARALLEL_LEVEL:-4}"
uv venv --python 3.12
uv pip install -r requirements/build/cpu.txt --index-strategy unsafe-best-match
uv build --wheel --no-build-isolation -o artifacts/dist
ls -l artifacts/dist/*.whl
@@ -29,11 +29,7 @@ if python3 -c "import torch; assert torch.version.hip" 2>/dev/null; then
TORCH_INDEX_URL=""
fi
else
if [ "${TORCH_NIGHTLY:-0}" = "1" ]; then
TORCH_INDEX_URL="https://download.pytorch.org/whl/nightly/cu130"
else
TORCH_INDEX_URL="https://download.pytorch.org/whl/cu130"
fi
TORCH_INDEX_URL="https://download.pytorch.org/whl/cu130"
fi
echo ">>> Using PyTorch index: ${TORCH_INDEX_URL:-PyPI default}"
+38 -301
View File
@@ -15,9 +15,9 @@ set -euo pipefail
DEFAULT_REPO_SLUG="vllm-project/vllm"
DEFAULT_CI_HCL_SOURCE="docker/ci-rocm.hcl"
DEFAULT_CI_BASE_CONTENT_FILES="requirements/common.txt requirements/rocm.txt requirements/test/rocm.txt docker/Dockerfile.rocm_base docker/ci-rocm.hcl docker/docker-bake-rocm.hcl tools/install_torchcodec_rocm.sh tools/install_protoc.sh rust-toolchain.toml tests/vllm_test_utils .buildkite/scripts/ci-bake-rocm.sh .buildkite/scripts/rocm/build-ci-base.sh"
DEFAULT_CI_BASE_CONTENT_FILES="requirements/common.txt requirements/rocm.txt requirements/test/rocm.txt docker/Dockerfile.rocm_base docker/ci-rocm.hcl docker/docker-bake-rocm.hcl tools/install_torchcodec_rocm.sh tests/vllm_test_utils .buildkite/scripts/ci-bake-rocm.sh"
DEFAULT_CI_BASE_DOCKERFILE="docker/Dockerfile.rocm"
DEFAULT_CI_BASE_DOCKERFILE_STAGES="base rust_toolchain_input_0 rust_toolchain_input_1 rust-toolchain-input rust-toolchain build_rixl build_rocshmem build_deepep mori_base ci_base"
DEFAULT_CI_BASE_DOCKERFILE_STAGES="base build_rixl build_rocshmem build_deepep mori_base ci_base"
DEFAULT_CI_BASE_METADATA_VERSION="1"
IMAGE_EXISTED_BEFORE_BUILD=0
@@ -285,7 +285,7 @@ get_content_arg_names() {
fi | awk 'NF && !seen[$0]++'
}
compute_ci_base_content_hash_once() {
compute_ci_base_content_hash() {
local -a content_paths=()
local -a content_args=()
local dockerfile="${CI_BASE_DOCKERFILE:-}"
@@ -301,8 +301,7 @@ compute_ci_base_content_hash_once() {
if [[ -n "${dockerfile}" ]]; then
printf 'dockerfile:%s\n' "${dockerfile}"
printf 'resolved-build-args:\n'
hash_dockerfile_arg_values "${dockerfile}" "${content_args[@]}" \
|| return 1
hash_dockerfile_arg_values "${dockerfile}" "${content_args[@]}"
if [[ -n "${stages}" ]]; then
printf 'dockerfile-stages:%s\n' "${stages}"
if [[ -f "${dockerfile}" ]]; then
@@ -315,53 +314,6 @@ compute_ci_base_content_hash_once() {
} | sha256sum | cut -d' ' -f1
}
compute_ci_base_content_hash() {
local attempts="${CI_BASE_HASH_ATTEMPTS:-3}"
local delay_secs="${CI_BASE_HASH_RETRY_DELAY:-5}"
local attempt=0
local hash=""
local failed=0
local -a hashes=()
if [[ ! "${attempts}" =~ ^[1-9][0-9]*$ ]]; then
echo "Invalid CI_BASE_HASH_ATTEMPTS: ${attempts}" >&2
return 1
fi
if [[ ! "${delay_secs}" =~ ^[0-9]+$ ]]; then
echo "Invalid CI_BASE_HASH_RETRY_DELAY: ${delay_secs}" >&2
return 1
fi
for ((attempt = 1; attempt <= attempts; attempt++)); do
if ! hash=$(compute_ci_base_content_hash_once); then
echo "ci_base content hash calculation ${attempt}/${attempts} failed" >&2
failed=1
else
hashes+=("${hash}")
echo "ci_base content hash calculation ${attempt}/${attempts}: ${hash}" >&2
fi
if ((attempt < attempts)); then
sleep "${delay_secs}"
fi
done
if ((failed)) || ((${#hashes[@]} != attempts)); then
echo "Could not calculate a reliable ci_base content hash" >&2
return 1
fi
for hash in "${hashes[@]:1}"; do
if [[ "${hash}" != "${hashes[0]}" ]]; then
echo "ci_base content hash changed between calculations" >&2
printf ' observed: %s\n' "${hashes[@]}" >&2
return 1
fi
done
printf '%s\n' "${hashes[0]}"
}
extract_dockerfile_arg_default() {
local dockerfile="$1"
local arg_name="$2"
@@ -414,11 +366,7 @@ hash_dockerfile_arg_values() {
printf 'arg:%s=%s\n' "${arg_name}" "${arg_value:-<empty>}"
if [[ "${arg_name}" == "BASE_IMAGE" && -n "${arg_value}" ]]; then
digest=$(resolve_image_digest "${arg_value}")
if [[ -z "${digest}" ]]; then
echo "Failed to resolve digest for BASE_IMAGE=${arg_value}" >&2
return 1
fi
printf 'arg:%s.digest=%s\n' "${arg_name}" "${digest}"
printf 'arg:%s.digest=%s\n' "${arg_name}" "${digest:-unknown}"
fi
done
}
@@ -445,16 +393,6 @@ should_upload_wheel_artifacts() {
|| "${TARGET}" == *"artifact"* ]]
}
set_buildkite_metadata() {
local key="$1"
local value="$2"
[[ -n "${value}" ]] || return 0
if command -v buildkite-agent >/dev/null 2>&1; then
buildkite-agent meta-data set "${key}" "${value}" || true
fi
}
get_remote_image_label() {
local image_ref="$1"
local label_key="$2"
@@ -795,8 +733,6 @@ configure_ci_base_image_refs() {
if is_ci_base_target; then
IMAGE_TAG="${primary_tag}"
CI_BASE_IMAGE="${primary_tag}"
export CI_BASE_IMAGE
export IMAGE_TAG
echo "ci_base primary image tag: ${CI_BASE_IMAGE_TAG}"
@@ -814,10 +750,6 @@ configure_ci_base_image_refs() {
echo "ci_base stable alias will not be pushed for this build"
echo "Set NIGHTLY=1 on ${CI_BASE_STABLE_BRANCH:-main} to refresh ${stable_tag}"
fi
set_buildkite_metadata "rocm-ci-base-image" "${CI_BASE_IMAGE_TAG}"
set_buildkite_metadata "rocm-ci-base-image-content" "${content_tag}"
set_buildkite_metadata "rocm-ci-base-image-commit" "${CI_BASE_IMAGE_TAG_COMMIT_REF:-}"
set_buildkite_metadata "rocm-ci-base-image-stable" "${CI_BASE_IMAGE_TAG_STABLE:-}"
return 0
fi
@@ -1263,24 +1195,12 @@ uses_rocm_csrc_cache() {
esac
}
uses_rocm_rust_cache() {
case "${TARGET}" in
rust-rocm-ci|test-rocm-ci|test-rocm-ci-with-wheel|test-rocm-ci-with-artifacts|export-wheel-rocm)
return 0
;;
*)
return 1
;;
esac
}
compute_rocm_csrc_content_hash() {
local bake_dir=""
local dockerfile_rocm=""
local -a content_paths=(
"requirements/common.txt"
"requirements/rocm.txt"
"pyproject.toml"
"setup.py"
"CMakeLists.txt"
"cmake"
@@ -1324,56 +1244,6 @@ compute_rocm_csrc_content_hash_if_needed() {
echo "ROCm csrc content cache ref: ${ROCM_CSRC_CONTENT_CACHE_REF}"
}
compute_rocm_rust_content_hash() {
local bake_dir=""
local dockerfile_rocm=""
local -a content_paths=(
"requirements/build/rust.txt"
"rust/Cargo.lock"
"rust/Cargo.toml"
"rust/proto"
"rust/src"
"rust-toolchain.toml"
"tools/build_rust.py"
"tools/install_protoc.sh"
"build_rust.sh"
)
local -a content_args=()
bake_dir=$(dirname "${VLLM_BAKE_FILE}")
dockerfile_rocm="${bake_dir}/Dockerfile.rocm"
mapfile -t content_args < <(
get_content_arg_names "${dockerfile_rocm}" "base rust_toolchain_input_0 rust_toolchain_input_1 rust-toolchain-input rust_input_0 rust_input_1 rust-input rust-toolchain rust-build" "${ROCM_RUST_CONTENT_ARGS:-}"
)
{
printf 'rust-input-files-hash:%s\n' "$(compute_content_hash "${content_paths[@]}")"
printf 'dockerfile:%s\n' "${dockerfile_rocm}"
printf 'resolved-build-args:\n'
hash_dockerfile_arg_values "${dockerfile_rocm}" "${content_args[@]}"
printf 'dockerfile-stages:base rust_toolchain_input_0 rust_toolchain_input_1 rust-toolchain-input rust_input_0 rust_input_1 rust-input rust-toolchain rust-build\n'
if [[ -f "${dockerfile_rocm}" ]]; then
hash_dockerfile_stages "${dockerfile_rocm}" "base rust_toolchain_input_0 rust_toolchain_input_1 rust-toolchain-input rust_input_0 rust_input_1 rust-input rust-toolchain rust-build"
else
printf 'missing:%s\n' "${dockerfile_rocm}"
fi
} | sha256sum | cut -d' ' -f1
}
compute_rocm_rust_content_hash_if_needed() {
local cache_repo="${DOCKERHUB_CACHE_REPO:-rocm/vllm-ci-cache}"
if [[ "${ROCM_RUST_CONTENT_CACHE:-1}" == "0" ]] || ! uses_rocm_rust_cache; then
return 0
fi
ROCM_RUST_CONTENT_HASH=$(compute_rocm_rust_content_hash)
ROCM_RUST_CONTENT_CACHE_REF="${cache_repo}:rust-rocm-input-${ROCM_RUST_CONTENT_HASH}"
export ROCM_RUST_CONTENT_HASH
export ROCM_RUST_CONTENT_CACHE_REF
echo "ROCm Rust content cache ref: ${ROCM_RUST_CONTENT_CACHE_REF}"
}
write_hcl_string_list_entries() {
local indent="$1"
local value=""
@@ -1431,7 +1301,6 @@ write_rocm_build_arg_override() {
"${CI_BASE_DOCKERFILE_STAGES:-${DEFAULT_CI_BASE_DOCKERFILE_STAGES}}" \
"${CI_BASE_CONTENT_ARGS:-}"
get_content_arg_names "${dockerfile_rocm}" "base csrc-build" "${ROCM_CSRC_CONTENT_ARGS:-}"
get_content_arg_names "${dockerfile_rocm}" "base rust_toolchain_input_0 rust_toolchain_input_1 rust-toolchain-input rust_input_0 rust_input_1 rust-input rust-toolchain rust-build" "${ROCM_RUST_CONTENT_ARGS:-}"
} | awk 'NF && !seen[$0]++'
)
@@ -1480,133 +1349,46 @@ validate_cache_export_mode() {
esac
}
validate_content_cache_export_mode() {
local mode="$1"
local env_name="$2"
case "${mode}" in
missing|always|never)
;;
*)
echo "Error: ${env_name} must be one of: missing, always, never"
exit 1
;;
esac
}
should_export_content_cache_ref() {
local cache_ref="$1"
local cache_name="$2"
local mode="${ROCM_CONTENT_CACHE_EXPORT_MODE:-missing}"
case "${mode}" in
always)
echo "${cache_name} content cache export mode is always; exporting ${cache_ref}"
return 0
;;
never)
echo "${cache_name} content cache export mode is never; not exporting ${cache_ref}"
return 1
;;
missing|"")
if docker buildx imagetools inspect "${cache_ref}" >/dev/null 2>&1; then
echo "${cache_name} content cache exists; not re-exporting ${cache_ref}"
return 1
fi
echo "${cache_name} content cache missing; will export ${cache_ref}"
return 0
;;
*)
echo "Error: ROCM_CONTENT_CACHE_EXPORT_MODE must be one of: missing, always, never"
exit 1
;;
esac
}
write_rocm_cache_override() {
local cache_repo="${DOCKERHUB_CACHE_REPO:-rocm/vllm-ci-cache}"
local content_cache_export_mode="${ROCM_CONTENT_CACHE_EXPORT_MODE:-missing}"
local csrc_cache_to_mode="${ROCM_CSRC_CACHE_TO_MODE:-max}"
local rust_cache_to_mode="${ROCM_RUST_CACHE_TO_MODE:-max}"
local rocm_cache_to_mode="${ROCM_FINAL_CACHE_TO_MODE:-min}"
local -a csrc_content_cache_from=()
local -a rust_content_cache_from=()
local -a combined_content_cache_from=()
local -a content_cache_from=()
local -a csrc_cache_to=()
local -a rust_cache_to=()
local -a rocm_cache_to=()
local -a export_wheel_cache_to=()
local export_csrc_cache=1
local export_rust_cache=1
if ! uses_rocm_csrc_cache && ! uses_rocm_rust_cache; then
if ! uses_rocm_csrc_cache; then
return 0
fi
validate_content_cache_export_mode \
"${content_cache_export_mode}" \
"ROCM_CONTENT_CACHE_EXPORT_MODE"
validate_cache_export_mode "${csrc_cache_to_mode}" "ROCM_CSRC_CACHE_TO_MODE"
validate_cache_export_mode "${rust_cache_to_mode}" "ROCM_RUST_CACHE_TO_MODE"
validate_cache_export_mode "${rocm_cache_to_mode}" "ROCM_FINAL_CACHE_TO_MODE"
echo "ROCm content cache export mode: ${content_cache_export_mode}"
echo "ROCm csrc cache export mode: ${csrc_cache_to_mode}"
echo "ROCm Rust cache export mode: ${rust_cache_to_mode}"
echo "ROCm final image cache export mode: ${rocm_cache_to_mode}"
if [[ -n "${ROCM_CSRC_CONTENT_CACHE_REF:-}" ]]; then
csrc_content_cache_from+=("type=registry,ref=${ROCM_CSRC_CONTENT_CACHE_REF}")
if should_export_content_cache_ref "${ROCM_CSRC_CONTENT_CACHE_REF}" "ROCm csrc"; then
csrc_cache_to+=(
"type=registry,ref=${ROCM_CSRC_CONTENT_CACHE_REF},mode=${csrc_cache_to_mode},ignore-error=true"
)
else
export_csrc_cache=0
fi
content_cache_from+=("type=registry,ref=${ROCM_CSRC_CONTENT_CACHE_REF}")
csrc_cache_to+=(
"type=registry,ref=${ROCM_CSRC_CONTENT_CACHE_REF},mode=${csrc_cache_to_mode},ignore-error=true"
)
fi
if [[ -n "${ROCM_RUST_CONTENT_CACHE_REF:-}" ]]; then
rust_content_cache_from+=("type=registry,ref=${ROCM_RUST_CONTENT_CACHE_REF}")
if should_export_content_cache_ref "${ROCM_RUST_CONTENT_CACHE_REF}" "ROCm Rust"; then
rust_cache_to+=(
"type=registry,ref=${ROCM_RUST_CONTENT_CACHE_REF},mode=${rust_cache_to_mode},ignore-error=true"
)
else
export_rust_cache=0
fi
fi
combined_content_cache_from=("${csrc_content_cache_from[@]}" "${rust_content_cache_from[@]}")
# Docker Hub cache exports are best-effort. A cache-only target failure can
# otherwise cancel the sibling image target before its manifest is pushed.
if [[ -n "${BUILDKITE_COMMIT:-}" ]]; then
if [[ ${export_csrc_cache} -eq 1 ]]; then
csrc_cache_to+=(
"type=registry,ref=${cache_repo}:csrc-rocm-${BUILDKITE_COMMIT},mode=${csrc_cache_to_mode},ignore-error=true"
)
fi
if [[ ${export_rust_cache} -eq 1 ]]; then
rust_cache_to+=(
"type=registry,ref=${cache_repo}:rust-rocm-${BUILDKITE_COMMIT},mode=${rust_cache_to_mode},ignore-error=true"
)
fi
csrc_cache_to+=(
"type=registry,ref=${cache_repo}:csrc-rocm-${BUILDKITE_COMMIT},mode=${csrc_cache_to_mode},ignore-error=true"
)
rocm_cache_to+=(
"type=registry,ref=${cache_repo}:rocm-${BUILDKITE_COMMIT},mode=${rocm_cache_to_mode},ignore-error=true"
)
fi
if [[ -n "${ROCM_CACHE_BRANCH_TAG:-}" ]]; then
if [[ ${export_csrc_cache} -eq 1 ]]; then
csrc_cache_to+=(
"type=registry,ref=${cache_repo}:csrc-rocm-branch-${ROCM_CACHE_BRANCH_TAG},mode=${csrc_cache_to_mode},ignore-error=true"
)
fi
if [[ ${export_rust_cache} -eq 1 ]]; then
rust_cache_to+=(
"type=registry,ref=${cache_repo}:rust-rocm-branch-${ROCM_CACHE_BRANCH_TAG},mode=${rust_cache_to_mode},ignore-error=true"
)
fi
csrc_cache_to+=(
"type=registry,ref=${cache_repo}:csrc-rocm-branch-${ROCM_CACHE_BRANCH_TAG},mode=${csrc_cache_to_mode},ignore-error=true"
)
rocm_cache_to+=(
"type=registry,ref=${cache_repo}:rocm-branch-${ROCM_CACHE_BRANCH_TAG},mode=${rocm_cache_to_mode},ignore-error=true"
)
@@ -1624,7 +1406,7 @@ target "csrc-rocm-ci" {
cache-from = concat(
get_cache_from_rocm_csrc(),
EOF
write_hcl_string_list " " "${csrc_content_cache_from[@]}"
write_hcl_string_list " " "${content_cache_from[@]}"
cat <<EOF
)
EOF
@@ -1632,23 +1414,11 @@ EOF
cat <<EOF
}
target "rust-rocm-ci" {
cache-from = concat(
get_cache_from_rocm_rust(),
EOF
write_hcl_string_list " " "${rust_content_cache_from[@]}"
cat <<EOF
)
EOF
write_hcl_string_list_attr " " "cache-to" "${rust_cache_to[@]}"
cat <<EOF
}
target "test-rocm-ci" {
cache-from = concat(
get_cache_from_rocm(),
EOF
write_hcl_string_list " " "${combined_content_cache_from[@]}"
write_hcl_string_list " " "${content_cache_from[@]}"
cat <<EOF
)
EOF
@@ -1660,7 +1430,7 @@ target "export-wheel-rocm" {
cache-from = concat(
get_cache_from_rocm(),
EOF
write_hcl_string_list " " "${combined_content_cache_from[@]}"
write_hcl_string_list " " "${content_cache_from[@]}"
cat <<EOF
)
EOF
@@ -1950,8 +1720,8 @@ confirm_remote_image_push() {
fi
if [[ -z "${remote_revision}" \
&& ${IMAGE_EXISTED_BEFORE_BUILD} -eq 0 ]] \
&& image_tag_is_commit_scoped; then
&& ${IMAGE_EXISTED_BEFORE_BUILD} -eq 0 \
&& image_tag_is_commit_scoped ]]; then
echo "Remote image exists under a commit-scoped tag; accepting push despite missing revision label."
return 0
fi
@@ -2009,10 +1779,7 @@ seed_dependency_caches_if_needed() {
echo "--- :docker: Seeding ${target}"
echo "Expected cache ref: ${cache_ref}"
docker buildx bake \
"${BAKE_FILES[@]}" \
--progress "${BUILDKIT_PROGRESS:-plain}" \
"${target}"
docker buildx bake "${BAKE_FILES[@]}" --progress plain "${target}"
verify_dependency_cache_ref "${cache_ref}"
done
}
@@ -2040,10 +1807,7 @@ run_bake() {
local build_rc=0
echo "--- :docker: Building ${TARGET}"
docker buildx bake \
"${BAKE_FILES[@]}" \
--progress "${BUILDKIT_PROGRESS:-plain}" \
"${BAKE_TARGETS[@]}" || build_rc=$?
docker buildx bake "${BAKE_FILES[@]}" --progress plain "${BAKE_TARGETS[@]}" || build_rc=$?
if [[ ${build_rc} -eq 0 ]]; then
echo "--- :white_check_mark: Build complete"
@@ -2081,57 +1845,36 @@ upload_wheel_artifacts_if_present() {
local wheel_dir="./wheel-export"
local artifact_dir="artifacts/vllm-rocm-install"
local archive_name="vllm-rocm-install.tar.gz"
local metadata_dir="${wheel_dir}/.vllm-ci-artifact"
local native_base_image=""
local whl=""
local whl_name=""
local -a wheels=()
if ! should_upload_wheel_artifacts; then
return 0
fi
if [[ -d "${wheel_dir}" ]]; then
mapfile -t wheels < <(find "${wheel_dir}" -maxdepth 1 -type f -name '*.whl' -print)
fi
if [[ ${#wheels[@]} -ne 1 ]]; then
echo "Expected exactly one ROCm wheel in ${wheel_dir}; found ${#wheels[@]}" >&2
return 1
fi
whl="${wheels[0]}"
whl_name=$(basename "${whl}")
native_base_image="${CI_BASE_IMAGE_TAG_COMMIT_REF:-${CI_BASE_IMAGE:-}}"
if [[ -z "${native_base_image}" ]]; then
echo "Native ROCm artifact requires a ci_base image reference" >&2
return 1
if [[ ! -d "${wheel_dir}" ]] || ! ls "${wheel_dir}"/*.whl >/dev/null 2>&1; then
echo "No ROCm wheel artifacts found in ${wheel_dir}"
return 0
fi
echo "--- :package: Uploading ROCm vLLM install artifact"
rm -rf "${artifact_dir}" "${metadata_dir}"
mkdir -p "${artifact_dir}" "${metadata_dir}"
printf '%s\n' "${BUILDKITE_COMMIT:-local}" > "${metadata_dir}/commit.txt"
printf '%s\n' "${native_base_image}" > "${metadata_dir}/native-base-image.txt"
printf '%s\n' "${CI_BASE_IMAGE:-}" > "${metadata_dir}/ci-base-image.txt"
printf '%s\n' "${IMAGE_TAG:-}" > "${metadata_dir}/fallback-image.txt"
printf '%s\n' "${whl_name}" > "${metadata_dir}/wheel-filename.txt"
mkdir -p "${artifact_dir}"
tar -C "${wheel_dir}" -czf "${artifact_dir}/${archive_name}" .
(
cd "${artifact_dir}"
sha256sum "${archive_name}" > "${archive_name}.sha256"
)
echo "Created ${archive_name}: $(du -sh "${artifact_dir}/${archive_name}" | cut -f1)"
cp "${metadata_dir}"/*.txt "${artifact_dir}/"
cp "${whl}" "${artifact_dir}/${whl_name}"
echo "Copied ${whl_name}: $(du -sh "${artifact_dir}/${whl_name}" | cut -f1)"
printf '%s\n' "${CI_BASE_IMAGE:-}" > "${artifact_dir}/ci-base-image.txt"
printf '%s\n' "${IMAGE_TAG:-}" > "${artifact_dir}/fallback-image.txt"
for whl in "${wheel_dir}"/*.whl; do
[[ -f "${whl}" ]] || continue
whl_name=$(basename "${whl}")
cp "${whl}" "${artifact_dir}/${whl_name}"
echo "Copied ${whl_name}: $(du -sh "${artifact_dir}/${whl_name}" | cut -f1)"
done
if command -v buildkite-agent >/dev/null 2>&1; then
buildkite-agent artifact upload "${artifact_dir}/*" || return 1
buildkite-agent artifact upload "${artifact_dir}/*"
echo "ROCm vLLM install artifacts uploaded to ${artifact_dir}/"
elif [[ "${BUILDKITE:-false}" == "true" ]]; then
echo "buildkite-agent not found; cannot upload required ROCm artifacts" >&2
return 1
else
echo "Not in Buildkite, skipping artifact upload"
fi
@@ -2155,7 +1898,6 @@ main() {
compute_dependency_cache_keys
write_ci_base_label_override
compute_rocm_csrc_content_hash_if_needed
compute_rocm_rust_content_hash_if_needed
write_rocm_cache_override
resolve_ci_base_dependency_targets
print_bake_config
@@ -2163,11 +1905,6 @@ main() {
echo "BAKE_PRINT_ONLY=1 set; skipping build"
return 0
fi
if should_upload_wheel_artifacts; then
# wheel-export is an output directory, not a BuildKit cache. Starting
# clean prevents a failed/retried export from packaging a stale wheel.
rm -rf ./wheel-export
fi
seed_dependency_caches_if_needed
run_bake
upload_wheel_artifacts_if_present
@@ -45,10 +45,8 @@ $PYTHON .buildkite/scripts/generate-nightly-index.py --version "$SUBPATH" --curr
echo "Uploading indices to $S3_COMMIT_PREFIX"
aws s3 cp --recursive "$INDICES_OUTPUT_DIR/" "$S3_COMMIT_PREFIX"
# copy to /nightly/ only when enabled for a main branch build that is not a PR
if [[ "${UPDATE_NIGHTLY_INDEX:-1}" == "1" && \
"$BUILDKITE_BRANCH" == "main" && \
"$BUILDKITE_PULL_REQUEST" == "false" ]]; then
# copy to /nightly/ only if it is on the main branch and not a PR
if [[ "$BUILDKITE_BRANCH" == "main" && "$BUILDKITE_PULL_REQUEST" == "false" ]]; then
echo "Uploading indices to overwrite /nightly/"
aws s3 cp --recursive "$INDICES_OUTPUT_DIR/" "s3://$BUCKET/nightly/"
fi
@@ -69,7 +67,7 @@ pure_version="${version%%+*}"
echo "Pure version (without variant): $pure_version"
# re-generate and copy to /<pure_version>/ only if it does not have "dev" in the version
if [[ "${UPDATE_VERSION_INDEX:-1}" == "1" && "$version" != *"dev"* ]]; then
if [[ "$version" != *"dev"* ]]; then
echo "Re-generating indices for /$pure_version/"
rm -rf "${INDICES_OUTPUT_DIR:?}"
mkdir -p "$INDICES_OUTPUT_DIR"
+21 -398
View File
@@ -1,7 +1,7 @@
#!/bin/bash
# This script runs ROCm tests either directly in a native CI pod or inside the
# corresponding Docker container. Multi-node tests continue to use Docker.
# This script runs tests inside the corresponding ROCm docker container.
# It handles both single-node and multi-node test configurations.
#
# Multi-node detection: Instead of matching on fragile group names, we detect
# multi-node jobs structurally by looking for the bracket command syntax
@@ -28,34 +28,6 @@
###############################################################################
set -o pipefail
: "${BUILDKIT_PROGRESS:=plain}"
: "${TERM:=xterm-256color}"
: "${FORCE_COLOR:=1}"
: "${CLICOLOR_FORCE:=1}"
: "${PY_COLORS:=1}"
: "${ROCM_DOCKER_TTY:=1}"
: "${PYTHONFAULTHANDLER:=1}"
: "${PYTEST_TIMEOUT:=2100}"
if [[ " ${PYTEST_ADDOPTS:-} " != *" --color"* ]]; then
PYTEST_ADDOPTS="${PYTEST_ADDOPTS:+${PYTEST_ADDOPTS} }--color=yes"
fi
if [[ " ${PYTEST_ADDOPTS:-} " != *" --durations="* ]]; then
PYTEST_ADDOPTS="${PYTEST_ADDOPTS:+${PYTEST_ADDOPTS} }--durations=25"
fi
if [[ " ${PYTEST_ADDOPTS:-} " != *" --durations-min="* ]]; then
PYTEST_ADDOPTS="${PYTEST_ADDOPTS:+${PYTEST_ADDOPTS} }--durations-min=1.0"
fi
# Dump stacks after 15 minutes, then stop an individual test after 35 minutes.
if [[ " ${PYTEST_ADDOPTS:-} " != *" faulthandler_timeout="* ]]; then
PYTEST_ADDOPTS="${PYTEST_ADDOPTS:+${PYTEST_ADDOPTS} }-o faulthandler_timeout=900"
fi
if [[ " ${PYTEST_ADDOPTS:-} " != *" --timeout-method="* &&
" ${PYTEST_ADDOPTS:-} " != *" --timeout-method "* ]]; then
PYTEST_ADDOPTS="${PYTEST_ADDOPTS:+${PYTEST_ADDOPTS} }--timeout-method=thread"
fi
export BUILDKIT_PROGRESS TERM FORCE_COLOR CLICOLOR_FORCE PY_COLORS PYTEST_ADDOPTS PYTEST_TIMEOUT ROCM_DOCKER_TTY
export PYTHONFAULTHANDLER
# 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.
@@ -70,28 +42,6 @@ report_docker_usage() {
docker system df || true
}
clear_ci_orchestration_env() {
unset -v \
VLLM_TEST_GROUP_NAME \
VLLM_CI_REQUIRE_PERSISTENT_HF_CACHE \
VLLM_CI_ARTIFACT_STEP \
VLLM_TEST_CACHE \
VLLM_CI_EXECUTION_MODE \
VLLM_CI_WORKSPACE \
VLLM_CI_REQUIRE_WORKSPACE_MOUNT \
VLLM_TEST_COMMANDS \
VLLM_CI_BRANCH \
VLLM_CI_BASE_IMAGE \
VLLM_CI_FALLBACK_IMAGE \
VLLM_CI_DOCKER_DISABLED \
VLLM_CI_ARTIFACT_GLOB \
VLLM_CI_ARTIFACT_CHECKSUM_GLOB \
VLLM_CI_EXPECTED_GPU_COUNT \
VLLM_CI_USE_ARTIFACTS \
VLLM_CI_RESULTS_ROOT \
VLLM_ALLOW_DEPRECATED_BEAM_SEARCH
}
cleanup_network() {
local max_nodes=${NUM_NODES:-2}
for node in $(seq 0 $((max_nodes - 1))); do
@@ -184,11 +134,7 @@ prepare_artifact_image() {
fi
cp "${wheel_dir}"/*.whl "${context_dir}/wheels/" || return 1
tar -C "${wheel_dir}" \
--exclude='*.whl' \
--exclude='.vllm-ci-artifact' \
--exclude='./.vllm-ci-artifact' \
-cf - . \
tar -C "${wheel_dir}" --exclude='*.whl' -cf - . \
| tar -C "${workspace_dir}" -xf - || return 1
cat > "${context_dir}/Dockerfile" <<'EOF'
ARG BASE_IMAGE
@@ -203,7 +149,6 @@ EOF
echo "--- Building local ROCm test image"
docker build \
--pull=false \
--progress "${BUILDKIT_PROGRESS}" \
--build-arg "BASE_IMAGE=${base_image}" \
-t "${artifact_image}" \
"${context_dir}" || return 1
@@ -211,257 +156,6 @@ EOF
return 0
}
is_native_runtime() {
[[ "${AMD_CI_RUNTIME:-}" == "native" || "${NATIVE_CI:-}" == "true" ]]
}
validate_native_workspace() {
local workspace_dir="${VLLM_CI_WORKSPACE:-/vllm-workspace}"
local workspace_real=""
local checkout_real=""
local workspace_mount=""
mkdir -p "${workspace_dir}" || return 1
workspace_real=$(readlink -m "${workspace_dir}") || return 1
if [[ -n "${BUILDKITE_BUILD_CHECKOUT_PATH:-}" ]]; then
checkout_real=$(readlink -m "${BUILDKITE_BUILD_CHECKOUT_PATH}") || return 1
if [[ "${checkout_real}" == "${workspace_real}" \
|| "${checkout_real}" == "${workspace_real}/"* \
|| "${workspace_real}" == "${checkout_real}/"* ]]; then
echo "Refusing to replace ${workspace_real}; it overlaps the Buildkite checkout ${checkout_real}" >&2
return 1
fi
fi
if [[ "${VLLM_CI_REQUIRE_WORKSPACE_MOUNT:-1}" == "1" ]]; then
if ! command -v findmnt >/dev/null 2>&1; then
echo "findmnt is required to verify the native workspace mount" >&2
return 1
fi
workspace_mount=$(findmnt -n -T "${workspace_real}" -o TARGET 2>/dev/null || true)
if [[ "$(readlink -m "${workspace_mount:-/}")" != "${workspace_real}" ]]; then
echo "Native CI requires a dedicated volume mounted at ${workspace_real}" >&2
return 1
fi
fi
}
prepare_native_workspace() {
if [[ "${VLLM_CI_USE_ARTIFACTS:-0}" != "1" ]]; then
echo "Native CI requires VLLM_CI_USE_ARTIFACTS=1"
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
validate_native_workspace || return 1
local artifact_glob="${VLLM_CI_ARTIFACT_GLOB:-artifacts/vllm-rocm-install/vllm-rocm-install.tar.gz}"
local artifact_checksum_glob="${VLLM_CI_ARTIFACT_CHECKSUM_GLOB:-${artifact_glob}.sha256}"
local artifact_step="${VLLM_CI_ARTIFACT_STEP:-image-build-amd}"
local archive=""
local checksum=""
local download_dir=""
local metadata_dir=""
local recorded_base=""
local recorded_commit=""
local recorded_wheel=""
local workspace_dir="${VLLM_CI_WORKSPACE:-/vllm-workspace}"
local wheel_dir=""
local attempt=0
local attempt_dir=""
local -a archives=()
local -a checksums=()
local -a wheels=()
artifact_work_dir=$(mktemp -d -t vllm-rocm-artifact.XXXXXX) || return 1
wheel_dir="${artifact_work_dir}/wheels"
mkdir -p "${wheel_dir}" || return 1
echo "--- Downloading ROCm wheel artifact from ${artifact_step} (native in-pod)"
for attempt in 1 2 3; do
attempt_dir="${artifact_work_dir}/download-${attempt}"
rm -rf "${attempt_dir}" || return 1
mkdir -p "${attempt_dir}" || return 1
if buildkite-agent artifact download \
"${artifact_glob}" "${attempt_dir}" --step "${artifact_step}" \
&& buildkite-agent artifact download \
"${artifact_checksum_glob}" "${attempt_dir}" --step "${artifact_step}"; then
download_dir="${attempt_dir}"
break
fi
echo "Artifact download attempt ${attempt}/3 failed"
if [[ "${attempt}" -lt 3 ]]; then
sleep $((attempt * 2))
fi
done
if [[ -z "${download_dir}" ]]; then
echo "Failed to download ${artifact_glob} and ${artifact_checksum_glob} from ${artifact_step}"
return 1
fi
mapfile -t archives < <(
find "${download_dir}" -name "vllm-rocm-install.tar.gz" -type f -print
)
mapfile -t checksums < <(
find "${download_dir}" -name "vllm-rocm-install.tar.gz.sha256" -type f -print
)
if [[ ${#archives[@]} -ne 1 || ${#checksums[@]} -ne 1 ]]; then
echo "Expected exactly one ROCm archive and checksum; found ${#archives[@]} archive(s) and ${#checksums[@]} checksum(s)" >&2
return 1
fi
archive="${archives[0]}"
checksum="${checksums[0]}"
if [[ "$(dirname "${archive}")" != "$(dirname "${checksum}")" ]]; then
echo "ROCm archive and checksum were downloaded to different directories" >&2
return 1
fi
(
cd "$(dirname "${archive}")"
sha256sum -c "$(basename "${checksum}")"
) || return 1
tar --no-same-owner -xzf "${archive}" -C "${wheel_dir}" || return 1
mapfile -t wheels < <(
find "${wheel_dir}" -maxdepth 1 -type f -name '*.whl' -print
)
if [[ ${#wheels[@]} -ne 1 ]]; then
echo "ROCm artifact must contain exactly one top-level wheel; found ${#wheels[@]}" >&2
return 1
fi
metadata_dir="${wheel_dir}/.vllm-ci-artifact"
for metadata_file in commit.txt native-base-image.txt wheel-filename.txt; do
if [[ ! -s "${metadata_dir}/${metadata_file}" ]]; then
echo "ROCm artifact metadata is missing ${metadata_file}" >&2
return 1
fi
done
for metadata_file in ci-base-image.txt fallback-image.txt; do
if [[ ! -f "${metadata_dir}/${metadata_file}" ]]; then
echo "ROCm artifact metadata is missing ${metadata_file}" >&2
return 1
fi
done
recorded_commit=$(tr -d '\r\n' < "${metadata_dir}/commit.txt")
recorded_base=$(tr -d '\r\n' < "${metadata_dir}/native-base-image.txt")
recorded_wheel=$(tr -d '\r\n' < "${metadata_dir}/wheel-filename.txt")
if [[ -z "${BUILDKITE_COMMIT:-}" || "${recorded_commit}" != "${BUILDKITE_COMMIT}" ]]; then
echo "ROCm artifact commit ${recorded_commit} does not match ${BUILDKITE_COMMIT:-unset}" >&2
return 1
fi
if [[ -z "${VLLM_CI_BASE_IMAGE:-}" || "${recorded_base}" != "${VLLM_CI_BASE_IMAGE}" ]]; then
echo "ROCm artifact base ${recorded_base} does not match ${VLLM_CI_BASE_IMAGE:-unset}" >&2
return 1
fi
if [[ "${recorded_wheel}" != "$(basename "${wheels[0]}")" ]]; then
echo "ROCm artifact wheel manifest ${recorded_wheel} does not match $(basename "${wheels[0]}")" >&2
return 1
fi
for required_dir in tests .buildkite requirements; do
if [[ ! -d "${wheel_dir}/${required_dir}" ]]; then
echo "ROCm wheel artifact did not contain ${required_dir}/" >&2
return 1
fi
done
echo "--- Installing ROCm wheel into pod environment"
python3 -m pip install --no-deps --force-reinstall "${wheels[0]}" || return 1
echo "--- Preparing ${workspace_dir} from artifact"
find "${workspace_dir}" -mindepth 1 -maxdepth 1 -exec rm -rf -- {} + || return 1
tar -C "${wheel_dir}" \
--exclude='*.whl' \
--exclude='.vllm-ci-artifact' \
--exclude='./.vllm-ci-artifact' \
-cf - . | tar --no-same-owner -C "${workspace_dir}" -xf - || return 1
if [[ ! -d "${workspace_dir}/tests" ]]; then
echo "Failed to stage the native test workspace" >&2
return 1
fi
return 0
}
initialize_native_environment() {
local job_id="${BUILDKITE_JOB_ID:-${BUILDKITE_PARALLEL_JOB:-local}}"
local job_id_suffix=""
local native_root=""
local hf_mount=""
if [[ "$(id -u)" -ne 0 ]]; then
echo "Native ROCm CI currently requires the ci_base container to run as root" >&2
return 1
fi
job_id="${job_id//[^A-Za-z0-9_.-]/_}"
job_id_suffix="${job_id##*-}"
job_id_suffix="${job_id_suffix:0:12}"
native_root="/tmp/vllm-native-${job_id}"
TMPDIR="/tmp/vllm-${job_id_suffix}/tmp"
VLLM_RPC_BASE_PATH="/tmp"
: "${TORCHINDUCTOR_CACHE_DIR:=${native_root}/cache/torchinductor}"
: "${TRITON_CACHE_DIR:=${native_root}/cache/triton}"
: "${VLLM_CACHE_ROOT:=${native_root}/cache/vllm}"
: "${XDG_CACHE_HOME:=${native_root}/cache/xdg}"
: "${HF_HOME:=/home/buildkite-agent/huggingface}"
: "${HF_HUB_DOWNLOAD_TIMEOUT:=300}"
: "${HF_HUB_ETAG_TIMEOUT:=60}"
export TMPDIR VLLM_RPC_BASE_PATH
export TORCHINDUCTOR_CACHE_DIR TRITON_CACHE_DIR VLLM_CACHE_ROOT XDG_CACHE_HOME
export HF_HOME HF_HUB_DOWNLOAD_TIMEOUT HF_HUB_ETAG_TIMEOUT
export PYTORCH_ROCM_ARCH=""
mkdir -p "${TMPDIR}" \
"${TORCHINDUCTOR_CACHE_DIR}" \
"${TRITON_CACHE_DIR}" \
"${VLLM_CACHE_ROOT}" \
"${XDG_CACHE_HOME}" \
"${HF_HOME}" || return 1
if [[ "${VLLM_CI_REQUIRE_PERSISTENT_HF_CACHE:-0}" == "1" ]]; then
if ! command -v findmnt >/dev/null 2>&1; then
echo "findmnt is required to verify the native Hugging Face cache mount" >&2
return 1
fi
hf_mount=$(findmnt -n -T "${HF_HOME}" -o TARGET 2>/dev/null || true)
if [[ -z "${hf_mount}" || "${hf_mount}" == "/" ]]; then
echo "Native CI requires a persistent volume mounted at or above ${HF_HOME}" >&2
return 1
fi
fi
}
run_native_preflight() {
local expected_gpus="${VLLM_CI_EXPECTED_GPU_COUNT:-1}"
if [[ ! "${expected_gpus}" =~ ^[0-9]+$ ]]; then
echo "Invalid VLLM_CI_EXPECTED_GPU_COUNT=${expected_gpus}" >&2
return 1
fi
python3 -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" || return 1
if [[ "${expected_gpus}" == "0" ]]; then
echo "Native CPU-only AMD job: skipping ROCm device validation"
return 0
fi
echo "--- ROCm info"
rocminfo || return 1
VLLM_CI_EXPECTED_GPU_COUNT="${expected_gpus}" python3 - <<'PY'
import os
import torch
expected = int(os.environ["VLLM_CI_EXPECTED_GPU_COUNT"])
assert torch.version.hip, "PyTorch is not a ROCm build"
assert torch.cuda.is_available(), "ROCm GPU is not available to PyTorch"
actual = torch.cuda.device_count()
assert actual == expected, f"Expected {expected} ROCm GPU(s), found {actual}"
PY
}
is_multi_node() {
local cmds="$1"
# Primary signal: NUM_NODES environment variable set by the pipeline
@@ -644,58 +338,7 @@ re_quote_pytest_markers() {
# Main
###############################################################################
if is_native_runtime; then
echo "--- Native in-pod ROCm CI (AMD_CI_RUNTIME=${AMD_CI_RUNTIME:-unset}, NATIVE_CI=${NATIVE_CI:-unset})"
artifact_work_dir=""
cleanup_native_workspace() {
if [[ -n "${artifact_work_dir}" ]]; then
rm -rf "${artifact_work_dir}"
fi
}
trap cleanup_native_workspace EXIT
if [[ -n "${VLLM_TEST_COMMANDS:-}" ]]; then
commands="${VLLM_TEST_COMMANDS}"
commands_source="env"
else
commands="$*"
commands_source="argv"
if [[ -z "$commands" ]]; then
echo "Error: No test commands provided for native CI." >&2
exit 1
fi
fi
if [[ "$commands_source" == "argv" ]]; then
commands=$(re_quote_pytest_markers "$commands")
fi
if is_multi_node "$commands"; then
echo "Native CI does not support multi-node jobs yet."
exit 1
fi
if ! initialize_native_environment; then
echo "Failed to initialize the native test environment"
exit 1
fi
if ! prepare_native_workspace; then
echo "Failed to prepare native test workspace"
exit 1
fi
export PYTHONPATH="${VLLM_CI_WORKSPACE:-/vllm-workspace}"
echo "Native test commands: $commands"
run_native_preflight || exit 1
# Keep AMD CI orchestration variables out of vLLM's runtime environment.
clear_ci_orchestration_env
/bin/bash -o pipefail -c "${commands}"
handle_pytest_exit "$?"
fi
# --- GPU initialization for legacy Docker execution ---
# --- GPU initialization ---
echo "--- ROCm info"
rocminfo
@@ -797,27 +440,23 @@ fi
echo "Final commands: $commands"
standalone_merge_base_env=()
if [[ "$commands" == *python_only_compile.sh* ]]; then
# The ROCm test image often ships /vllm-workspace without .git. Resolve the
# wheels.vllm.ai commit from the agent checkout for this test only.
vllm_standalone_merge_base=""
checkout="${BUILDKITE_BUILD_CHECKOUT_PATH:-}"
if [[ -z "${checkout}" || ! -d "${checkout}" ]]; then
checkout="."
fi
# Pass safe.directory per-command because Buildkite uses mixed user IDs.
if git -c "safe.directory=${checkout}" -C "${checkout}" rev-parse --is-inside-work-tree >/dev/null 2>&1; then
vllm_standalone_merge_base="$(
git -c "safe.directory=${checkout}" -C "${checkout}" merge-base HEAD origin/main 2>/dev/null || true
)"
fi
if [[ -z "${vllm_standalone_merge_base}" ]]; then
vllm_standalone_merge_base="${BUILDKITE_COMMIT:-}"
fi
echo "INFO: passing CI_STANDALONE_MERGE_BASE into container: ${vllm_standalone_merge_base}"
standalone_merge_base_env=(-e "CI_STANDALONE_MERGE_BASE=${vllm_standalone_merge_base}")
# The ROCm test image often ships /vllm-workspace without .git (artifact tarball unpack).
# tests/standalone_tests/python_only_compile.sh uses merge-base(HEAD, origin/main) for
# wheels.vllm.ai; compute on the agent (full git checkout) and pass into the container.
vllm_standalone_merge_base=""
checkout="${BUILDKITE_BUILD_CHECKOUT_PATH:-}"
if [[ -z "${checkout}" || ! -d "${checkout}" ]]; then
checkout="."
fi
if git -C "${checkout}" rev-parse --is-inside-work-tree >/dev/null 2>&1; then
vllm_standalone_merge_base="$(
git -C "${checkout}" merge-base HEAD origin/main 2>/dev/null || true
)"
fi
if [[ -z "${vllm_standalone_merge_base}" ]]; then
vllm_standalone_merge_base="${BUILDKITE_COMMIT:-}"
fi
echo "INFO: passing VLLM_STANDALONE_MERGE_BASE into container: ${vllm_standalone_merge_base}"
MYPYTHONPATH="/vllm-workspace"
@@ -848,7 +487,6 @@ else
fi
# --- Route: multi-node vs single-node ---
clear_ci_orchestration_env
if is_multi_node "$commands"; then
echo "--- Multi-node job detected"
export DCKR_VER=$(docker --version | sed 's/Docker version \(.*\), build .*/\1/')
@@ -895,13 +533,6 @@ if is_multi_node "$commands"; then
else
echo "--- Single-node job"
echo "Render devices: $BUILDKITE_AGENT_META_DATA_RENDER_DEVICES"
docker_run_terminal_args=(-i)
if [[ "${ROCM_DOCKER_TTY}" == "1" ]]; then
docker_run_terminal_args+=(-t)
echo "Docker interactive stdin: enabled; TTY allocation: enabled"
else
echo "Docker interactive stdin: enabled; TTY allocation: disabled"
fi
ulimit_core_hard=$(ulimit -H -c)
if [[ "$ulimit_core_hard" == "unlimited" ]]; then
@@ -918,7 +549,6 @@ else
fi
docker run \
"${docker_run_terminal_args[@]}" \
--device /dev/kfd $BUILDKITE_AGENT_META_DATA_RENDER_DEVICES \
$RDMA_FLAGS \
--network=host \
@@ -933,13 +563,6 @@ else
-e AWS_SECRET_ACCESS_KEY \
-e BUILDKITE_PARALLEL_JOB \
-e BUILDKITE_PARALLEL_JOB_COUNT \
-e TERM \
-e FORCE_COLOR \
-e CLICOLOR_FORCE \
-e PY_COLORS \
-e PYTHONFAULTHANDLER \
-e PYTEST_ADDOPTS \
-e PYTEST_TIMEOUT \
-v "${HF_CACHE}:${HF_MOUNT}" \
-e "HF_HOME=${HF_MOUNT}" \
-e "PYTHONPATH=${MYPYTHONPATH}" \
@@ -949,7 +572,7 @@ else
-e "VLLM_CACHE_ROOT=${CONTAINER_CACHE_ROOT}/vllm" \
-e "XDG_CACHE_HOME=${CONTAINER_CACHE_ROOT}/xdg" \
-e "PYTORCH_ROCM_ARCH=" \
"${standalone_merge_base_env[@]}" \
-e "VLLM_STANDALONE_MERGE_BASE=${vllm_standalone_merge_base}" \
--name "${container_name}" \
"${image_name}" \
/bin/bash -c "${CONTAINER_PREFLIGHT} && ${commands}"
@@ -38,11 +38,7 @@ function cpu_tests() {
pytest -x -v -s tests/kernels/attention/test_cpu_attn.py
pytest -x -v -s tests/kernels/core/test_cpu_activation.py
pytest -x -v -s tests/kernels/moe/test_cpu_fused_moe.py
pytest -x -v -s tests/kernels/mamba/cpu/test_cpu_gdn_ops.py
pytest -x -v -s tests/kernels/moe/test_cpu_int4_moe.py
pytest -x -v -s tests/kernels/mamba/test_cpu_short_conv.py
pytest -x -v -s tests/kernels/mamba/test_causal_conv1d.py
pytest -x -v -s tests/kernels/mamba/test_mamba_ssm.py"
pytest -x -v -s tests/kernels/mamba/cpu/test_cpu_gdn_ops.py"
# skip tests requiring model downloads if HF_TOKEN is not set
# due to rate-limits
@@ -66,6 +62,7 @@ function cpu_tests() {
set -e
pytest -x -v -s tests/quantization/test_compressed_tensors.py::test_compressed_tensors_w8a8_logprobs"
# basic online serving
docker exec cpu-test bash -c '
set -e
@@ -99,4 +96,3 @@ function cpu_tests() {
# All of CPU tests are expected to be finished less than 40 mins.
export -f cpu_tests
timeout 2h bash -c cpu_tests
@@ -21,7 +21,6 @@ case "${test_suite}" in
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 TheBloke/TinyLlama-1.1B-Chat-v0.3-AWQ --block-size 64 --enforce-eager
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
@@ -35,7 +34,6 @@ case "${test_suite}" in
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/e2e/general/test_correctness_sliding_window.py --deselect="tests/v1/e2e/general/test_correctness_sliding_window.py::test_sliding_window_retrieval[True-1-5-google/gemma-3-1b-it]"
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
;;
@@ -360,7 +360,7 @@ export HF_TOKEN ZE_AFFINITY_MASK
--ipc=host \
--privileged \
-v /dev/dri/by-path:/dev/dri/by-path \
-v "/data/huggingface:/root/.cache/huggingface" \
-v "${HOME}/.cache/huggingface:/root/.cache/huggingface" \
--entrypoint='' \
-e HF_TOKEN \
-e ZE_AFFINITY_MASK \
@@ -85,7 +85,7 @@ RUN pip config set global.index-url http://cache-service-vllm.nginx-pypi-cache.s
# Install for pytest to make the docker build cache layer always valid
RUN --mount=type=cache,target=/root/.cache/pip \
pip install pytest>=6.0 'modelscope<1.38'
pip install pytest>=6.0 modelscope
WORKDIR /workspace/vllm
@@ -130,22 +130,6 @@ docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${ROCM_BASE_CACHE_KEY}-rocm
docker push vllm/vllm-openai-rocm:latest-base
docker push vllm/vllm-openai-rocm:v${RELEASE_VERSION}-base
# ---- XPU ----
docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-xpu
docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-xpu vllm/vllm-openai-xpu:latest-x86_64
docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-xpu vllm/vllm-openai-xpu:v${RELEASE_VERSION}-x86_64
docker push vllm/vllm-openai-xpu:latest-x86_64
docker push vllm/vllm-openai-xpu:v${RELEASE_VERSION}-x86_64
docker manifest rm vllm/vllm-openai-xpu:latest || true
docker manifest rm vllm/vllm-openai-xpu:v${RELEASE_VERSION} || true
docker manifest create vllm/vllm-openai-xpu:latest vllm/vllm-openai-xpu:latest-x86_64 --amend
docker manifest create vllm/vllm-openai-xpu:v${RELEASE_VERSION} vllm/vllm-openai-xpu:v${RELEASE_VERSION}-x86_64 --amend
docker manifest push vllm/vllm-openai-xpu:latest
docker manifest push vllm/vllm-openai-xpu:v${RELEASE_VERSION}
# ---- CPU ----
# CPU images are behind separate block steps and may not have been built.
# All-or-nothing: inspect both arches first, then either publish everything
-32
View File
@@ -1,32 +0,0 @@
#!/usr/bin/env bash
# Build the ROCm ci_base image, optionally from a freshly rebuilt ROCm base.
set -euo pipefail
metadata_get() {
local key="$1"
if command -v buildkite-agent >/dev/null 2>&1; then
buildkite-agent meta-data get "${key}" 2>/dev/null || true
fi
}
main() {
local base_refreshed=""
base_refreshed="$(metadata_get rocm-base-refresh)"
if [[ "${base_refreshed}" == "1" ]]; then
export BASE_IMAGE
export CI_BASE_PUSH_STABLE_TAG
BASE_IMAGE="$(metadata_get rocm-base-image)"
CI_BASE_PUSH_STABLE_TAG="$(metadata_get rocm-base-push-stable-tag)"
CI_BASE_PUSH_STABLE_TAG="${CI_BASE_PUSH_STABLE_TAG:-0}"
echo "Using refreshed ROCm base image for ci_base: ${BASE_IMAGE}"
echo "Push stable ci_base tag: ${CI_BASE_PUSH_STABLE_TAG}"
fi
bash .buildkite/scripts/ci-bake-rocm.sh ci-base-rocm-ci-with-deps
}
main "$@"
@@ -1,68 +0,0 @@
#!/usr/bin/env bash
# Build the ROCm CI test image or wheel artifact.
#
# When Dockerfile.rocm_base changes, always build the full image so downstream
# ROCm tests can validate the freshly rebuilt base -> ci_base -> ci image chain.
set -euo pipefail
metadata_get() {
local key="$1"
if command -v buildkite-agent >/dev/null 2>&1; then
buildkite-agent meta-data get "${key}" 2>/dev/null || true
fi
}
use_ci_base_if_present() {
local ci_base_image=""
ci_base_image="$(metadata_get rocm-ci-base-image)"
if [[ -z "${ci_base_image}" ]]; then
return 1
fi
export CI_BASE_IMAGE="${ci_base_image}"
echo "Using ROCm ci_base image selected by the preceding build step: ${CI_BASE_IMAGE}"
}
use_refreshed_base_if_present() {
local base_refreshed=""
base_refreshed="$(metadata_get rocm-base-refresh)"
if [[ "${base_refreshed}" != "1" ]]; then
return 1
fi
export BASE_IMAGE
export IMAGE_TAG_LATEST
BASE_IMAGE="$(metadata_get rocm-base-image)"
IMAGE_TAG_LATEST="$(metadata_get rocm-ci-image-descriptive)"
echo "Using refreshed ROCm base image for test image: ${BASE_IMAGE}"
if [[ -n "${IMAGE_TAG_LATEST}" ]]; then
echo "Also tagging full ROCm CI image as: ${IMAGE_TAG_LATEST}"
fi
return 0
}
main() {
local base_refreshed=0
use_ci_base_if_present || true
if use_refreshed_base_if_present; then
base_refreshed=1
fi
if [[ "${ROCM_CI_ARTIFACT_ONLY:-0}" == "1" && "${base_refreshed}" != "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
return
fi
bash .buildkite/scripts/ci-bake-rocm.sh test-rocm-ci-with-wheel
}
main "$@"
@@ -1,513 +0,0 @@
#!/usr/bin/env bash
# Build and publish a fresh ROCm base image when Dockerfile.rocm_base changes.
#
# Normal AMD CI builds should not pay for this path. The script no-ops unless
# docker/Dockerfile.rocm_base changed relative to the branch base, the previous
# main commit, or ROCM_BASE_REFRESH_FORCE=1 is set.
set -euo pipefail
DOCKERFILE="${ROCM_BASE_DOCKERFILE:-docker/Dockerfile.rocm_base}"
BASE_REPO="${ROCM_BASE_IMAGE_REPO:-rocm/vllm-dev}"
CI_IMAGE_REPO="${ROCM_CI_IMAGE_REPO:-rocm/vllm-ci}"
BUILDER_NAME="${ROCM_BASE_BUILDER_NAME:-vllm-rocm-base-builder}"
DEFAULT_ROCM_BASE_METADATA_VERSION="1"
DEFAULT_ROCM_BASE_CONTENT_FILES="${DOCKERFILE}"
DEFAULT_ROCM_BASE_CONTENT_ARGS="BASE_IMAGE TRITON_BRANCH TRITON_REPO PYTORCH_BRANCH PYTORCH_REPO PYTORCH_VISION_BRANCH PYTORCH_VISION_REPO PYTORCH_AUDIO_BRANCH PYTORCH_AUDIO_REPO FA_BRANCH FA_REPO AITER_BRANCH AITER_REPO MORI_BRANCH MORI_REPO PYTORCH_ROCM_ARCH PYTHON_VERSION USE_SCCACHE"
metadata_set() {
local key="$1"
local value="$2"
[[ -n "${value}" ]] || return 0
if command -v buildkite-agent >/dev/null 2>&1; then
buildkite-agent meta-data set "${key}" "${value}" || true
fi
}
compute_content_hash() {
local path=""
local file=""
for path in "$@"; do
if [[ -d "${path}" ]]; then
while IFS= read -r -d '' file; do
printf 'file:%s\n' "${file}"
sha256sum "${file}"
done < <(find "${path}" -type f -print0 | sort -z)
elif [[ -f "${path}" ]]; then
printf 'file:%s\n' "${path}"
sha256sum "${path}"
else
printf 'missing:%s\n' "${path}"
fi
done | sha256sum | cut -d' ' -f1
}
clean_docker_tag() {
local input="$1"
echo "${input}" | sed 's/[^a-zA-Z0-9._-]/_/g' | cut -c1-128
}
tag_component() {
local input="$1"
local max_chars="${2:-24}"
clean_docker_tag "${input:-unknown}" | cut -c1-"${max_chars}"
}
extract_arg_default() {
local arg_name="$1"
sed -n -E "s/^[[:space:]]*ARG[[:space:]]+${arg_name}=\"?([^\"[:space:]]+)\"?.*/\\1/p" \
"${DOCKERFILE}" | head -1
}
resolve_image_digest() {
local image_ref="$1"
docker buildx imagetools inspect "${image_ref}" 2>/dev/null \
| sed -n -E 's/^Digest:[[:space:]]+//p' \
| head -1 || true
}
resolve_rocm_base_arg_value() {
local arg_name="$1"
local use_sccache="$2"
case "${arg_name}" in
USE_SCCACHE)
printf '%s\n' "${use_sccache}"
;;
*)
extract_arg_default "${arg_name}"
;;
esac
}
hash_rocm_base_arg_values() {
local use_sccache="$1"
local base_image_digest="$2"
local arg_name=""
local arg_value=""
shift 2 || true
for arg_name in "$@"; do
[[ -n "${arg_name}" ]] || continue
arg_value=$(resolve_rocm_base_arg_value "${arg_name}" "${use_sccache}")
printf 'arg:%s=%s\n' "${arg_name}" "${arg_value:-<empty>}"
if [[ "${arg_name}" == "BASE_IMAGE" && -n "${arg_value}" ]]; then
printf 'arg:%s.digest=%s\n' "${arg_name}" "${base_image_digest:-unknown}"
fi
done
}
rocm_version_from_base_image() {
local base_image="$1"
local version=""
version="$(sed -n -E 's/.*:([0-9]+\.[0-9]+(\.[0-9]+)?)-.*/\1/p' <<<"${base_image}")"
tag_component "${version:-${base_image}}" 16
}
git_diff_changed_base() {
local range="$1"
[[ -n "$(git diff --name-only "${range}" -- "${DOCKERFILE}" 2>/dev/null)" ]]
}
short_git_ref() {
local ref="$1"
git rev-parse --short "${ref}" 2>/dev/null || printf '%s\n' "${ref}"
}
extract_arg_default_from_ref() {
local ref="$1"
local arg_name="$2"
local content=""
content="$(git show "${ref}:${DOCKERFILE}" 2>/dev/null || true)"
sed -n -E "s/^[[:space:]]*ARG[[:space:]]+${arg_name}=\"?([^\"[:space:]]+)\"?.*/\\1/p" \
<<<"${content}" | head -1
}
log_arg_default_changes() {
local old_ref="$1"
local new_ref="$2"
local content_args="${ROCM_BASE_CONTENT_ARGS:-${DEFAULT_ROCM_BASE_CONTENT_ARGS}}"
local arg_name=""
local old_value=""
local new_value=""
local changed=0
echo "Changed ROCm base ARG defaults:"
for arg_name in ${content_args}; do
old_value="$(extract_arg_default_from_ref "${old_ref}" "${arg_name}")"
new_value="$(extract_arg_default_from_ref "${new_ref}" "${arg_name}")"
if [[ "${old_value}" != "${new_value}" ]]; then
echo " - ${arg_name}: ${old_value:-<unset>} -> ${new_value:-<unset>}"
changed=1
fi
done
if [[ "${changed}" == "0" ]]; then
echo " - none detected; Dockerfile instructions changed outside tracked ARG defaults"
fi
}
log_arg_line_diff() {
local range="$1"
local arg_diff=""
arg_diff="$(
git diff --unified=0 "${range}" -- "${DOCKERFILE}" 2>/dev/null \
| awk '/^[+-][[:space:]]*ARG[[:space:]]/ && $0 !~ /^(---|\+\+\+)/ { print " " $0 }' \
|| true
)"
if [[ -n "${arg_diff}" ]]; then
echo "Changed Dockerfile ARG lines:"
printf '%s\n' "${arg_diff}"
fi
}
log_rocm_base_change_check() {
local context="$1"
local range="$2"
local old_ref="$3"
local old_short=""
local head_short=""
old_short="$(short_git_ref "${old_ref}")"
head_short="$(short_git_ref HEAD)"
echo "--- :mag: ROCm base refresh check"
echo "Context: ${context}"
echo "Dockerfile: ${DOCKERFILE}"
echo "Base revision: ${old_short}"
echo "Head revision: ${head_short}"
echo "Git diff range: ${range}"
}
log_rocm_base_rebuild_reason() {
local context="$1"
local range="$2"
local old_ref="$3"
local changed_files=""
log_rocm_base_change_check "${context}" "${range}" "${old_ref}"
changed_files="$(git diff --name-only "${range}" -- "${DOCKERFILE}" 2>/dev/null || true)"
echo "Changed files:"
if [[ -n "${changed_files}" ]]; then
sed 's/^/ - /' <<<"${changed_files}"
else
echo " - ${DOCKERFILE}"
fi
log_arg_default_changes "${old_ref}" HEAD
log_arg_line_diff "${range}"
echo "Decision: rebuilding ROCm base image because ${DOCKERFILE} changed."
}
rocm_base_changed_in_range() {
local context="$1"
local range="$2"
local old_ref="$3"
if git_diff_changed_base "${range}"; then
log_rocm_base_rebuild_reason "${context}" "${range}" "${old_ref}"
return 0
fi
log_rocm_base_change_check "${context}" "${range}" "${old_ref}"
echo "Decision: ROCm base refresh not required; ${DOCKERFILE} is unchanged."
return 1
}
rocm_base_changed() {
local base_branch="${BUILDKITE_PULL_REQUEST_BASE_BRANCH:-main}"
local base_ref="refs/remotes/origin/${base_branch}"
local merge_base=""
if [[ "${ROCM_BASE_REFRESH_SKIP:-0}" == "1" ]]; then
echo "ROCM_BASE_REFRESH_SKIP=1 set; skipping ROCm base refresh"
return 1
fi
if [[ "${ROCM_BASE_REFRESH_FORCE:-0}" == "1" ]]; then
echo "ROCM_BASE_REFRESH_FORCE=1 set; refreshing ROCm base image"
return 0
fi
if ! git rev-parse --is-inside-work-tree >/dev/null 2>&1; then
echo "Not in a git checkout; skipping ROCm base refresh unless forced"
return 1
fi
if [[ "${BUILDKITE_PULL_REQUEST:-false}" != "false" ]]; then
git fetch --no-tags --depth=200 origin \
"+refs/heads/${base_branch}:${base_ref}" >/dev/null 2>&1 || true
merge_base=$(git merge-base HEAD "${base_ref}" 2>/dev/null || true)
if [[ -z "${merge_base}" ]]; then
echo "Unable to determine merge base with PR base ${base_ref}; skipping ROCm base refresh unless forced"
return 1
fi
if rocm_base_changed_in_range \
"pull request build against ${base_ref}" \
"${merge_base}...HEAD" \
"${merge_base}"; then
return 0
fi
elif [[ "${BUILDKITE_BRANCH:-}" == "${ROCM_BASE_STABLE_BRANCH:-main}" ]] \
&& git rev-parse --verify HEAD~1 >/dev/null 2>&1; then
if rocm_base_changed_in_range \
"stable branch build; comparing against previous ${ROCM_BASE_STABLE_BRANCH:-main} commit" \
"HEAD~1..HEAD" \
"HEAD~1"; then
return 0
fi
else
git fetch --no-tags --depth=200 origin \
"+refs/heads/${base_branch}:${base_ref}" >/dev/null 2>&1 || true
merge_base=$(git merge-base HEAD "${base_ref}" 2>/dev/null || true)
if [[ -z "${merge_base}" ]]; then
echo "Unable to determine merge base with branch base ${base_ref}; skipping ROCm base refresh unless forced"
return 1
fi
if rocm_base_changed_in_range \
"branch build against ${base_ref}" \
"${merge_base}...HEAD" \
"${merge_base}"; then
return 0
fi
fi
return 1
}
should_push_stable_tag() {
if [[ "${BUILDKITE_PULL_REQUEST:-false}" != "false" ]]; then
return 1
fi
if [[ "${ROCM_BASE_PUSH_STABLE_TAG:-}" == "1" ]]; then
return 0
fi
if [[ "${ROCM_BASE_PUSH_STABLE_TAG:-}" == "0" ]]; then
return 1
fi
[[ "${BUILDKITE_PULL_REQUEST:-false}" == "false" \
&& "${BUILDKITE_BRANCH:-}" == "${ROCM_BASE_STABLE_BRANCH:-main}" ]]
}
setup_builder() {
echo "--- :buildkite: Setting up buildx builder for ROCm base"
if docker buildx inspect "${BUILDER_NAME}" >/dev/null 2>&1; then
docker buildx use "${BUILDER_NAME}"
else
docker buildx create --name "${BUILDER_NAME}" --driver docker-container --use
fi
docker buildx inspect --bootstrap
}
compute_base_content_hash() {
local use_sccache="$1"
local base_image_digest="$2"
local content_files="${ROCM_BASE_CONTENT_FILES:-${DEFAULT_ROCM_BASE_CONTENT_FILES}}"
local content_args="${ROCM_BASE_CONTENT_ARGS:-${DEFAULT_ROCM_BASE_CONTENT_ARGS}}"
local -a content_paths=()
local -a content_arg_names=()
read -r -a content_paths <<< "${content_files}"
read -r -a content_arg_names <<< "${content_args}"
{
printf 'content-files-hash:%s\n' "$(compute_content_hash "${content_paths[@]}")"
printf 'dockerfile:%s\n' "${DOCKERFILE}"
printf 'resolved-build-args:\n'
hash_rocm_base_arg_values \
"${use_sccache}" "${base_image_digest}" "${content_arg_names[@]}"
} | sha256sum | cut -d' ' -f1
}
build_base_image() {
local use_sccache="${ROCM_BASE_USE_SCCACHE:-${USE_SCCACHE:-0}}"
local base_hash=""
local build_date=""
local build_suffix=""
local base_image_arg=""
local base_image_digest=""
local rocm_version=""
local triton_arg=""
local pytorch_arg=""
local pytorch_vision_arg=""
local pytorch_audio_arg=""
local fa_arg=""
local aiter_arg=""
local mori_arg=""
local python_version_arg=""
local pytorch_rocm_arch_arg=""
local pytorch_branch=""
local aiter_branch=""
local dependency_summary=""
local descriptor=""
local ci_descriptor=""
local descriptive_tag=""
local stable_tag="${BASE_REPO}:base"
local ci_descriptive_tag=""
local content_files="${ROCM_BASE_CONTENT_FILES:-${DEFAULT_ROCM_BASE_CONTENT_FILES}}"
local content_args="${ROCM_BASE_CONTENT_ARGS:-${DEFAULT_ROCM_BASE_CONTENT_ARGS}}"
local content_files_hash=""
local metadata_version="${ROCM_BASE_METADATA_VERSION:-${DEFAULT_ROCM_BASE_METADATA_VERSION}}"
local -a tags=()
local -a no_cache_args=()
local -a sccache_args=()
local -a content_paths=()
if [[ ! -f "${DOCKERFILE}" ]]; then
echo "Error: ROCm base Dockerfile not found: ${DOCKERFILE}" >&2
exit 1
fi
build_date="${ROCM_BASE_TAG_DATE:-$(date -u +%Y%m%d)}"
if [[ -n "${BUILDKITE_BUILD_NUMBER:-}" ]]; then
build_suffix="_bk_${BUILDKITE_BUILD_NUMBER}"
fi
base_image_arg="$(extract_arg_default BASE_IMAGE)"
base_image_digest="$(resolve_image_digest "${base_image_arg}")"
read -r -a content_paths <<< "${content_files}"
content_files_hash="$(compute_content_hash "${content_paths[@]}")"
base_hash=$(compute_base_content_hash "${use_sccache}" "${base_image_digest}")
rocm_version="$(rocm_version_from_base_image "${base_image_arg}")"
triton_arg="$(extract_arg_default TRITON_BRANCH)"
pytorch_arg="$(extract_arg_default PYTORCH_BRANCH)"
pytorch_vision_arg="$(extract_arg_default PYTORCH_VISION_BRANCH)"
pytorch_audio_arg="$(extract_arg_default PYTORCH_AUDIO_BRANCH)"
fa_arg="$(extract_arg_default FA_BRANCH)"
aiter_arg="$(extract_arg_default AITER_BRANCH)"
mori_arg="$(extract_arg_default MORI_BRANCH)"
python_version_arg="$(extract_arg_default PYTHON_VERSION)"
pytorch_rocm_arch_arg="$(extract_arg_default PYTORCH_ROCM_ARCH)"
pytorch_branch="$(tag_component "${pytorch_arg}" 16)"
aiter_branch="$(tag_component "${aiter_arg}" 24)"
dependency_summary="base=${base_image_arg},rocm=${rocm_version},python=${python_version_arg},pytorch=${pytorch_arg},torchvision=${pytorch_vision_arg},torchaudio=${pytorch_audio_arg},triton=${triton_arg},flash-attn=${fa_arg},aiter=${aiter_arg},mori=${mori_arg},pytorch-rocm-arch=${pytorch_rocm_arch_arg}"
descriptor="$(clean_docker_tag "base_custom_aiter_${aiter_branch}_torch_${pytorch_branch}_${build_date}${build_suffix}")"
ci_descriptor="$(clean_docker_tag "ci_custom_aiter_${aiter_branch}_torch_${pytorch_branch}_${build_date}${build_suffix}")"
descriptive_tag="${BASE_REPO}:${descriptor}"
ci_descriptive_tag="${CI_IMAGE_REPO}:${ci_descriptor}"
tags=(-t "${descriptive_tag}")
if should_push_stable_tag; then
tags+=(-t "${stable_tag}")
metadata_set "rocm-base-push-stable-tag" "1"
else
metadata_set "rocm-base-push-stable-tag" "0"
fi
if [[ "${ROCM_BASE_NO_CACHE:-1}" == "1" ]]; then
no_cache_args=(--no-cache)
fi
for env_name in \
SCCACHE_DOWNLOAD_URL \
SCCACHE_ENDPOINT \
SCCACHE_BUCKET_NAME \
SCCACHE_REGION_NAME \
SCCACHE_S3_NO_CREDENTIALS; do
if [[ -n "${!env_name:-}" ]]; then
sccache_args+=(--build-arg "${env_name}=${!env_name}")
fi
done
echo "--- :docker: Building ROCm base image"
echo "Dockerfile: ${DOCKERFILE}"
echo "Descriptive tag: ${descriptive_tag}"
echo "Stable tag: ${stable_tag} ($(should_push_stable_tag && echo enabled || echo disabled))"
echo "Content hash: ${base_hash}"
echo "Dependency summary: ${dependency_summary}"
echo "USE_SCCACHE: ${use_sccache}"
docker buildx build \
"${no_cache_args[@]}" \
--pull \
--progress "${BUILDKIT_PROGRESS:-plain}" \
--file "${DOCKERFILE}" \
--build-arg "USE_SCCACHE=${use_sccache}" \
"${sccache_args[@]}" \
--label "org.opencontainers.image.source=https://github.com/vllm-project/vllm" \
--label "org.opencontainers.image.vendor=vLLM" \
--label "org.opencontainers.image.title=vLLM ROCm base" \
--label "org.opencontainers.image.revision=${BUILDKITE_COMMIT:-}" \
--label "vllm.rocm_base.metadata_version=${metadata_version}" \
--label "vllm.rocm_base.content_hash=${base_hash}" \
--label "vllm.rocm_base.content_files_hash=${content_files_hash}" \
--label "vllm.rocm_base.dockerfile=${DOCKERFILE}" \
--label "vllm.rocm_base.image.descriptive=${descriptive_tag}" \
--label "vllm.rocm_base.image.stable=${stable_tag}" \
--label "vllm.rocm_base.git_commit=${BUILDKITE_COMMIT:-}" \
--label "vllm.rocm_base.stable_branch=${ROCM_BASE_STABLE_BRANCH:-main}" \
--label "vllm.rocm_base.descriptor=${descriptor}" \
--label "vllm.rocm_base.dependency_summary=${dependency_summary}" \
--label "vllm.rocm_base.base_image=${base_image_arg}" \
--label "vllm.rocm_base.base_image_digest=${base_image_digest}" \
--label "vllm.rocm_base.dependency.rocm=${rocm_version}" \
--label "vllm.rocm_base.dependency.python=${python_version_arg}" \
--label "vllm.rocm_base.dependency.pytorch=${pytorch_arg}" \
--label "vllm.rocm_base.dependency.torchvision=${pytorch_vision_arg}" \
--label "vllm.rocm_base.dependency.torchaudio=${pytorch_audio_arg}" \
--label "vllm.rocm_base.dependency.triton=${triton_arg}" \
--label "vllm.rocm_base.dependency.flash_attention=${fa_arg}" \
--label "vllm.rocm_base.dependency.aiter=${aiter_arg}" \
--label "vllm.rocm_base.dependency.mori=${mori_arg}" \
--label "vllm.rocm_base.pytorch_rocm_arch=${pytorch_rocm_arch_arg}" \
"${tags[@]}" \
--push \
.
docker buildx imagetools inspect "${descriptive_tag}" >/dev/null
metadata_set "rocm-base-refresh" "1"
metadata_set "rocm-base-image" "${descriptive_tag}"
metadata_set "rocm-base-image-descriptive" "${descriptive_tag}"
metadata_set "rocm-base-image-stable" "${stable_tag}"
metadata_set "rocm-base-image-ci-descriptive" "${ci_descriptive_tag}"
metadata_set "rocm-base-metadata-version" "${metadata_version}"
metadata_set "rocm-base-content-hash" "${base_hash}"
metadata_set "rocm-base-content-files-hash" "${content_files_hash}"
metadata_set "rocm-base-content-files" "${content_files}"
metadata_set "rocm-base-content-args" "${content_args}"
metadata_set "rocm-base-base-image-digest" "${base_image_digest}"
metadata_set "rocm-base-dockerfile" "${DOCKERFILE}"
metadata_set "rocm-base-descriptor" "${descriptor}"
metadata_set "rocm-base-dependency-summary" "${dependency_summary}"
metadata_set "rocm-base-dependency-rocm" "${rocm_version}"
metadata_set "rocm-base-dependency-python" "${python_version_arg}"
metadata_set "rocm-base-dependency-pytorch" "${pytorch_arg}"
metadata_set "rocm-base-dependency-torchvision" "${pytorch_vision_arg}"
metadata_set "rocm-base-dependency-torchaudio" "${pytorch_audio_arg}"
metadata_set "rocm-base-dependency-triton" "${triton_arg}"
metadata_set "rocm-base-dependency-flash-attention" "${fa_arg}"
metadata_set "rocm-base-dependency-aiter" "${aiter_arg}"
metadata_set "rocm-base-dependency-mori" "${mori_arg}"
metadata_set "rocm-base-pytorch-rocm-arch" "${pytorch_rocm_arch_arg}"
metadata_set "rocm-ci-image-descriptive" "${ci_descriptive_tag}"
echo "--- :white_check_mark: ROCm base image published"
echo "Use BASE_IMAGE=${descriptive_tag} for downstream ROCm CI builds"
}
main() {
metadata_set "rocm-base-refresh" "0"
if ! rocm_base_changed; then
echo "ROCm base Dockerfile did not change; skipping base image refresh"
return 0
fi
setup_builder
build_base_image
}
main "$@"
@@ -1,32 +0,0 @@
#!/usr/bin/env bash
# Fast structural smoke test for the full ROCm CI image.
set -euo pipefail
image_ref="${VLLM_CI_SMOKE_IMAGE:-rocm/vllm-ci:${BUILDKITE_COMMIT:?BUILDKITE_COMMIT is required}}"
docker run --rm --network=none --entrypoint /bin/bash "${image_ref}" -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
import vllm
print(torch.__version__)
print(vllm.__version__)
PY
echo AMD image smoke OK
'
+1 -3
View File
@@ -109,9 +109,7 @@ run_nodes() {
if [ "$node" -ne 0 ]; then
docker exec -d "node$node" /bin/bash -c "cd $WORKING_DIR ; ${COMMANDS[$node]}"
else
# Allocate a TTY (-t -i) for the foreground head node so its output
# keeps ANSI color in the Buildkite log (see run-amd-test.sh).
docker exec -t -i "node$node" /bin/bash -c "cd $WORKING_DIR ; ${COMMANDS[$node]}"
docker exec "node$node" /bin/bash -c "cd $WORKING_DIR ; ${COMMANDS[$node]}"
fi
done
}
@@ -8,12 +8,7 @@ if [[ "$MODE" != "style-clippy" && "$MODE" != "test" ]]; then
exit 2
fi
if ROOT_DIR="$(git rev-parse --show-toplevel 2>/dev/null)"; then
:
else
SCRIPT_DIR="$(cd -- "$(dirname -- "${BASH_SOURCE[0]}")" && pwd -P)"
ROOT_DIR="$(cd -- "${SCRIPT_DIR}/../.." && pwd -P)"
fi
ROOT_DIR="$(git rev-parse --show-toplevel)"
cd "$ROOT_DIR"
export CARGO_TERM_COLOR="${CARGO_TERM_COLOR:-always}"
@@ -21,21 +16,16 @@ export CARGO_HOME="${CARGO_HOME:-$HOME/.cargo}"
export RUSTUP_HOME="${RUSTUP_HOME:-$HOME/.rustup}"
export PATH="$CARGO_HOME/bin:$PATH"
PROTOC_VERSION="${PROTOC_VERSION:-31.1}"
CARGO_BINSTALL_VERSION="${CARGO_BINSTALL_VERSION:-1.20.1}"
UV_VERSION="${UV_VERSION:-0.11.28}"
PYO3_PYTHON_VERSION="${PYO3_PYTHON_VERSION:-3.12}"
CARGO_SORT_VERSION_REQ="${CARGO_SORT_VERSION_REQ:-2}"
CARGO_DENY_VERSION_REQ="${CARGO_DENY_VERSION_REQ:-0.20}"
CARGO_NEXTEST_VERSION_REQ="${CARGO_NEXTEST_VERSION_REQ:-0.9}"
CARGO_LLVM_COV_VERSION="${CARGO_LLVM_COV_VERSION:-0.8.7}"
log_section() {
echo "--- $*"
}
install_protoc() {
if command -v protoc >/dev/null 2>&1; then
return
fi
local version="${PROTOC_VERSION:-31.1}"
local arch
case "$(uname -m)" in
x86_64)
@@ -50,17 +40,16 @@ install_protoc() {
;;
esac
local url="https://github.com/protocolbuffers/protobuf/releases/download/v${PROTOC_VERSION}/protoc-${PROTOC_VERSION}-linux-${arch}.zip"
local url="https://github.com/protocolbuffers/protobuf/releases/download/v${version}/protoc-${version}-linux-${arch}.zip"
local tmp_dir
tmp_dir="$(mktemp -d)"
log_section "Installing protoc ${PROTOC_VERSION}"
log_section "Installing protoc ${version}"
curl -L --proto '=https' --tlsv1.2 -sSf "$url" -o "$tmp_dir/protoc.zip"
mkdir -p "$CARGO_HOME/bin"
unzip -q "$tmp_dir/protoc.zip" bin/protoc 'include/*' -d "$CARGO_HOME"
chmod +x "$CARGO_HOME/bin/protoc"
rm -rf "$tmp_dir"
protoc --version
}
rust_toolchain() {
@@ -81,60 +70,66 @@ install_rust_toolchain() {
}
install_cargo_binstall() {
log_section "Installing cargo-binstall ${CARGO_BINSTALL_VERSION}"
if command -v cargo-binstall >/dev/null 2>&1; then
return
fi
log_section "Installing cargo-binstall"
curl -L --proto '=https' --tlsv1.2 -sSf \
"https://raw.githubusercontent.com/cargo-bins/cargo-binstall/v${CARGO_BINSTALL_VERSION}/install-from-binstall-release.sh" \
| env BINSTALL_VERSION="$CARGO_BINSTALL_VERSION" bash
cargo-binstall -V
https://raw.githubusercontent.com/cargo-bins/cargo-binstall/main/install-from-binstall-release.sh \
| bash
}
install_cargo_sort() {
log_section "Installing cargo-sort ${CARGO_SORT_VERSION_REQ}"
cargo binstall --no-confirm --force "cargo-sort@${CARGO_SORT_VERSION_REQ}"
if command -v cargo-sort >/dev/null 2>&1; then
return
fi
log_section "Installing cargo-sort"
install_cargo_binstall
cargo binstall --no-confirm cargo-sort
}
install_cargo_deny() {
log_section "Installing cargo-deny ${CARGO_DENY_VERSION_REQ}"
cargo binstall --no-confirm --force "cargo-deny@${CARGO_DENY_VERSION_REQ}"
if command -v cargo-deny >/dev/null 2>&1; then
return
fi
log_section "Installing cargo-deny"
install_cargo_binstall
cargo binstall --no-confirm cargo-deny
}
install_cargo_nextest() {
log_section "Installing cargo-nextest ${CARGO_NEXTEST_VERSION_REQ}"
cargo binstall \
--no-confirm \
--force \
--secure \
"cargo-nextest@${CARGO_NEXTEST_VERSION_REQ}"
}
if command -v cargo-nextest >/dev/null 2>&1; then
return
fi
install_cargo_llvm_cov() {
log_section "Installing cargo-llvm-cov ${CARGO_LLVM_COV_VERSION}"
local toolchain
toolchain="$(rust_toolchain)"
rustup component add --toolchain "$toolchain" llvm-tools-preview
cargo binstall \
--no-confirm \
--force \
--secure \
"cargo-llvm-cov@${CARGO_LLVM_COV_VERSION}"
log_section "Installing cargo-nextest"
install_cargo_binstall
cargo binstall --no-confirm --secure cargo-nextest
}
install_uv() {
log_section "Installing uv ${UV_VERSION}"
curl -L --proto '=https' --tlsv1.2 -sSf \
"https://github.com/astral-sh/uv/releases/download/${UV_VERSION}/uv-installer.sh" \
if command -v uv >/dev/null 2>&1; then
return
fi
log_section "Installing uv"
curl -LsSf --proto '=https' --tlsv1.2 https://astral.sh/uv/install.sh \
| env UV_INSTALL_DIR="$CARGO_HOME/bin" sh
uv --version
}
setup_pyo3_python() {
log_section "Installing Python ${PYO3_PYTHON_VERSION} for PyO3 tests"
uv python install "$PYO3_PYTHON_VERSION"
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 \
"$PYO3_PYTHON_VERSION")"
"$python_version")"
export PYO3_PYTHON
local python_libdir
@@ -156,7 +151,6 @@ PY
}
run_style_clippy() {
install_cargo_binstall
install_cargo_sort
install_cargo_deny
@@ -169,8 +163,8 @@ run_style_clippy() {
log_section "Checking Rust dependency bans"
cargo deny \
--manifest-path rust/Cargo.toml \
--config rust/deny.toml \
check \
--config rust/deny.toml \
bans
log_section "Running clippy"
@@ -187,43 +181,15 @@ run_style_clippy() {
run_tests() {
install_uv
setup_pyo3_python
install_cargo_binstall
install_cargo_nextest
install_cargo_llvm_cov
log_section "Running cargo nextest with Rust coverage"
mkdir -p artifacts
export LLVM_PROFILE_FILE_NAME="vllm-rust-unit-%4m.profraw"
cargo llvm-cov clean \
--manifest-path rust/Cargo.toml \
--profraw-only
set +e
cargo llvm-cov nextest \
log_section "Running cargo nextest"
cargo nextest run \
--manifest-path rust/Cargo.toml \
--workspace \
--all-features \
--locked \
--no-fail-fast \
--no-clean \
--lcov \
--output-path artifacts/rust-unit.lcov \
--ignore-filename-regex='/\.cargo/(registry|git)/|/rustc/|/target/'
local coverage_rc=$?
local upload_rc=0
if [[ $coverage_rc -eq 0 ]]; then
# shellcheck source=.buildkite/scripts/rust-coverage.sh
source .buildkite/scripts/rust-coverage.sh
rust_coverage_upload artifacts/rust-unit.lcov rust-unit
upload_rc=$?
fi
set -e
if [[ $coverage_rc -ne 0 ]]; then
return "$coverage_rc"
fi
return "$upload_rc"
--no-fail-fast
}
install_protoc
-182
View File
@@ -1,182 +0,0 @@
#!/bin/sh
RUST_CODECOV_VERSION="v11.3.1"
RUST_CODECOV_SHA256="ca1d64196d2d34771084afe76ea657d581bf628e31d993ff8e52ea09cc88a56d"
rust_coverage_repo_root() {
if [ -f /vllm-workspace/.buildkite/scripts/rust-coverage.sh ]; then
printf '%s\n' /vllm-workspace
elif [ -n "${BUILDKITE_BUILD_CHECKOUT_PATH:-}" ] \
&& [ -d "$BUILDKITE_BUILD_CHECKOUT_PATH" ]; then
printf '%s\n' "$BUILDKITE_BUILD_CHECKOUT_PATH"
else
git rev-parse --show-toplevel
fi
}
rust_coverage_start() {
RUST_COVERAGE_FLAG=${1:?coverage flag is required}
RUST_COVERAGE_DIR="/tmp/vllm-rust-coverage/${BUILDKITE_JOB_ID:-local}"
export RUST_COVERAGE_FLAG RUST_COVERAGE_DIR
mkdir -p "$RUST_COVERAGE_DIR"
LLVM_PROFILE_FILE="$RUST_COVERAGE_DIR/rust-%4m.profraw"
export LLVM_PROFILE_FILE
trap rust_coverage_finalize 0
}
rust_coverage_objects() {
rust_cov_objects_manifest="$(dirname "$(command -v llvm-cov)")/../objects"
python3 - "$rust_cov_objects_manifest" <<'PY'
from pathlib import Path
import sys
for relative in Path(sys.argv[1]).read_text().splitlines():
for entry in sys.path:
path = Path(entry or ".").resolve() / relative
if path.is_file():
print(path)
break
else:
raise RuntimeError(f"installed Rust coverage object was not found: {relative}")
PY
}
rust_coverage_collect() {
rust_cov_collect_flag=${1:?coverage flag is required}
rust_cov_collect_lcov="$RUST_COVERAGE_DIR/$rust_cov_collect_flag.lcov"
rust_cov_collect_objects=$(rust_coverage_objects) || return 1
rust_cov_collect_primary=
set --
while IFS= read -r rust_cov_collect_object; do
if [ -z "$rust_cov_collect_primary" ]; then
rust_cov_collect_primary=$rust_cov_collect_object
else
set -- "$@" "--object=$rust_cov_collect_object"
fi
done <<EOF
$rust_cov_collect_objects
EOF
llvm-profdata merge \
-sparse \
"$RUST_COVERAGE_DIR"/*.profraw \
-o "$RUST_COVERAGE_DIR/merged.profdata" || return 1
llvm-cov export \
"$rust_cov_collect_primary" \
"$@" \
--format=lcov \
--instr-profile="$RUST_COVERAGE_DIR/merged.profdata" \
--ignore-filename-regex='/\.cargo/(registry|git)/|/rustc/|/target/' \
> "$rust_cov_collect_lcov" || return 1
RUST_COVERAGE_LCOV=$rust_cov_collect_lcov
export RUST_COVERAGE_LCOV
}
rust_coverage_upload() {
rust_cov_upload_lcov=${1:?LCOV path is required}
rust_cov_upload_flag=${2:?coverage flag is required}
rust_cov_upload_repo_root=$(rust_coverage_repo_root) || return 1
if [ "$(uname -m)" != "x86_64" ]; then
echo "Rust coverage upload currently supports x86_64 CI agents" >&2
return 1
fi
rust_cov_upload_codecov_dir=$(mktemp -d /tmp/codecov-bin.XXXXXX) \
|| return 1
curl -fsSL \
"https://github.com/codecov/codecov-cli/releases/download/${RUST_CODECOV_VERSION}/codecovcli_linux" \
-o "$rust_cov_upload_codecov_dir/codecov" || return 1
echo "$RUST_CODECOV_SHA256 $rust_cov_upload_codecov_dir/codecov" \
| sha256sum -c - || return 1
chmod +x "$rust_cov_upload_codecov_dir/codecov" || return 1
rust_cov_upload_slug="vllm-project/vllm"
if [ -n "${BUILDKITE_PULL_REQUEST:-}" ] \
&& [ "${BUILDKITE_PULL_REQUEST}" != "false" ] \
&& [ -n "${BUILDKITE_PULL_REQUEST_REPO:-}" ]; then
rust_cov_upload_slug=$(echo "$BUILDKITE_PULL_REQUEST_REPO" \
| sed -E 's#(git@|https?://)([^/:]+)[:/]([^/]+/[^/.]+)(\.git)?$#\3#')
case "$rust_cov_upload_slug" in
*/*) ;;
*) rust_cov_upload_slug="vllm-project/vllm" ;;
esac
fi
rust_cov_upload_branch=${BUILDKITE_BRANCH:?BUILDKITE_BRANCH is required}
if [ -z "${CODECOV_TOKEN:-}" ]; then
# Codecov accepts tokenless public uploads on unprotected branch names.
# A colon-separated prefix keeps feature-branch and fork uploads from
# requiring a repository secret.
if [ -n "${BUILDKITE_PULL_REQUEST:-}" ] \
&& [ "${BUILDKITE_PULL_REQUEST}" != "false" ]; then
rust_cov_upload_branch="pr${BUILDKITE_PULL_REQUEST}:$rust_cov_upload_branch"
else
rust_cov_upload_branch="buildkite:$rust_cov_upload_branch"
fi
fi
set --
set -- "$@" upload-process
set -- "$@" --file "$rust_cov_upload_lcov"
# LCOV paths are mapped server-side by codecov.yml. Skip the CLI's local
# source-line fix scanning, which is unrelated to path mapping.
set -- "$@" --disable-search --disable-file-fixes
set -- "$@" --fail-on-error --git-service github
set -- "$@" --build "${BUILDKITE_BUILD_NUMBER:?BUILDKITE_BUILD_NUMBER is required}"
set -- "$@" --branch "$rust_cov_upload_branch"
set -- "$@" --sha "${BUILDKITE_COMMIT:?BUILDKITE_COMMIT is required}"
set -- "$@" --slug "$rust_cov_upload_slug"
set -- "$@" --flag "$rust_cov_upload_flag"
set -- "$@" --name "${rust_cov_upload_flag}-${BUILDKITE_JOB_ID:?BUILDKITE_JOB_ID is required}"
set -- "$@" --dir "$rust_cov_upload_repo_root"
set -- "$@" --network-root-folder "$rust_cov_upload_repo_root"
if [ -n "${BUILDKITE_PULL_REQUEST:-}" ] \
&& [ "${BUILDKITE_PULL_REQUEST}" != "false" ]; then
set -- "$@" --pr "$BUILDKITE_PULL_REQUEST"
fi
rust_cov_upload_log="$rust_cov_upload_codecov_dir/codecov.log"
# E2E steps run from tests/, so execute from the repository root to resolve
# codecov.yml and repository paths consistently.
(
cd "$rust_cov_upload_repo_root" || exit 1
"$rust_cov_upload_codecov_dir/codecov" "$@"
) >"$rust_cov_upload_log" 2>&1
rust_cov_upload_rc=$?
cat "$rust_cov_upload_log"
# v11.3.1 can log API failures while returning zero even with
# --fail-on-error. Preserve the strict CI contract explicitly.
if grep -aEq 'error.* -- ' "$rust_cov_upload_log"; then
echo "Codecov CLI reported an upload error" >&2
rust_cov_upload_rc=1
fi
rm -rf "$rust_cov_upload_codecov_dir"
return "$rust_cov_upload_rc"
}
rust_coverage_finalize() {
rust_cov_finalize_test_rc=$?
trap - 0
set +e
rust_coverage_collect "$RUST_COVERAGE_FLAG"
rust_cov_finalize_collect_rc=$?
rust_cov_finalize_upload_rc=0
if [ "$rust_cov_finalize_collect_rc" -eq 0 ]; then
rust_coverage_upload "$RUST_COVERAGE_LCOV" "$RUST_COVERAGE_FLAG"
rust_cov_finalize_upload_rc=$?
fi
find "$RUST_COVERAGE_DIR" -type f -name '*.profraw' -delete
if [ "$rust_cov_finalize_test_rc" -ne 0 ]; then
exit "$rust_cov_finalize_test_rc"
fi
if [ "$rust_cov_finalize_collect_rc" -ne 0 ]; then
exit "$rust_cov_finalize_collect_rc"
fi
exit "$rust_cov_finalize_upload_rc"
}
@@ -18,10 +18,6 @@ wait_for_server() {
MODEL="Qwen/Qwen3-30B-A3B-FP8"
BACK="allgather_reducescatter"
if command -v rocm-smi &> /dev/null || [[ -d /opt/rocm ]] || [[ -n "${ROCM_PATH:-}" ]]; then
# Disable MOE padding for ROCm since it is causing eplb to fail.
export VLLM_ROCM_MOE_PADDING=0
fi
cleanup() {
if [[ -n "${SERVER_PID:-}" ]] && kill -0 "${SERVER_PID}" 2>/dev/null; then
+11 -32
View File
@@ -23,7 +23,6 @@ NC='\033[0m' # No Color
# Default configuration
PIPELINE="ci"
DRY_RUN=true
TORCH_NIGHTLY=false
usage() {
cat <<EOF
@@ -35,14 +34,12 @@ Sets RUN_ALL=1 and NIGHTLY=1 environment variables.
SAFETY: Dry-run by default. Use --execute to actually trigger a build.
Options:
--execute Actually trigger the build (default: dry-run)
--pipeline Buildkite pipeline slug (default: ${PIPELINE})
--commit Override commit SHA (default: current HEAD)
--branch Override branch name (default: current branch)
--message Custom build message (default: auto-generated)
--torch-nightly Also build and run the full suite against torch nightly
(sets TORCH_NIGHTLY=1)
--help Show this help message
--execute Actually trigger the build (default: dry-run)
--pipeline Buildkite pipeline slug (default: ${PIPELINE})
--commit Override commit SHA (default: current HEAD)
--branch Override branch name (default: current branch)
--message Custom build message (default: auto-generated)
--help Show this help message
Prerequisites:
- bk CLI installed: brew tap buildkite/buildkite && brew install buildkite/buildkite/bk
@@ -52,7 +49,6 @@ Examples:
$(basename "$0") # Dry-run, show what would happen
$(basename "$0") --execute # Actually trigger the build
$(basename "$0") --pipeline ci-shadow # Dry-run with different pipeline
$(basename "$0") --torch-nightly # Dry-run a full torch-nightly run
EOF
exit 1
}
@@ -100,10 +96,6 @@ while [[ $# -gt 0 ]]; do
MESSAGE="$2"
shift 2
;;
--torch-nightly)
TORCH_NIGHTLY=true
shift
;;
--help|-h)
usage
;;
@@ -179,17 +171,11 @@ if [[ $(echo "$REMOTE_BRANCHES" | wc -l) -gt 5 ]]; then
fi
echo ""
# Environment variables passed to the build.
BUILD_ENV=("RUN_ALL=1" "NIGHTLY=1")
if [[ "$TORCH_NIGHTLY" == true ]]; then
BUILD_ENV+=("TORCH_NIGHTLY=1")
fi
log_info "Pipeline: ${PIPELINE}"
log_info "Branch: ${BRANCH}"
log_info "Commit: ${COMMIT}"
log_info "Message: ${MESSAGE}"
log_info "Environment: ${BUILD_ENV[*]}"
log_info "Environment: RUN_ALL=1, NIGHTLY=1"
echo ""
# Build the command
@@ -201,10 +187,9 @@ CMD=(bk build create
--commit "${COMMIT}"
--branch "${BRANCH}"
--message "${MESSAGE}"
--env "RUN_ALL=1"
--env "NIGHTLY=1"
)
for env_var in "${BUILD_ENV[@]}"; do
CMD+=(--env "${env_var}")
done
if [[ "$DRY_RUN" == true ]]; then
echo "=========================================="
@@ -225,14 +210,8 @@ if [[ "$DRY_RUN" == true ]]; then
echo " --commit '$(escape_for_shell "${COMMIT}")' \\"
echo " --branch '$(escape_for_shell "${BRANCH}")' \\"
echo " --message '$(escape_for_shell "${MESSAGE}")' \\"
last_idx=$(( ${#BUILD_ENV[@]} - 1 ))
for i in "${!BUILD_ENV[@]}"; do
if [[ $i -eq $last_idx ]]; then
echo " --env '$(escape_for_shell "${BUILD_ENV[$i]}")'"
else
echo " --env '$(escape_for_shell "${BUILD_ENV[$i]}")' \\"
fi
done
echo " --env 'RUN_ALL=1' \\"
echo " --env 'NIGHTLY=1'"
echo ""
echo "=========================================="
echo -e "${YELLOW}To actually trigger this build, run:${NC}"
+4 -12
View File
@@ -6,14 +6,8 @@ set -ex
# manylinux platform tag with auditwheel.
# Index generation is handled separately by generate-and-upload-nightly-index.sh.
# auditwheel is Linux-only; macOS wheels already carry a valid tag, so skip the
# manylinux retag for them.
WHEEL_PLATFORM="${VLLM_WHEEL_PLATFORM:-linux}"
if [[ "$WHEEL_PLATFORM" == "linux" ]]; then
# shellcheck source=lib/manylinux.sh
source .buildkite/scripts/lib/manylinux.sh
fi
# shellcheck source=lib/manylinux.sh
source .buildkite/scripts/lib/manylinux.sh
BUCKET="vllm-wheels"
SUBPATH=$BUILDKITE_COMMIT
@@ -33,10 +27,8 @@ wheel="${wheel_files[0]}"
# ========= detect manylinux tag and rename ==========
if [[ "$WHEEL_PLATFORM" == "linux" ]]; then
wheel="$(apply_manylinux_tag "$wheel")"
echo "Renamed wheel to: $wheel"
fi
wheel="$(apply_manylinux_tag "$wheel")"
echo "Renamed wheel to: $wheel"
# Extract the version from the wheel
version=$(unzip -p "$wheel" '**/METADATA' | grep '^Version: ' | cut -d' ' -f2)
+3 -3
View File
@@ -113,8 +113,8 @@ $PYTHON .buildkite/scripts/generate-nightly-index.py \
echo "Uploading indices to $S3_COMMIT_PREFIX"
aws s3 cp --recursive "$INDICES_OUTPUT_DIR/" "$S3_COMMIT_PREFIX"
# Only scheduled nightly builds should update the moving nightly index.
if [[ "${NIGHTLY:-0}" == "1" ]]; then
# Update rocm/nightly/ if on main branch and not a PR
if [[ "$BUILDKITE_BRANCH" == "main" && "$BUILDKITE_PULL_REQUEST" == "false" ]] || [[ "$NIGHTLY" == "1" ]]; then
echo "Updating rocm/nightly/ index..."
aws s3 cp --recursive "$INDICES_OUTPUT_DIR/" "s3://$BUCKET/rocm/nightly/"
fi
@@ -147,7 +147,7 @@ echo ""
echo "Install command (by commit):"
echo " pip install vllm --extra-index-url https://${BUCKET}.s3.amazonaws.com/$ROCM_SUBPATH/"
echo ""
if [[ "${NIGHTLY:-0}" == "1" ]]; then
if [[ "$BUILDKITE_BRANCH" == "main" ]] || [[ "$NIGHTLY" == "1" ]]; then
echo "Install command (nightly):"
echo " pip install vllm --extra-index-url https://${BUCKET}.s3.amazonaws.com/rocm/nightly/"
fi
@@ -1,13 +0,0 @@
#!/bin/bash
set -euo pipefail
REGISTRY="public.ecr.aws/q9t5s3a7"
REPO="vllm-release-repo"
ARCH_TAG="${BUILDKITE_COMMIT}-$(uname -m)-xpu"
PLATFORM_TAG="${BUILDKITE_COMMIT}-xpu"
aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin ${REGISTRY}
docker manifest rm ${REGISTRY}/${REPO}:${PLATFORM_TAG} || true
docker manifest create ${REGISTRY}/${REPO}:${PLATFORM_TAG} ${REGISTRY}/${REPO}:${ARCH_TAG} --amend
docker manifest push ${REGISTRY}/${REPO}:${PLATFORM_TAG}
@@ -1,21 +0,0 @@
#!/bin/bash
set -ex
ORIG_TAG_NAME="$BUILDKITE_COMMIT"
REPO="vllm/vllm-openai-xpu"
echo "Pushing original XPU tag ${ORIG_TAG_NAME}-xpu to nightly tags in ${REPO}"
aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7
docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:"$ORIG_TAG_NAME"-x86_64-xpu
docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:"$ORIG_TAG_NAME"-x86_64-xpu ${REPO}:nightly-x86_64
docker push ${REPO}:nightly-x86_64
docker manifest rm ${REPO}:nightly || true
docker manifest rm ${REPO}:nightly-"$BUILDKITE_COMMIT" || true
docker manifest create ${REPO}:nightly ${REPO}:nightly-x86_64 --amend
docker manifest create ${REPO}:nightly-"$BUILDKITE_COMMIT" ${REPO}:nightly-x86_64 --amend
docker manifest push ${REPO}:nightly
docker manifest push ${REPO}:nightly-"$BUILDKITE_COMMIT"
+95 -636
View File
File diff suppressed because it is too large Load Diff
+5 -7
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: V1 attention (H100-MI300)
key: v1-attention-h100-mi300
timeout_in_minutes: 85
timeout_in_minutes: 30
device: h100
source_file_dependencies:
- vllm/config/attention.py
@@ -12,12 +12,11 @@ steps:
- vllm/v1/attention
- tests/v1/attention
commands:
- pytest -v -s v1/attention --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
parallelism: 2
- pytest -v -s v1/attention
mirror:
amd:
device: mi325_1
timeout_in_minutes: 95
timeout_in_minutes: 70
depends_on:
- image-build-amd
source_file_dependencies:
@@ -31,7 +30,7 @@ steps:
- label: V1 attention (B200)
key: v1-attention-b200
timeout_in_minutes: 80
timeout_in_minutes: 30
device: b200-k8s
source_file_dependencies:
- vllm/config/attention.py
@@ -39,5 +38,4 @@ steps:
- vllm/v1/attention
- tests/v1/attention
commands:
- pytest -v -s v1/attention --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
parallelism: 2
- pytest -v -s v1/attention
+2 -2
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Basic Correctness
key: basic-correctness
timeout_in_minutes: 45
timeout_in_minutes: 30
device: h200_18gb
source_file_dependencies:
- vllm/
@@ -19,6 +19,6 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 70
timeout_in_minutes: 50
depends_on:
- image-build-amd
+2 -3
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Benchmarks CLI Test
key: benchmarks-cli-test
timeout_in_minutes: 30
timeout_in_minutes: 20
device: h200_18gb
source_file_dependencies:
- vllm/
@@ -13,7 +13,6 @@ steps:
- pytest -v -s benchmarks/
mirror:
amd:
dind: false
device: mi300_1
depends_on:
- image-build-amd
@@ -24,7 +23,7 @@ steps:
num_gpus: 2
optional: true
working_dir: "/vllm-workspace/"
timeout_in_minutes: 20
timeout_in_minutes: 10
source_file_dependencies:
- benchmarks/attention_benchmarks/
- vllm/v1/attention/
+11 -11
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Sequence Parallel Correctness Tests (2 GPUs)
key: sequence-parallel-correctness-tests-2-gpus
timeout_in_minutes: 80
timeout_in_minutes: 50
working_dir: "/vllm-workspace/"
num_devices: 2
source_file_dependencies:
@@ -19,7 +19,7 @@ steps:
- label: Sequence Parallel Correctness Tests (2xH100)
key: sequence-parallel-correctness-tests-2xh100
timeout_in_minutes: 75
timeout_in_minutes: 50
working_dir: "/vllm-workspace/"
device: h100
optional: true
@@ -30,7 +30,7 @@ steps:
- label: AsyncTP Correctness Tests (2xH100)
key: asynctp-correctness-tests-2xh100
timeout_in_minutes: 30
timeout_in_minutes: 50
working_dir: "/vllm-workspace/"
device: h100
optional: true
@@ -41,7 +41,7 @@ steps:
- label: AsyncTP Correctness Tests (B200)
key: asynctp-correctness-tests-b200
timeout_in_minutes: 30
timeout_in_minutes: 50
working_dir: "/vllm-workspace/"
device: b200-k8s
optional: true
@@ -52,7 +52,7 @@ steps:
- label: Distributed Compile Unit Tests (2xH100)
key: distributed-compile-unit-tests-2xh100
timeout_in_minutes: 45
timeout_in_minutes: 20
working_dir: "/vllm-workspace/"
device: h100
num_devices: 2
@@ -66,7 +66,7 @@ steps:
- label: Fusion and Compile Unit Tests (2xB200)
key: fusion-and-compile-unit-tests-2xb200
timeout_in_minutes: 30
timeout_in_minutes: 20
working_dir: "/vllm-workspace/"
device: b200-k8s
source_file_dependencies:
@@ -96,7 +96,7 @@ steps:
- label: Fusion E2E Quick (H100)
key: fusion-e2e-quick-h100
timeout_in_minutes: 25
timeout_in_minutes: 15
working_dir: "/vllm-workspace/"
device: h100
num_devices: 1
@@ -115,7 +115,7 @@ steps:
- label: Fusion E2E Config Sweep (H100)
key: fusion-e2e-config-sweep-h100
timeout_in_minutes: 25
timeout_in_minutes: 30
working_dir: "/vllm-workspace/"
device: h100
num_devices: 1
@@ -149,7 +149,7 @@ steps:
- label: Fusion E2E TP2 Quick (H100)
key: fusion-e2e-tp2-quick-h100
timeout_in_minutes: 35
timeout_in_minutes: 20
working_dir: "/vllm-workspace/"
device: h100
num_devices: 2
@@ -167,7 +167,7 @@ steps:
- label: Fusion E2E TP2 AR-RMS Config Sweep (H100)
key: fusion-e2e-tp2-ar-rms-config-sweep-h100
timeout_in_minutes: 30
timeout_in_minutes: 40
working_dir: "/vllm-workspace/"
device: h100
num_devices: 2
@@ -207,7 +207,7 @@ steps:
- label: Fusion E2E TP2 (B200)
key: fusion-e2e-tp2-b200
timeout_in_minutes: 45
timeout_in_minutes: 20
working_dir: "/vllm-workspace/"
device: b200-k8s
num_devices: 2
+3 -4
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Platform Tests
key: platform-tests
timeout_in_minutes: 20
timeout_in_minutes: 15
device: h200_18gb
source_file_dependencies:
- vllm/envs.py
@@ -18,9 +18,8 @@ steps:
- pytest -v -s cuda/test_platform_no_cuda_init.py
- label: Cudagraph
device: h200_35gb
key: cudagraph
timeout_in_minutes: 30
timeout_in_minutes: 20
source_file_dependencies:
- tests/v1/cudagraph
- vllm/v1/cudagraph_dispatcher.py
@@ -29,4 +28,4 @@ steps:
commands:
- pytest -v -s v1/cudagraph/test_cudagraph_dispatch.py
- pytest -v -s v1/cudagraph/test_cudagraph_mode.py
- pytest -v -s v1/cudagraph/test_breakable_cudagraph.py
- pytest -v -s v1/cudagraph/test_breakable_cudagraph.py
+13 -53
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Distributed NixlConnector PD accuracy (4 GPUs)
key: distributed-nixlconnector-pd-accuracy-4-gpus
timeout_in_minutes: 55
timeout_in_minutes: 30
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
@@ -15,9 +15,8 @@ steps:
- bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
mirror:
amd:
dind: false
device: mi300_4
timeout_in_minutes: 85
timeout_in_minutes: 110
depends_on:
- image-build-amd
source_file_dependencies:
@@ -30,7 +29,7 @@ steps:
- label: Distributed FlashInfer NixlConnector PD accuracy (4 GPUs)
key: distributed-flashinfer-nixlconnector-pd-accuracy-4-gpus
timeout_in_minutes: 55
timeout_in_minutes: 30
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
@@ -40,19 +39,6 @@ steps:
- bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
- FLASHINFER=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
- label: Push NixlConnector PP prefill PD accuracy (4 GPUs)
key: push-nixlconnector-pp-prefill-pd-accuracy-4-gpus
timeout_in_minutes: 30
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
- vllm/distributed/kv_transfer/kv_connector/v1/nixl/
- tests/v1/kv_connector/nixl_integration/
- tests/v1/kv_connector/nixl_push_integration/
commands:
- bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
- bash v1/kv_connector/nixl_push_integration/config_sweep_accuracy_test.sh
- label: DP EP Distributed NixlConnector PD accuracy tests (4 GPUs)
key: dp-ep-distributed-nixlconnector-pd-accuracy-tests-4-gpus
timeout_in_minutes: 30
@@ -66,9 +52,8 @@ steps:
- DP_EP=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
mirror:
amd:
dind: false
device: mi300_4
timeout_in_minutes: 60
timeout_in_minutes: 50
depends_on:
- image-build-amd
source_file_dependencies:
@@ -81,7 +66,7 @@ steps:
- label: CrossLayer KV layout Distributed NixlConnector PD accuracy tests (4 GPUs)
key: crosslayer-kv-layout-distributed-nixlconnector-pd-accuracy-tests-4-gpus
timeout_in_minutes: 55
timeout_in_minutes: 30
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
@@ -92,9 +77,8 @@ steps:
- CROSS_LAYERS_BLOCKS=True bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
mirror:
amd:
dind: false
device: mi300_4
timeout_in_minutes: 85
timeout_in_minutes: 110
depends_on:
- image-build-amd
source_file_dependencies:
@@ -107,7 +91,7 @@ steps:
- label: Hybrid SSM NixlConnector PD accuracy tests (4 GPUs)
key: hybrid-ssm-nixlconnector-pd-accuracy-tests-4-gpus
timeout_in_minutes: 60
timeout_in_minutes: 25
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
@@ -118,9 +102,8 @@ steps:
- HYBRID_SSM=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
mirror:
amd:
dind: false
device: mi300_4
timeout_in_minutes: 80
timeout_in_minutes: 60
depends_on:
- image-build-amd
source_file_dependencies:
@@ -147,7 +130,7 @@ steps:
- label: MultiConnector (Nixl+Offloading) PD accuracy (2 GPUs)
key: multiconnector-nixl-offloading-pd-accuracy-2-gpus
timeout_in_minutes: 40
timeout_in_minutes: 30
working_dir: "/vllm-workspace/tests"
num_devices: 2
source_file_dependencies:
@@ -162,7 +145,7 @@ steps:
- label: NixlConnector PD + Spec Decode acceptance (2 GPUs)
key: nixlconnector-pd-spec-decode-acceptance-2-gpus
timeout_in_minutes: 45
timeout_in_minutes: 30
device: a100
working_dir: "/vllm-workspace/tests"
num_devices: 2
@@ -175,9 +158,8 @@ steps:
- bash v1/kv_connector/nixl_integration/config_sweep_spec_decode_test.sh
mirror:
amd:
dind: false
device: mi300_2
timeout_in_minutes: 70
timeout_in_minutes: 60
depends_on:
- image-build-amd
source_file_dependencies:
@@ -187,11 +169,11 @@ steps:
- vllm/platforms/rocm.py
commands:
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
- KV_CACHE_MEMORY_BYTES=8G ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/config_sweep_spec_decode_test.sh
- ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/config_sweep_spec_decode_test.sh
- label: MultiConnector (Nixl+Offloading) PD edge cases (2 GPUs)
key: multiconnector-nixl-offloading-pd-edge-cases-2-gpus
timeout_in_minutes: 25
timeout_in_minutes: 30
working_dir: "/vllm-workspace/tests"
num_devices: 2
source_file_dependencies:
@@ -203,25 +185,3 @@ steps:
commands:
- bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
- bash v1/kv_connector/nixl_integration/run_multi_connector_edge_case_test.sh
# P TP 4 - D DPEP 4 test case for DSv4-Flash
- label: DSv4-Flash Disaggregated DP EP
key: dsv4-flash-disaggregated
timeout_in_minutes: 60
device: h200
optional: true
working_dir: "/vllm-workspace/tests"
num_devices: 8
env:
ENABLE_HMA_FLAG: "1"
DP_EP: "1"
GPU_MEMORY_UTILIZATION: "0.85"
PREFILLER_TP_SIZE: "4"
DECODER_TP_SIZE: "4"
PREFILL_BLOCK_SIZE: "256"
DECODE_BLOCK_SIZE: "256"
MODEL_NAMES: "deepseek-ai/DeepSeek-V4-Flash"
VLLM_SERVE_EXTRA_ARGS: "--trust-remote-code,--kv-cache-dtype,fp8"
commands:
- bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
- bash v1/kv_connector/nixl_integration/run_accuracy_test.sh
+10 -11
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Distributed Comm Ops
key: distributed-comm-ops
timeout_in_minutes: 25
timeout_in_minutes: 20
working_dir: "/vllm-workspace/tests"
num_devices: 2
source_file_dependencies:
@@ -18,7 +18,7 @@ steps:
- label: Distributed DP Tests (2 GPUs)
key: distributed-dp-tests-2-gpus
timeout_in_minutes: 35
timeout_in_minutes: 20
working_dir: "/vllm-workspace/tests"
num_devices: 2
source_file_dependencies:
@@ -39,7 +39,6 @@ steps:
- DP_SIZE=2 pytest -v -s entrypoints/openai/test_multi_api_servers.py
mirror:
amd:
dind: false
device: mi300_2
depends_on:
- image-build-amd
@@ -56,7 +55,7 @@ steps:
- label: Distributed Compile + RPC Tests (2 GPUs)
key: distributed-compile-rpc-tests-2-gpus
timeout_in_minutes: 65
timeout_in_minutes: 20
working_dir: "/vllm-workspace/tests"
num_devices: 2
source_file_dependencies:
@@ -79,7 +78,7 @@ steps:
- label: Distributed Torchrun + Shutdown Tests (2 GPUs)
key: distributed-torchrun-shutdown-tests-2-gpus
timeout_in_minutes: 30
timeout_in_minutes: 20
working_dir: "/vllm-workspace/tests"
num_devices: 2
source_file_dependencies:
@@ -134,7 +133,7 @@ steps:
- label: Distributed DP Tests (4 GPUs)
key: distributed-dp-tests-4-gpus
timeout_in_minutes: 45
timeout_in_minutes: 30
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
@@ -155,7 +154,7 @@ steps:
- label: Distributed Compile + Comm (4 GPUs)
key: distributed-compile-comm-4-gpus
timeout_in_minutes: 70
timeout_in_minutes: 30
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
@@ -177,7 +176,7 @@ steps:
- label: Distributed Tests (8xH100)
key: distributed-tests-8xh100
timeout_in_minutes: 20
timeout_in_minutes: 10
device: h100
num_devices: 8
working_dir: "/vllm-workspace/tests"
@@ -213,7 +212,7 @@ steps:
- label: Distributed Tests (2xH100-2xMI300)
key: distributed-tests-2xh100-2xmi300
timeout_in_minutes: 30
timeout_in_minutes: 15
device: h100
optional: true
working_dir: "/vllm-workspace/"
@@ -260,7 +259,7 @@ steps:
- label: Pipeline + Context Parallelism (4 GPUs)
key: pipeline-context-parallelism-4-gpus
timeout_in_minutes: 55
timeout_in_minutes: 60
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
@@ -275,7 +274,7 @@ steps:
- label: RayExecutorV2 (4 GPUs)
key: rayexecutorv2-4-gpus
timeout_in_minutes: 45
timeout_in_minutes: 60
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
+1 -1
View File
@@ -3,7 +3,7 @@ depends_on:
- image-build-cpu
steps:
- label: Docker Build Metadata
timeout_in_minutes: 20
timeout_in_minutes: 10
device: cpu-small
source_file_dependencies:
- .buildkite/release-pipeline.yaml
+5 -5
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: DeepSeek V2-Lite Sync EPLB Accuracy (4xH100)
key: deepseek-v2-lite-sync-eplb-accuracy-4xh100
timeout_in_minutes: 25
timeout_in_minutes: 60
device: h100
optional: true
num_devices: 4
@@ -14,7 +14,7 @@ steps:
- label: Qwen3-30B-A3B-FP8-block Sync EPLB Accuracy (4xH100)
key: qwen3-30b-a3b-fp8-block-sync-eplb-accuracy-4xh100
timeout_in_minutes: 25
timeout_in_minutes: 60
device: h100
optional: true
num_devices: 4
@@ -24,7 +24,7 @@ steps:
- label: Qwen3-30B-A3B-FP8-block Sync EPLB Accuracy (2xB200)
key: qwen3-30b-a3b-fp8-block-sync-eplb-accuracy-2xb200
timeout_in_minutes: 20
timeout_in_minutes: 60
device: b200-k8s
optional: true
num_devices: 2
@@ -34,7 +34,7 @@ steps:
- label: Qwen3-30B-A3B-FP8 DP4 Async EPLB Accuracy
key: qwen3-30b-a3b-fp8-dp4-async-eplb-accuracy
timeout_in_minutes: 25
timeout_in_minutes: 60
device: h100
optional: true
num_devices: 4
@@ -44,7 +44,7 @@ steps:
- label: DeepSeek V2-Lite Prefetch Offload Accuracy (H100)
key: deepseek-v2-lite-prefetch-offload-accuracy-h100
timeout_in_minutes: 20
timeout_in_minutes: 60
device: h100
optional: true
num_devices: 1
+13 -14
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Engine
key: engine
timeout_in_minutes: 30
timeout_in_minutes: 15
device: h200_18gb
source_file_dependencies:
- vllm/compilation/
@@ -29,13 +29,13 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 50
timeout_in_minutes: 60
depends_on:
- image-build-amd
- label: Engine (1 GPU)
key: engine-1-gpu
timeout_in_minutes: 45
timeout_in_minutes: 30
source_file_dependencies:
- vllm/v1/engine/
- tests/v1/engine/
@@ -45,13 +45,13 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 55
timeout_in_minutes: 40
depends_on:
- image-build-amd
- label: e2e Scheduling (1 GPU)
key: e2e-scheduling-1-gpu
timeout_in_minutes: 35
timeout_in_minutes: 30
device: h200_18gb
source_file_dependencies:
- vllm/v1/
@@ -60,15 +60,15 @@ steps:
- pytest -v -s v1/e2e/general/test_async_scheduling.py
mirror:
amd:
device: mi325_1
timeout_in_minutes: 70
device: mi250_1
timeout_in_minutes: 60
depends_on:
- image-build-amd
- label: e2e Core (1 GPU)
device: h200_35gb
key: e2e-core-1-gpu
timeout_in_minutes: 40
timeout_in_minutes: 30
source_file_dependencies:
- vllm/v1/
- tests/v1/e2e/general/
@@ -76,8 +76,8 @@ steps:
- pytest -v -s v1/e2e/general --ignore v1/e2e/general/test_async_scheduling.py
mirror:
amd:
device: mi325_1
timeout_in_minutes: 60
device: mi250_1
timeout_in_minutes: 35
depends_on:
- image-build-amd
source_file_dependencies:
@@ -87,7 +87,7 @@ steps:
- label: V1 e2e (2 GPUs)
key: v1-e2e-2-gpus
timeout_in_minutes: 25 # TODO: Fix timeout after we have more confidence in the test stability
timeout_in_minutes: 60 # TODO: Fix timeout after we have more confidence in the test stability
optional: true
num_devices: 2
source_file_dependencies:
@@ -114,14 +114,13 @@ steps:
- pytest -v -s v1/e2e/spec_decode/test_spec_decode.py -k "tensor_parallelism"
mirror:
amd:
dind: false
device: mi300_2
depends_on:
- image-build-amd
- label: V1 e2e (4 GPUs)
key: v1-e2e-4-gpus
timeout_in_minutes: 20 # TODO: Fix timeout after we have more confidence in the test stability
timeout_in_minutes: 60 # TODO: Fix timeout after we have more confidence in the test stability
optional: true
num_devices: 4
source_file_dependencies:
@@ -149,7 +148,7 @@ steps:
- label: V1 e2e (4xH100)
key: v1-e2e-4xh100
timeout_in_minutes: 35
timeout_in_minutes: 60
device: h100
num_devices: 4
optional: true
+10 -17
View File
@@ -3,9 +3,8 @@ depends_on:
- image-build
steps:
- label: Entrypoints Unit Tests
device: h200_35gb
key: entrypoints-unit-tests
timeout_in_minutes: 25
timeout_in_minutes: 10
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/entrypoints
@@ -16,9 +15,8 @@ steps:
- pytest -v -s entrypoints/weight_transfer
- label: Entrypoints Integration (LLM)
device: h200_35gb
key: entrypoints-integration-llm
timeout_in_minutes: 60
timeout_in_minutes: 40
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
@@ -39,7 +37,7 @@ steps:
- label: Entrypoints Integration (API Server)
key: entrypoints-integration-api-server
device: h200_35gb
timeout_in_minutes: 50
timeout_in_minutes: 130
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
@@ -57,9 +55,8 @@ steps:
- image-build-amd
- label: Entrypoints Integration (API Server OpenAI - Part 1)
device: h200_35gb
key: entrypoints-integration-api-server-openai-part-1
timeout_in_minutes: 45
timeout_in_minutes: 50
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
@@ -71,14 +68,13 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 65
timeout_in_minutes: 80
depends_on:
- image-build-amd
- label: Entrypoints Integration (API Server OpenAI - Part 2)
device: h200_35gb
key: entrypoints-integration-api-server-openai-part-2
timeout_in_minutes: 45
timeout_in_minutes: 50
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
@@ -96,7 +92,6 @@ steps:
- image-build-amd
- label: Entrypoints Integration (API Server Generate)
device: h200_35gb
key: entrypoints-integration-api-server-generate
timeout_in_minutes: 50
working_dir: "/vllm-workspace/tests"
@@ -114,12 +109,11 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 65
timeout_in_minutes: 60
depends_on:
- image-build-amd
- label: Entrypoints Integration (Responses API)
device: h200_35gb
key: entrypoints-integration-responses-api
timeout_in_minutes: 50
working_dir: "/vllm-workspace/tests"
@@ -132,7 +126,7 @@ steps:
- label: Entrypoints Integration (Speech to Text)
device: h200_35gb
key: entrypoints-integration-speech_to_text
timeout_in_minutes: 45
timeout_in_minutes: 50
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
@@ -144,7 +138,7 @@ steps:
- label: Entrypoints Integration (Multimodal)
device: h200_35gb
key: entrypoints-integration-multimodal
timeout_in_minutes: 45
timeout_in_minutes: 50
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
@@ -154,7 +148,6 @@ steps:
- pytest -v -s entrypoints/multimodal
- label: Entrypoints Integration (Pooling)
device: h200_35gb
key: entrypoints-integration-pooling
timeout_in_minutes: 50
working_dir: "/vllm-workspace/tests"
@@ -167,7 +160,7 @@ steps:
- label: OpenAI API Correctness
key: openai-api-correctness
timeout_in_minutes: 20
timeout_in_minutes: 30
device: h200_18gb
source_file_dependencies:
- csrc/
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: EPLB Algorithm
key: eplb-algorithm
timeout_in_minutes: 20
timeout_in_minutes: 15
device: h200_18gb
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
@@ -16,7 +16,6 @@ steps:
- pytest -v -s distributed/test_eplb_utils.py
mirror:
amd:
dind: false
device: mi300_1
depends_on:
- image-build-amd
@@ -28,7 +27,7 @@ steps:
- label: EPLB Execution # 17min
key: eplb-execution
timeout_in_minutes: 25
timeout_in_minutes: 27
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
@@ -40,7 +39,7 @@ steps:
- label: Elastic EP Scaling Test
key: elastic-ep-scaling-test
timeout_in_minutes: 30
timeout_in_minutes: 20
device: h100
working_dir: "/vllm-workspace/tests"
num_devices: 4
+24 -40
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: vLLM IR Tests
key: vllm-ir-tests
timeout_in_minutes: 35
timeout_in_minutes: 10
device: h200_18gb
working_dir: "/vllm-workspace/"
source_file_dependencies:
@@ -15,21 +15,19 @@ steps:
- pytest -v -s tests/kernels/ir
- label: Kernels Core Operation Test
device: h200_35gb
key: kernels-core-operation-test
timeout_in_minutes: 120
timeout_in_minutes: 75
source_file_dependencies:
- 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 --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
parallelism: 3
- 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
- label: Kernels MiniMax Reduce RMS Test (2 GPUs)
key: kernels-minimax-reduce-rms-test-2-gpus
timeout_in_minutes: 20
timeout_in_minutes: 15
num_devices: 2
device: h100
source_file_dependencies:
@@ -43,7 +41,7 @@ steps:
- label: Deepseek V4 Kernel Test (H100)
key: deepseek-v4-kernel-test-h100
timeout_in_minutes: 30
timeout_in_minutes: 15
device: h100
source_file_dependencies:
- csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu
@@ -56,7 +54,7 @@ steps:
- label: Deepseek V4 Kernel Test (B200)
key: deepseek-v4-kernel-test-b200
timeout_in_minutes: 20
timeout_in_minutes: 15
device: b200-k8s
source_file_dependencies:
- csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu
@@ -67,7 +65,7 @@ steps:
- label: Kernels Attention Test %N
key: kernels-attention-test
timeout_in_minutes: 65
timeout_in_minutes: 35
source_file_dependencies:
- csrc/attention/
- vllm/v1/attention
@@ -81,7 +79,7 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 90
timeout_in_minutes: 55
depends_on:
- image-build-amd
source_file_dependencies:
@@ -108,7 +106,7 @@ steps:
- label: Kernels Quantization Test %N
key: kernels-quantization-test
timeout_in_minutes: 60
timeout_in_minutes: 90
source_file_dependencies:
- csrc/quantization/
- vllm/model_executor/layers/quantization
@@ -133,7 +131,7 @@ steps:
- label: Kernels MoE Test %N
key: kernels-moe-test
timeout_in_minutes: 50
timeout_in_minutes: 25
source_file_dependencies:
- csrc/quantization/cutlass_w8a8/moe/
- csrc/moe/
@@ -149,7 +147,7 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 65
timeout_in_minutes: 50
source_file_dependencies:
- csrc/quantization/cutlass_w8a8/moe/
- csrc/moe/
@@ -164,9 +162,8 @@ steps:
- image-build-amd
- label: Kernels Mamba Test
device: h200_35gb
key: kernels-mamba-test
timeout_in_minutes: 40
timeout_in_minutes: 45
source_file_dependencies:
- csrc/mamba/
- tests/kernels/mamba
@@ -175,19 +172,19 @@ steps:
- pytest -v -s kernels/mamba
- label: Kernels KDA Test
timeout_in_minutes: 25
timeout_in_minutes: 20
device: h200_18gb
source_file_dependencies:
- vllm/third_party/flash_linear_attention/ops/kda.py
- vllm/third_party/flash_linear_attention/ops/chunk_delta_h.py
- vllm/third_party/flash_linear_attention/ops/l2norm.py
- vllm/model_executor/layers/fla/ops/kda.py
- vllm/model_executor/layers/fla/ops/chunk_delta_h.py
- vllm/model_executor/layers/fla/ops/l2norm.py
- tests/kernels/test_kda.py
commands:
- pytest -v -s kernels/test_kda.py
- label: Kernels DeepGEMM Test (H100)
key: kernels-deepgemm-test-h100
timeout_in_minutes: 35
timeout_in_minutes: 45
device: h100
num_devices: 1
source_file_dependencies:
@@ -214,7 +211,7 @@ steps:
- label: Kernels (B200)
key: kernels-b200
timeout_in_minutes: 80
timeout_in_minutes: 30
working_dir: "/vllm-workspace/"
device: b200-k8s
# optional: true
@@ -234,15 +231,6 @@ steps:
- vllm/v1/attention/backends/mla/flashinfer_mla.py
- vllm/v1/attention/selector.py
- vllm/platforms/cuda.py
- vllm/model_executor/kernels/linear/cute_dsl/ll_bf16.py
- vllm/model_executor/kernels/linear/cute_dsl/_ll_bf16_dotprod.py
- vllm/model_executor/kernels/linear/cute_dsl/_ll_bf16_splitk.py
- vllm/cute_utils/
- vllm/model_executor/layers/mamba/ops/gdn_chunk_cutedsl/
- vllm/model_executor/layers/fused_moe/router/bf16x3_router_gemm_cutedsl.py
- tests/kernels/mamba/test_gdn_prefill_cutedsl.py
- tests/kernels/test_bf16x3_router_gemm_cutedsl.py
- tests/kernels/test_ll_bf16_gemm.py
- tests/kernels/test_top_k_per_row.py
commands:
- nvidia-smi
@@ -271,28 +259,24 @@ steps:
- pytest -v -s tests/kernels/moe/test_flashinfer_moe.py
- pytest -v -s tests/kernels/moe/test_trtllm_nvfp4_moe.py
- pytest -v -s tests/kernels/moe/test_cutedsl_moe.py
- pytest -v -s tests/kernels/mamba/test_gdn_prefill_cutedsl.py
- pytest -v -s tests/kernels/test_bf16x3_router_gemm_cutedsl.py
- pytest -v -s tests/kernels/test_ll_bf16_gemm.py
# e2e
- pytest -v -s tests/models/quantization/test_nvfp4.py
- label: Kernels Helion Test
key: kernels-helion-test
timeout_in_minutes: 115
timeout_in_minutes: 30
device: h100
source_file_dependencies:
- vllm/utils/import_utils.py
- tests/kernels/helion/
commands:
- pip install helion==1.1.0
- pytest -v -s kernels/helion/ --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
parallelism: 2
- pytest -v -s kernels/helion/
- label: Kernels FP8 MoE Test (1xH100)
key: kernels-fp8-moe-test-1xh100
timeout_in_minutes: 40
timeout_in_minutes: 90
device: h100
num_devices: 1
optional: true
@@ -309,7 +293,7 @@ steps:
- label: Kernels FP8 MoE Test (2xH100)
key: kernels-fp8-moe-test-2xh100
timeout_in_minutes: 45
timeout_in_minutes: 90
device: h100
num_devices: 2
optional: true
@@ -319,7 +303,7 @@ steps:
- label: Kernels Fp4 MoE Test (B200)
key: kernels-fp4-moe-test-b200
timeout_in_minutes: 25
timeout_in_minutes: 60
device: b200-k8s
num_devices: 1
optional: true
@@ -332,7 +316,7 @@ steps:
- label: Kernels FusedMoE Layer Test (2 H100s)
key: kernels-fusedmoe-layer-test-2-h100s
timeout_in_minutes: 30
timeout_in_minutes: 90
device: h100
num_devices: 2
source_file_dependencies:
+24 -136
View File
@@ -5,7 +5,7 @@ steps:
- label: LM Eval Small Models
device: h200_35gb
key: lm-eval-small-models
timeout_in_minutes: 45
timeout_in_minutes: 75
source_file_dependencies:
- csrc/
- vllm/model_executor/layers/quantization
@@ -56,7 +56,7 @@ steps:
- label: LM Eval Small Models (1xB200)
key: lm-eval-small-models-1xb200
timeout_in_minutes: 50
timeout_in_minutes: 120
device: b200-k8s
optional: true
source_file_dependencies:
@@ -78,31 +78,9 @@ steps:
commands:
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-small-tp.txt
- label: LM Eval PCP (4xB200)
key: lm-eval-pcp-4xb200
timeout_in_minutes: 360
device: b200-k8s
num_devices: 4
optional: true
source_file_dependencies:
- csrc/
- tests/evals/gsm8k/configs/GLM-5.2-NVFP4-TP2-PCP2-EP.yaml
- tests/evals/gsm8k/configs/GLM-5.2-NVFP4-TP1-PCP4-EP.yaml
- tests/evals/gsm8k/configs/models-pcp.txt
- vllm/model_executor/layers/quantization
- vllm/config/parallel.py
- vllm/distributed/parallel_state.py
- vllm/model_executor/layers/attention/mla_attention.py
- vllm/model_executor/layers/attention/pcp.py
- vllm/v1/worker/gpu/model_runner.py
- vllm/v1/worker/gpu/pcp_manager.py
autorun_on_main: true
commands:
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-pcp.txt
- label: LM Eval Large Models EP (2xB200)
key: lm-eval-large-models-ep-2xb200
timeout_in_minutes: 60
timeout_in_minutes: 120
device: b200-k8s
optional: true
num_devices: 2
@@ -114,7 +92,7 @@ steps:
- label: LM Eval Qwen3.5 Models (2xB200)
key: lm-eval-qwen3-5-models-2xb200
timeout_in_minutes: 45
timeout_in_minutes: 120
device: b200-k8s
optional: true
num_devices: 2
@@ -125,13 +103,13 @@ steps:
- vllm/transformers_utils/configs/qwen3_5_moe.py
- vllm/model_executor/models/qwen3_next.py
- vllm/model_executor/models/qwen3_next_mtp.py
- vllm/third_party/flash_linear_attention/ops/
- vllm/model_executor/layers/fla/ops/
commands:
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-qwen35-blackwell.txt
- label: LM Eval Large Models (8xH200)
key: lm-eval-large-models-8xh200
timeout_in_minutes: 50
timeout_in_minutes: 60
device: h200
optional: true
num_devices: 8
@@ -139,9 +117,8 @@ steps:
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-h200.txt
mirror:
amd:
dind: false
device: mi300_8
timeout_in_minutes: 60
timeout_in_minutes: 180
depends_on:
- image-build-amd
commands:
@@ -173,9 +150,9 @@ 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 f16 (A100 - TEMPORARY)
key: lm-eval-humming-f16-a100
timeout_in_minutes: 75
- label: LM Eval Humming (A100 - TEMPORARY)
key: lm-eval-humming-a100
timeout_in_minutes: 30
device: a100
optional: true
num_devices: 1
@@ -183,29 +160,13 @@ steps:
- 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/
- vllm/model_executor/kernels/linear/
- 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 Act int8 (A100 - TEMPORARY)
key: lm-eval-humming-act-a100
timeout_in_minutes: 45
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/
- vllm/model_executor/kernels/linear/
commands:
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config-act-int8.txt
- label: LM Eval Humming f16 (H100 - TEMPORARY)
key: lm-eval-humming-f16-h100
timeout_in_minutes: 70
- label: LM Eval Humming (H100 - TEMPORARY)
key: lm-eval-humming-h100
timeout_in_minutes: 30
device: h100
optional: true
num_devices: 1
@@ -213,30 +174,14 @@ steps:
- 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/
- vllm/model_executor/kernels/linear/
- 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 Act fp8/int8 (H100 - TEMPORARY)
key: lm-eval-humming-act-h100
timeout_in_minutes: 70
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/
- vllm/model_executor/kernels/linear/
commands:
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config-act-fp8.txt
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config-act-int8.txt
- label: LM Eval Humming f16 (B200 - TEMPORARY)
key: lm-eval-humming-f16-b200
timeout_in_minutes: 50
- label: LM Eval Humming (B200 - TEMPORARY)
key: lm-eval-humming-b200
timeout_in_minutes: 30
device: b200-k8s
optional: true
num_devices: 1
@@ -244,30 +189,14 @@ steps:
- 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/
- vllm/model_executor/kernels/linear/
- 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 Act fp8/int8 (B200 - TEMPORARY)
key: lm-eval-humming-act-b200
timeout_in_minutes: 50
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/
- vllm/model_executor/kernels/linear/
commands:
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config-act-fp8.txt
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config-act-int8.txt
- label: LM Eval TurboQuant KV Cache
key: lm-eval-turboquant-kv-cache
timeout_in_minutes: 55
timeout_in_minutes: 75
device: h200_18gb
source_file_dependencies:
- vllm/model_executor/layers/quantization/turboquant/
@@ -279,7 +208,7 @@ steps:
- label: GPQA Eval (GPT-OSS) (2xH100)
key: gpqa-eval-gpt-oss-2xh100
timeout_in_minutes: 35
timeout_in_minutes: 120
device: h100
optional: true
num_devices: 2
@@ -293,7 +222,7 @@ steps:
- label: GPQA Eval (GPT-OSS) (2xB200)
key: gpqa-eval-gpt-oss-2xb200
timeout_in_minutes: 30
timeout_in_minutes: 120
device: b200-k8s
optional: true
num_devices: 2
@@ -307,7 +236,7 @@ steps:
- label: GPQA Eval (GPT-OSS) (DGX Spark)
key: gpqa-eval-gpt-oss-spark
timeout_in_minutes: 35
timeout_in_minutes: 120
device: dgx-spark
optional: true
num_devices: 1
@@ -321,50 +250,9 @@ 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-spark.txt
- label: LM Eval KV-Offload (1xH200)
key: kv-offload-small
timeout_in_minutes: 30
device: h200_35gb
source_file_dependencies:
- vllm/distributed/kv_transfer/kv_connector/v1/offloading/
- vllm/distributed/kv_transfer/kv_connector/v1/simple_cpu_offload_connector.py
- vllm/v1/kv_offload/
- vllm/v1/simple_kv_offload/
- tests/evals/gsm8k/test_gsm8k_offloading.py
commands:
- pytest -s -v evals/gsm8k/test_gsm8k_offloading.py -k "nemotron-h-8b or gemma-4-e4b-it"
- label: LM Eval KV-Offload (2xH100)
key: kv-offload-medium
timeout_in_minutes: 30
device: h100
num_devices: 2
source_file_dependencies:
- vllm/distributed/kv_transfer/kv_connector/v1/offloading/
- vllm/distributed/kv_transfer/kv_connector/v1/simple_cpu_offload_connector.py
- vllm/v1/kv_offload/
- vllm/v1/simple_kv_offload/
- tests/evals/gsm8k/test_gsm8k_offloading.py
commands:
- pytest -s -v evals/gsm8k/test_gsm8k_offloading.py -k "qwen3.5-35b"
- label: LM Eval KV-Offload (4xH100)
key: kv-offload-large
timeout_in_minutes: 40
device: h100
num_devices: 4
source_file_dependencies:
- vllm/distributed/kv_transfer/kv_connector/v1/offloading/
- vllm/distributed/kv_transfer/kv_connector/v1/simple_cpu_offload_connector.py
- vllm/v1/kv_offload/
- vllm/v1/simple_kv_offload/
- tests/evals/gsm8k/test_gsm8k_offloading.py
commands:
- pytest -s -v evals/gsm8k/test_gsm8k_offloading.py -k "deepseek-v4-flash"
- label: MRCR Eval Small Models
device: h200_35gb
timeout_in_minutes: 25
timeout_in_minutes: 30
source_file_dependencies:
- tests/evals/mrcr/
commands:
+3 -3
View File
@@ -5,7 +5,7 @@ steps:
- label: LoRA %N
device: h200_35gb
key: lora
timeout_in_minutes: 40
timeout_in_minutes: 30
source_file_dependencies:
- vllm/lora
- tests/lora
@@ -16,7 +16,7 @@ steps:
amd:
device: mi325_1
working_dir: "/vllm-workspace/tests"
timeout_in_minutes: 65
timeout_in_minutes: 60
source_file_dependencies:
- vllm/lora
- tests/lora
@@ -27,7 +27,7 @@ steps:
- label: LoRA TP (Distributed)
key: lora-tp-distributed
timeout_in_minutes: 60
timeout_in_minutes: 30
num_devices: 4
source_file_dependencies:
- vllm/lora
+17 -23
View File
@@ -5,7 +5,7 @@ steps:
- label: V1 Spec Decode
device: h200_35gb
key: v1-spec-decode
timeout_in_minutes: 40
timeout_in_minutes: 30
source_file_dependencies:
- vllm/config/
- vllm/distributed/
@@ -23,15 +23,14 @@ steps:
- pytest -v -s -m 'not slow_test' v1/spec_decode
mirror:
amd:
dind: false
device: mi300_1
timeout_in_minutes: 75
timeout_in_minutes: 65
depends_on:
- image-build-amd
- label: V1 Sample + Logits
key: v1-sample-logits
timeout_in_minutes: 45
timeout_in_minutes: 30
device: h200_18gb
source_file_dependencies:
- vllm/config/
@@ -64,9 +63,8 @@ steps:
- image-build-amd
- label: V1 Core + KV + Metrics
device: h200_35gb
key: v1-core-kv-metrics
timeout_in_minutes: 80
timeout_in_minutes: 30
source_file_dependencies:
- vllm/config/
- vllm/distributed/
@@ -91,7 +89,6 @@ steps:
- tests/v1/simple_kv_offload
- tests/v1/worker
- tests/v1/kv_connector/unit
- tests/v1/ec_connector/unit
- tests/v1/metrics
- tests/entrypoints/openai/correctness/test_lmeval.py
commands:
@@ -104,15 +101,14 @@ steps:
- pytest -v -s v1/simple_kv_offload
- pytest -v -s v1/worker
- pytest -v -s -m 'not cpu_test' v1/kv_connector/unit
- pytest -v -s -m 'not cpu_test' v1/ec_connector/unit
- pytest -v -s -m 'not cpu_test' v1/metrics
# Integration test for streaming correctness (requires special branch).
- pip install -U git+https://github.com/vllm-project/lm-evaluation-harness.git@streaming-api
- 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: 75
timeout_in_minutes: 60
depends_on:
- image-build-amd
@@ -176,7 +172,7 @@ steps:
- label: Regression
key: regression
timeout_in_minutes: 30
timeout_in_minutes: 20
device: h200_18gb
source_file_dependencies:
- vllm/config/
@@ -192,14 +188,14 @@ steps:
- vllm/v1/
- tests/test_regression
commands:
- pip install 'modelscope<1.38'
- pip install modelscope
- pytest -v -s test_regression.py
working_dir: "/vllm-workspace/tests" # optional
- label: Examples
device: h200_35gb
key: examples
timeout_in_minutes: 40
timeout_in_minutes: 45
working_dir: "/vllm-workspace/examples"
source_file_dependencies:
- vllm/entrypoints
@@ -241,7 +237,7 @@ steps:
- label: Metrics, Tracing (2 GPUs)
key: metrics-tracing-2-gpus
timeout_in_minutes: 25
timeout_in_minutes: 20
num_devices: 2
source_file_dependencies:
- vllm/config/
@@ -285,7 +281,7 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 45
timeout_in_minutes: 20
depends_on:
- image-build-amd
source_file_dependencies:
@@ -296,7 +292,7 @@ steps:
- label: Async Engine, Inputs, Utils, Worker
device: h200_35gb
key: async-engine-inputs-utils-worker
timeout_in_minutes: 25
timeout_in_minutes: 50
source_file_dependencies:
- vllm/assets/
- vllm/config/
@@ -323,7 +319,7 @@ steps:
key: async-engine-inputs-utils-worker-config-cpu
depends_on:
- image-build-cpu
timeout_in_minutes: 65
timeout_in_minutes: 30
source_file_dependencies:
- vllm/assets/
- vllm/config/
@@ -355,7 +351,6 @@ steps:
- tests/test_outputs.py
- tests/test_pooling_params.py
- tests/test_ray_env.py
- tests/test_sampling_params.py
- tests/multimodal
- tests/renderers
- tests/standalone_tests/lazy_imports.py
@@ -373,7 +368,6 @@ steps:
- pytest -v -s test_outputs.py
- pytest -v -s test_pooling_params.py
- pytest -v -s test_ray_env.py
- pytest -v -s test_sampling_params.py
- pytest -v -s -m 'cpu_test' multimodal
- pytest -v -s renderers
- pytest -v -s reasoning
@@ -385,7 +379,7 @@ steps:
- label: Batch Invariance (A100)
key: batch-invariance-a100
timeout_in_minutes: 40
timeout_in_minutes: 30
device: a100
source_file_dependencies:
- vllm/v1/attention
@@ -399,7 +393,7 @@ steps:
- label: Batch Invariance (H100)
key: batch-invariance-h100
timeout_in_minutes: 40
timeout_in_minutes: 30
device: h100
source_file_dependencies:
- vllm/v1/attention
@@ -415,7 +409,7 @@ steps:
- label: Batch Invariance (B200)
key: batch-invariance-b200
timeout_in_minutes: 35
timeout_in_minutes: 30
device: b200-k8s
source_file_dependencies:
- vllm/v1/attention
@@ -434,7 +428,7 @@ steps:
- label: Acceptance Length Test (Large Models) # optional
device: h200_35gb
key: acceptance-length-test-large-models
timeout_in_minutes: 20
timeout_in_minutes: 25
gpu: h100
optional: true
num_gpus: 1
+1 -7
View File
@@ -3,16 +3,13 @@ depends_on:
- image-build
steps:
- label: Model Executor
device: h200_35gb
key: model-executor
timeout_in_minutes: 45
timeout_in_minutes: 35
source_file_dependencies:
- vllm/engine/arg_utils.py
- vllm/config/model.py
- vllm/model_executor
- vllm/model_executor/warmup
- tests/model_executor
- tests/model_executor/test_jit_warmup.py
- tests/entrypoints/openai/completion/test_tensorizer_entrypoint.py
commands:
- apt-get update && apt-get install -y curl libsodium23
@@ -28,7 +25,6 @@ steps:
- pytest -v -s entrypoints/openai/completion/test_tensorizer_entrypoint.py --timeout=900 --timeout-method=thread
mirror:
amd:
dind: false
device: mi300_1
depends_on:
- image-build-amd
@@ -36,9 +32,7 @@ steps:
- vllm/engine/arg_utils.py
- vllm/config/model.py
- vllm/model_executor
- vllm/model_executor/warmup
- tests/model_executor
- tests/model_executor/test_jit_warmup.py
- tests/entrypoints/openai/completion/test_tensorizer_entrypoint.py
- vllm/_aiter_ops.py
- vllm/platforms/rocm.py
+7 -5
View File
@@ -5,7 +5,7 @@ steps:
- label: Model Runner V2 Core Tests
device: h200_35gb
key: model-runner-v2-core-tests
timeout_in_minutes: 35
timeout_in_minutes: 45
source_file_dependencies:
- vllm/v1/worker/gpu/
- vllm/v1/worker/gpu_worker.py
@@ -18,7 +18,9 @@ steps:
- set -x
- export VLLM_USE_V2_MODEL_RUNNER=1
- pytest -v -s v1/engine/test_llm_engine.py -k "not test_engine_metrics"
- pytest -v -s v1/e2e/general/test_async_scheduling.py -k "not ngram"
# This requires eager until we sort out CG correctness issues.
# TODO: remove ENFORCE_EAGER here after https://github.com/vllm-project/vllm/pull/32936 is merged.
- ENFORCE_EAGER=1 pytest -v -s v1/e2e/general/test_async_scheduling.py -k "not ngram"
- pytest -v -s v1/e2e/general/test_context_length.py
- pytest -v -s v1/e2e/general/test_min_tokens.py
# Temporary hack filter to exclude ngram spec decoding based tests.
@@ -27,7 +29,7 @@ steps:
- label: Model Runner V2 Examples
device: h200_35gb
key: model-runner-v2-examples
timeout_in_minutes: 35
timeout_in_minutes: 45
working_dir: "/vllm-workspace/examples"
source_file_dependencies:
- vllm/v1/worker/gpu/
@@ -63,7 +65,7 @@ steps:
- label: Model Runner V2 Distributed (2 GPUs)
key: model-runner-v2-distributed-2-gpus
timeout_in_minutes: 30
timeout_in_minutes: 45
working_dir: "/vllm-workspace/tests"
num_devices: 2
source_file_dependencies:
@@ -84,7 +86,7 @@ steps:
- label: Model Runner V2 Pipeline Parallelism (4 GPUs)
key: model-runner-v2-pipeline-parallelism-4-gpus
timeout_in_minutes: 50
timeout_in_minutes: 60
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
+13 -9
View File
@@ -4,8 +4,9 @@ depends_on:
steps:
- label: Basic Models Tests (Initialization)
key: basic-models-tests-initialization
timeout_in_minutes: 25
timeout_in_minutes: 45
device: h200_18gb
torch_nightly: true
source_file_dependencies:
- vllm/
- tests/models/test_initialization.py
@@ -13,11 +14,13 @@ steps:
commands:
# Run a subset of model initialization tests
- pytest -v -s models/test_initialization.py::test_can_initialize_small_subset
mirror:
torch_nightly: {}
- label: Basic Models Tests (Extra Initialization) %N
device: h200_35gb
key: basic-models-tests-extra-initialization
timeout_in_minutes: 100
timeout_in_minutes: 45
source_file_dependencies:
- vllm/model_executor/models/
- tests/models/test_initialization.py
@@ -27,19 +30,21 @@ steps:
# subset of supported models (the complement of the small subset in the above
# test.) Also run if model initialization test file is modified
- pytest -v -s models/test_initialization.py -k 'not test_can_initialize_small_subset' --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
parallelism: 4
parallelism: 2
mirror:
torch_nightly: {}
- label: Basic Models Tests (Other)
device: h200_35gb
key: basic-models-tests-other
timeout_in_minutes: 35
timeout_in_minutes: 45
source_file_dependencies:
- vllm/
- tests/models/test_terratorch.py
- tests/models/transformers/test_backend.py
- tests/models/test_transformers.py
- tests/models/test_registry.py
commands:
- pytest -v -s models/test_terratorch.py models/transformers/test_backend.py models/test_registry.py
- pytest -v -s models/test_terratorch.py models/test_transformers.py models/test_registry.py
mirror:
amd:
device: mi325_1
@@ -50,12 +55,11 @@ steps:
key: basic-models-test-other-cpu
depends_on:
- image-build-cpu
timeout_in_minutes: 20
timeout_in_minutes: 10
source_file_dependencies:
- vllm/
- tests/models/test_utils.py
- tests/models/test_vision.py
- tests/models/transformers/fusers/
device: cpu-small
commands:
- pytest -v -s models/test_utils.py models/test_vision.py models/transformers/fusers/
- pytest -v -s models/test_utils.py models/test_vision.py
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Distributed Model Tests (2 GPUs)
key: distributed-model-tests-2-gpus
timeout_in_minutes: 60
timeout_in_minutes: 50
working_dir: "/vllm-workspace/tests"
num_devices: 2
source_file_dependencies:
@@ -17,7 +17,7 @@ steps:
- TARGET_TEST_SUITE=L4 pytest basic_correctness/ -v -s -m 'distributed(num_gpus=2)'
- CUDA_VISIBLE_DEVICES=0,1 pytest -v -s model_executor/model_loader/test_sharded_state_loader.py -m '(not slow_test)'
# Avoid importing model tests that cause CUDA reinitialization error
- pytest models/transformers/test_backend.py -v -s -m 'distributed(num_gpus=2)'
- pytest models/test_transformers.py -v -s -m 'distributed(num_gpus=2)'
- pytest models/language -v -s -m 'distributed(num_gpus=2)'
- pytest models/multimodal/generation/test_phi4siglip.py -v -s -m 'distributed(num_gpus=2)'
- pytest models/multimodal -v -s -m 'distributed(num_gpus=2)' --ignore models/multimodal/generation/test_whisper.py --ignore models/multimodal/generation/test_phi4siglip.py
+16 -30
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Language Models Tests (Standard)
key: language-models-tests-standard
timeout_in_minutes: 30
timeout_in_minutes: 25
device: h200_18gb
source_file_dependencies:
- vllm/
@@ -14,16 +14,15 @@ steps:
- pip freeze | grep -E 'torch'
- pytest -v -s models/language -m 'core_model and (not slow_test)'
mirror:
torch_nightly: {}
amd:
dind: false
device: mi300_1
depends_on:
- image-build-amd
- label: Language Models Tests (Extra Standard) %N
device: h200_35gb
key: language-models-tests-extra-standard
timeout_in_minutes: 40
timeout_in_minutes: 45
source_file_dependencies:
- vllm/model_executor/models/
- tests/models/language/pooling/test_embedding.py
@@ -36,8 +35,8 @@ steps:
- pytest -v -s models/language -m 'core_model and slow_test' --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
parallelism: 2
mirror:
torch_nightly: {}
amd:
dind: false
device: mi300_1
depends_on:
- image-build-amd
@@ -52,10 +51,10 @@ steps:
- tests/models/language/pooling/test_classification.py
- vllm/_aiter_ops.py
- vllm/platforms/rocm.py
- label: Language Models Tests (Hybrid) %N
device: h200_35gb
key: language-models-tests-hybrid
timeout_in_minutes: 65
timeout_in_minutes: 75
source_file_dependencies:
- vllm/
- tests/models/language/generation
@@ -64,13 +63,14 @@ steps:
# Note: also needed to run plamo2 model in vLLM
- uv pip install --system --no-build-isolation 'git+https://github.com/state-spaces/mamba@v2.3.0'
- uv pip install --system --no-build-isolation 'git+https://github.com/Dao-AILab/causal-conv1d@v1.6.0'
# Shard the hybrid language model tests that are numerically stable on Hopper.
- pytest -v -s models/language/generation -m hybrid_model -k 'not granite-4.0-tiny-preview' --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
# Shard hybrid language model tests
- pytest -v -s models/language/generation -m hybrid_model --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
parallelism: 2
mirror:
torch_nightly: {}
amd:
device: mi325_1
timeout_in_minutes: 70
timeout_in_minutes: 90
depends_on:
- image-build-amd
commands:
@@ -78,24 +78,10 @@ steps:
- uv pip install --system --no-build-isolation 'git+https://github.com/Dao-AILab/causal-conv1d@v1.6.0'
- pytest -v -s models/language/generation -m hybrid_model --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
# Granite 4 hybrid generation is sensitive to hardware-specific Triton SSD
# autotuning (https://github.com/vllm-project/vllm/issues/25194). Keep this one
# correctness test on L4 until its H200 output matches the Transformers reference.
- label: Language Models Tests (Granite L4 Compatibility)
key: language-models-tests-granite-l4-compatibility
timeout_in_minutes: 65
source_file_dependencies:
- vllm/
- tests/models/language/generation
commands:
- uv pip install --system --no-build-isolation 'git+https://github.com/state-spaces/mamba@v2.3.0'
- uv pip install --system --no-build-isolation 'git+https://github.com/Dao-AILab/causal-conv1d@v1.6.0'
- pytest -v -s models/language/generation -m hybrid_model -k 'granite-4.0-tiny-preview'
- label: Language Models Test (Extended Generation) # 80min
device: h200_35gb
key: language-models-test-extended-generation
timeout_in_minutes: 65
timeout_in_minutes: 110
optional: true
source_file_dependencies:
- vllm/
@@ -109,7 +95,7 @@ steps:
- label: Language Models Test (PPL)
key: language-models-test-ppl
timeout_in_minutes: 30
timeout_in_minutes: 110
device: h200_18gb
optional: true
source_file_dependencies:
@@ -118,10 +104,10 @@ steps:
commands:
- pytest -v -s models/language/generation_ppl_test
- label: Language Models Test (Extended Pooling)
- label: Language Models Test (Extended Pooling) # 36min
device: h200_35gb
key: language-models-test-extended-pooling
timeout_in_minutes: 120
timeout_in_minutes: 50
optional: true
source_file_dependencies:
- vllm/
@@ -131,13 +117,13 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 120
timeout_in_minutes: 100
depends_on:
- image-build-amd
- label: Language Models Test (MTEB)
key: language-models-test-mteb
timeout_in_minutes: 45
timeout_in_minutes: 110
device: h200_18gb
optional: true
source_file_dependencies:
+10 -14
View File
@@ -20,14 +20,13 @@ steps:
- label: "Multi-Modal Models (Standard) 2: qwen3 + gemma"
key: multi-modal-models-standard-2-qwen3-gemma
timeout_in_minutes: 50
timeout_in_minutes: 45
device: h200_18gb
source_file_dependencies:
- vllm/
- tests/models/multimodal
commands:
- pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "qwen3 or gemma"
- pytest -v -s models/multimodal/generation/test_mm_prefix_lm.py -m core_model
- pytest -v -s models/multimodal/generation/test_qwen2_5_vl.py -m core_model
mirror:
amd:
@@ -38,7 +37,7 @@ steps:
- label: "Multi-Modal Models (Standard) 3: llava + qwen2_vl"
device: h200_35gb
key: multi-modal-models-standard-3-llava-qwen2-vl
timeout_in_minutes: 40
timeout_in_minutes: 45
source_file_dependencies:
- vllm/
- tests/models/multimodal
@@ -54,12 +53,12 @@ steps:
- label: "Multi-Modal Models (Standard) 4: other + whisper"
device: h200_35gb
key: multi-modal-models-standard-4-other-whisper
timeout_in_minutes: 50
timeout_in_minutes: 45
source_file_dependencies:
- vllm/
- tests/models/multimodal
commands:
- 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_mm_prefix_lm.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 -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 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
@@ -69,23 +68,22 @@ steps:
depends_on:
- image-build-amd
- label: Multi-Modal Processor (CPU) %N
- label: Multi-Modal Processor (CPU)
key: multi-modal-processor-cpu
depends_on:
- image-build-cpu
timeout_in_minutes: 125
timeout_in_minutes: 60
source_file_dependencies:
- vllm/
- tests/models/multimodal
- tests/models/registry.py
device: cpu-medium
commands:
- pytest -v -s models/multimodal/processing --ignore models/multimodal/processing/test_tensor_schema.py --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
parallelism: 4
- pytest -v -s models/multimodal/processing --ignore models/multimodal/processing/test_tensor_schema.py
- label: Multi-Modal Processor # 44min
key: multi-modal-processor
timeout_in_minutes: 65
timeout_in_minutes: 60
device: h200_18gb
source_file_dependencies:
- vllm/
@@ -97,7 +95,7 @@ steps:
- label: Multi-Modal Accuracy Eval (Small Models) # 50min
device: h200_35gb
key: multi-modal-accuracy-eval-small-models
timeout_in_minutes: 30
timeout_in_minutes: 70
working_dir: "/vllm-workspace/.buildkite/lm-eval-harness"
source_file_dependencies:
- vllm/multimodal/
@@ -107,7 +105,6 @@ steps:
- pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-mm-small.txt --tp-size=1
mirror:
amd:
dind: false
device: mi300_1
depends_on:
- image-build-amd
@@ -119,7 +116,6 @@ steps:
- vllm/model_executor/model_loader/
- label: Multi-Modal Models (Extended Generation 1)
device: h200_35gb
key: multi-modal-models-extended-generation-1
optional: true
source_file_dependencies:
@@ -167,7 +163,7 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 75
timeout_in_minutes: 60
depends_on:
- image-build-amd
source_file_dependencies:
+1 -6
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Plugin Tests (2 GPUs)
key: plugin-tests-2-gpus
timeout_in_minutes: 35
timeout_in_minutes: 60
working_dir: "/vllm-workspace/tests"
num_devices: 2
source_file_dependencies:
@@ -37,11 +37,6 @@ steps:
- pytest -v -s plugins_tests/test_stats_logger_plugins.py
- pip uninstall dummy_stat_logger -y
# end stat_logger plugins test
# begin endpoint plugins test
- pip install -e ./plugins/vllm_add_dummy_endpoint_plugin
- pytest -v -s plugins_tests/test_endpoint_plugins.py
- pip uninstall vllm_add_dummy_endpoint_plugin -y
# end endpoint plugins test
# other tests continue here:
- pytest -v -s plugins_tests/test_scheduler_plugins.py
- pip install -e ./plugins/vllm_add_dummy_model
+6 -44
View File
@@ -5,7 +5,7 @@ steps:
- label: PyTorch Compilation Unit Tests
device: h200_35gb
key: pytorch-compilation-unit-tests
timeout_in_minutes: 90
timeout_in_minutes: 10
source_file_dependencies:
- vllm/__init__.py
- vllm/_aiter_ops.py
@@ -78,7 +78,7 @@ steps:
- label: PyTorch Compilation Passes Unit Tests
key: pytorch-compilation-passes-unit-tests
timeout_in_minutes: 45
timeout_in_minutes: 20
source_file_dependencies:
- vllm/__init__.py
- vllm/_aiter_ops.py
@@ -109,16 +109,14 @@ steps:
- pytest -s -v compile/passes --ignore compile/passes/distributed
mirror:
amd:
dind: false
device: mi300_1
timeout_in_minutes: 65
timeout_in_minutes: 180
depends_on:
- image-build-amd
- label: PyTorch Fullgraph Smoke Test
device: h200_35gb
key: pytorch-fullgraph-smoke-test
timeout_in_minutes: 90
timeout_in_minutes: 35
source_file_dependencies:
- vllm/__init__.py
- vllm/_aiter_ops.py
@@ -150,46 +148,11 @@ steps:
# as it is a heavy test that is covered in other steps.
# Use `find` to launch multiple instances of pytest so that
# they do not suffer from https://github.com/vllm-project/vllm/issues/28965
- "find compile/fullgraph/ -name 'test_*.py' -not -name 'test_full_cudagraph.py' -not -name 'test_full_graph.py' -print0 | xargs -0 -n1 -I{} pytest -s -v '{}'"
# Hopper-only DeepSeek-V2-Lite cases in this file require two 29.3-GiB model
# instances and cannot fit a 35GB MIG slice. L4 retains the original coverage:
# those SM90 cases skip while the architecture-compatible cases still run.
- label: PyTorch Fullgraph CUDAGraph (L4 Compatibility)
key: pytorch-fullgraph-cudagraph-l4-compatibility
timeout_in_minutes: 60
source_file_dependencies:
- vllm/__init__.py
- vllm/_aiter_ops.py
- vllm/_custom_ops.py
- vllm/compilation/
- vllm/config/
- vllm/distributed/
- vllm/engine/
- vllm/env_override.py
- vllm/envs.py
- vllm/forward_context.py
- vllm/inputs/
- vllm/ir/
- vllm/kernels/
- vllm/logger.py
- vllm/model_executor/
- vllm/multimodal/
- vllm/platforms/
- vllm/plugins/
- vllm/sampling_params.py
- vllm/sequence.py
- vllm/transformers_utils/
- vllm/triton_utils/
- vllm/utils/
- vllm/v1/
- tests/compile
commands:
- pytest -s -v compile/fullgraph/test_full_cudagraph.py
- "find compile/fullgraph/ -name 'test_*.py' -not -name 'test_full_graph.py' -print0 | xargs -0 -n1 -I{} pytest -s -v '{}'"
- label: PyTorch Fullgraph
key: pytorch-fullgraph
timeout_in_minutes: 40
timeout_in_minutes: 30
device: h200_18gb
source_file_dependencies:
- vllm/__init__.py
@@ -234,7 +197,6 @@ steps:
- bash standalone_tests/pytorch_nightly_dependency.sh
mirror:
amd:
dind: false
device: mi300_1
depends_on:
- image-build-amd
+6 -17
View File
@@ -3,11 +3,8 @@ depends_on:
- image-build
steps:
- label: Quantization
device: h200_35gb
key: quantization
timeout_in_minutes: 75
env:
VLLM_USE_V2_MODEL_RUNNER: "0"
timeout_in_minutes: 90
source_file_dependencies:
- csrc/
- vllm/model_executor/layers/quantization
@@ -22,15 +19,11 @@ steps:
# TODO(jerryzh168): resolve the above comment
- uv pip install --system torchao==0.17.0 --index-url https://download.pytorch.org/whl/cu130
- uv pip install --system conch-triton-kernels
# The SM90-only checkpoint currently contains a removed weight_chan_scale
# parameter. It was not exercised by the previous L4 job.
- VLLM_TEST_FORCE_LOAD_FORMAT=auto pytest -v -s quantization/ --ignore quantization/test_blackwell_moe.py -k 'not test_compressed_tensors_w4a8_fp8' --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
parallelism: 8
- VLLM_TEST_FORCE_LOAD_FORMAT=auto pytest -v -s quantization/ --ignore quantization/test_blackwell_moe.py
- label: Quantized Fusions
device: h200_35gb
key: quantized-fusions
timeout_in_minutes: 20
timeout_in_minutes: 30
source_file_dependencies:
- tests/fusion
- vllm/model_executor/layers/fusion
@@ -42,7 +35,7 @@ steps:
- label: Quantized MoE Test (B200)
key: quantized-moe-test-b200
timeout_in_minutes: 120
timeout_in_minutes: 60
working_dir: "/vllm-workspace/"
device: b200-k8s
source_file_dependencies:
@@ -59,14 +52,10 @@ steps:
- pytest -s -v tests/quantization/test_blackwell_moe.py
- label: Quantized Models Test
device: h200_35gb
key: quantized-models-test
timeout_in_minutes: 65
env:
VLLM_USE_V2_MODEL_RUNNER: "0"
timeout_in_minutes: 60
source_file_dependencies:
- vllm/model_executor/layers/quantization
- tests/models/quantization
commands:
- pytest -v -s models/quantization --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
parallelism: 3
- pytest -v -s models/quantization
+13 -51
View File
@@ -3,16 +3,11 @@ depends_on:
- image-build
steps:
- label: Rust Frontend OpenAI Coverage
timeout_in_minutes: 30
timeout_in_minutes: 90
device: h200_18gb
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- rust/
- build_rust.sh
- tools/build_rust.py
- rust-toolchain.toml
- .buildkite/scripts/rust-coverage.sh
- codecov.yml
- vllm/benchmarks/
- vllm/entrypoints/openai/
- vllm/entrypoints/serve/
@@ -20,116 +15,84 @@ steps:
- tests/utils.py
- tests/benchmarks/test_serve_cli.py
- tests/entrypoints/openai/chat_completion/test_chat_completion.py
- tests/entrypoints/openai/chat_completion/test_chat_logit_bias_validation.py
# - tests/entrypoints/openai/chat_completion/test_chat_logit_bias_validation.py
# - tests/entrypoints/openai/completion/test_prompt_validation.py
- tests/entrypoints/openai/completion/test_shutdown.py
- tests/entrypoints/openai/test_return_token_ids.py
- tests/entrypoints/openai/test_uds.py
# - tests/entrypoints/openai/test_return_token_ids.py
# - tests/entrypoints/openai/test_uds.py
- tests/v1/sample/test_logprobs_e2e.py
commands:
- . /vllm-workspace/.buildkite/scripts/rust-coverage.sh && rust_coverage_start rust-e2e
- export VLLM_USE_RUST_FRONTEND=1
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -v -s benchmarks/test_serve_cli.py -k "not insecure and not (test_bench_serve and not test_bench_serve_chat)"
- pytest -v -s entrypoints/openai/chat_completion/test_chat_completion.py -k "not test_invalid_json_schema and not test_invalid_regex"
- pytest -v -s entrypoints/openai/chat_completion/test_chat_logit_bias_validation.py -k "not multiple"
# - pytest -v -s entrypoints/openai/chat_completion/test_chat_logit_bias_validation.py -k "not invalid"
# - pytest -v -s entrypoints/openai/completion/test_prompt_validation.py -k "not prompt_embeds"
- pytest -v -s entrypoints/openai/completion/test_shutdown.py -k "not engine_failure and not test_abort_timeout_exits_quickly"
# test_comparison streams differently: Rust emits a separate first (prompt_token_ids) chunk and
# finish chunk without logprobs, while the test reads `logprobs.tokens` on every chunk.
- pytest -v -s entrypoints/openai/test_return_token_ids.py -k "not test_comparison"
- pytest -v -s entrypoints/openai/test_uds.py
# - pytest -v -s entrypoints/openai/test_return_token_ids.py
# - pytest -v -s entrypoints/openai/test_uds.py
- pytest -v -s v1/sample/test_logprobs_e2e.py -k "test_prompt_logprobs_e2e_server"
- label: Rust Frontend Serve/Admin Coverage
timeout_in_minutes: 25
timeout_in_minutes: 60
device: h200_18gb
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- rust/
- build_rust.sh
- tools/build_rust.py
- rust-toolchain.toml
- .buildkite/scripts/rust-coverage.sh
- codecov.yml
- vllm/entrypoints/openai/
- vllm/entrypoints/serve/
- vllm/v1/engine/
- tests/utils.py
- tests/entrypoints/serve/dev/rpc/test_collective_rpc.py
# - tests/entrypoints/serve/dev/rpc/test_collective_rpc.py
- tests/entrypoints/scale_out/token_in_token_out/test_serving_tokens.py
- tests/entrypoints/serve/instrumentator/test_basic.py
- tests/entrypoints/serve/instrumentator/test_metrics.py
# - tests/entrypoints/serve/dev/test_sleep.py
- tests/entrypoints/serve/tokenize/test_tokenization.py
commands:
- . /vllm-workspace/.buildkite/scripts/rust-coverage.sh && rust_coverage_start rust-e2e
- export VLLM_USE_RUST_FRONTEND=1
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/serve/dev/rpc/test_collective_rpc.py
# server_load can be flaky under the Rust frontend; keep it excluded for now.
# - pytest -v -s entrypoints/serve/dev/rpc/test_collective_rpc.py
- pytest -v -s entrypoints/serve/instrumentator/test_basic.py -k "not show_version and not server_load"
# test_generate_logprobs expects Python-style top_logprobs truncation (dedup sampled + cap at max(k, 1)).
- pytest -v -s entrypoints/scale_out/token_in_token_out/test_serving_tokens.py -k "not stream and not lora and not test_generate_logprobs and not stop_string_workflow"
- pytest -v -s entrypoints/serve/instrumentator/test_metrics.py -k "text and not show and not run_batch and not test_metrics_counts and not test_metrics_exist"
# - pytest -v -s entrypoints/serve/dev/test_sleep.py
# /tokenizer_info is not implemented in the Rust frontend (the CLI flag is accepted as a no-op).
- pytest -v -s entrypoints/serve/tokenize/test_tokenization.py -k "not tokenizer_info"
- label: Rust Frontend Core Correctness
timeout_in_minutes: 20
timeout_in_minutes: 30
device: h200_18gb
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- rust/
- build_rust.sh
- tools/build_rust.py
- rust-toolchain.toml
- .buildkite/scripts/rust-coverage.sh
- codecov.yml
- vllm/entrypoints/openai/
- tests/utils.py
- tests/entrypoints/openai/correctness/test_lmeval.py
commands:
- . /vllm-workspace/.buildkite/scripts/rust-coverage.sh && rust_coverage_start rust-e2e
- export VLLM_USE_RUST_FRONTEND=1
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine
- label: Rust Frontend Tool Use
device: h200_35gb
timeout_in_minutes: 25
timeout_in_minutes: 60
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- rust/
- build_rust.sh
- tools/build_rust.py
- rust-toolchain.toml
- .buildkite/scripts/rust-coverage.sh
- codecov.yml
- vllm/entrypoints/openai/
- vllm/tool_parsers/
- tests/utils.py
- tests/tool_use/
commands:
- . /vllm-workspace/.buildkite/scripts/rust-coverage.sh && rust_coverage_start rust-e2e
- export VLLM_USE_RUST_FRONTEND=1
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -v -s tool_use --ignore=tool_use/mistral --models llama3.2 -k "not test_response_format_with_tool_choice_required and not test_parallel_tool_calls_false and not test_tool_call_and_choice"
- label: Rust Frontend Distributed
timeout_in_minutes: 25
timeout_in_minutes: 30
num_devices: 4
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- rust/
- build_rust.sh
- tools/build_rust.py
- rust-toolchain.toml
- .buildkite/scripts/rust-coverage.sh
- codecov.yml
- vllm/distributed/
- vllm/engine/
- vllm/executor/
@@ -140,7 +103,6 @@ steps:
- tests/v1/distributed/test_hybrid_lb_dp.py
- tests/v1/distributed/test_internal_lb_dp.py
commands:
- . /vllm-workspace/.buildkite/scripts/rust-coverage.sh && rust_coverage_start rust-e2e
- export VLLM_USE_RUST_FRONTEND=1
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- export NCCL_CUMEM_HOST_ENABLE=0
@@ -4,7 +4,7 @@ steps:
- label: Rust Frontend Cargo Style + Clippy
key: rust-frontend-cargo-style-clippy
depends_on: []
timeout_in_minutes: 20
timeout_in_minutes: 30
device: cpu-medium
no_plugin: true
source_file_dependencies:
@@ -18,7 +18,7 @@ steps:
- label: Rust Frontend Cargo Tests
key: rust-frontend-cargo-tests
depends_on: []
timeout_in_minutes: 20
timeout_in_minutes: 30
device: cpu-medium
no_plugin: true
source_file_dependencies:
@@ -26,7 +26,5 @@ steps:
- rust-toolchain.toml
- .buildkite/test_areas/rust_frontend_cargo.yaml
- .buildkite/scripts/run-rust-frontend-cargo-ci.sh
- .buildkite/scripts/rust-coverage.sh
- codecov.yml
commands:
- .buildkite/scripts/run-rust-frontend-cargo-ci.sh test
+2 -2
View File
@@ -5,7 +5,7 @@ steps:
- label: Samplers Test
device: h200_35gb
key: samplers-test
timeout_in_minutes: 40
timeout_in_minutes: 75
source_file_dependencies:
- vllm/model_executor/layers
- vllm/sampling_metadata.py
@@ -19,7 +19,7 @@ steps:
- VLLM_USE_FLASHINFER_SAMPLER=1 pytest -v -s samplers
mirror:
amd:
device: mi325_1
device: mi250_1
depends_on:
- image-build-amd
commands:
+11 -11
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Spec Decode Eagle
key: spec-decode-eagle
timeout_in_minutes: 25
timeout_in_minutes: 30
device: h200_18gb
source_file_dependencies:
- vllm/v1/spec_decode/
@@ -15,7 +15,7 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 60
timeout_in_minutes: 45
depends_on:
- image-build-amd
source_file_dependencies:
@@ -29,7 +29,7 @@ steps:
- label: Spec Decode Eagle Nightly B200
key: spec-decode-eagle-nightly-b200
timeout_in_minutes: 25
timeout_in_minutes: 30
device: b200-k8s
optional: true
source_file_dependencies:
@@ -41,7 +41,7 @@ steps:
- label: Spec Decode Speculators + MTP
key: spec-decode-speculators-mtp
timeout_in_minutes: 20
timeout_in_minutes: 30
device: h200_18gb
source_file_dependencies:
- vllm/v1/spec_decode/
@@ -82,7 +82,7 @@ steps:
- label: Spec Decode Ngram + Suffix
key: spec-decode-ngram-suffix
timeout_in_minutes: 20
timeout_in_minutes: 30
device: h200_18gb
source_file_dependencies:
- vllm/v1/spec_decode/
@@ -93,7 +93,7 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 55
timeout_in_minutes: 65
# TODO(akaratza): Test after Torch >= 2.12 bump
soft_fail: true
depends_on:
@@ -109,7 +109,7 @@ steps:
- label: Spec Decode Draft Model
key: spec-decode-draft-model
timeout_in_minutes: 45
timeout_in_minutes: 30
device: h200_18gb
source_file_dependencies:
- vllm/v1/spec_decode/
@@ -120,7 +120,7 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 55
timeout_in_minutes: 50
depends_on:
- image-build-amd
source_file_dependencies:
@@ -134,7 +134,7 @@ steps:
- label: Spec Decode Draft Model Nightly B200
key: spec-decode-draft-model-nightly-b200
timeout_in_minutes: 40
timeout_in_minutes: 30
device: b200-k8s
optional: true
source_file_dependencies:
@@ -146,7 +146,7 @@ steps:
- label: Speculators Correctness
key: speculators-correctness
timeout_in_minutes: 30
timeout_in_minutes: 60
device: h100
optional: true
num_devices: 1
@@ -159,7 +159,7 @@ steps:
- pytest -v -s v1/spec_decode/test_speculators_correctness.py -m slow_test
- label: Spec Decode MTP hybrid (B200)
timeout_in_minutes: 20
timeout_in_minutes: 30
device: b200-k8s
optional: true
source_file_dependencies:
+1 -2
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Weight Loading Multiple GPU # 33min
key: weight-loading-multiple-gpu
timeout_in_minutes: 50
timeout_in_minutes: 45
working_dir: "/vllm-workspace/tests"
num_devices: 2
optional: true
@@ -15,7 +15,6 @@ steps:
- bash weight_loading/run_model_weight_loading_test.sh -c weight_loading/models.txt
mirror:
amd:
dind: false
device: mi300_2
depends_on:
- image-build-amd
+2 -3
View File
@@ -47,7 +47,6 @@
# Rust Frontend
/rust/ @BugenZhao @njhill
/rust/src/bench @esmeetu
/build_rust.sh @BugenZhao @njhill
/rust-toolchain.toml @BugenZhao @njhill
/.buildkite/test_areas/rust* @BugenZhao @njhill
@@ -120,7 +119,7 @@
# Transformers modeling backend
/vllm/model_executor/models/transformers @hmellor
/tests/models/transformers @hmellor
/tests/models/test_transformers.py @hmellor
# Docs
/docs/mkdocs @hmellor
@@ -173,7 +172,7 @@ mkdocs.yaml @hmellor
# Kernels
/vllm/v1/attention/ops/chunked_prefill_paged_decode.py @tdoublep
/vllm/v1/attention/ops/triton_unified_attention.py @tdoublep
/vllm/third_party/flash_linear_attention @ZJY0516 @vadiklyutiy
/vllm/model_executor/layers/fla @ZJY0516 @vadiklyutiy
# ROCm related: specify owner with write access to notify AMD folks for careful code review
/vllm/**/*rocm* @tjtanaa @dllehr-amd
+1 -1
View File
@@ -323,7 +323,7 @@ jobs:
// {users} will be replaced with @mentions
const ccConfig = {
rocm: {
users: ['hongxiayang', 'tjtanaa', 'vllmellm', 'giuseppegrossi'],
users: ['hongxiayang', 'tjtanaa', 'vllmellm'],
message: 'CC {users} for ROCm-related issue',
},
mistral: {
+2 -3
View File
@@ -28,8 +28,7 @@ jobs:
pull_number: context.payload.pull_request.number,
});
const readyLabels = ['ready', 'ready-run-all-tests'];
const hasReadyLabel = pr.labels.some(l => readyLabels.includes(l.name));
const hasReadyLabel = pr.labels.some(l => l.name === 'ready');
const hasVerifiedLabel = pr.labels.some(l => l.name === 'verified');
const { data: mergedPRs } = await github.rest.search.issuesAndPullRequests({
@@ -41,7 +40,7 @@ jobs:
if (hasReadyLabel || hasVerifiedLabel || mergedCount >= 4) {
core.info(`Check passed: verified label=${hasVerifiedLabel}, ready label=${hasReadyLabel}, 4+ merged PRs=${mergedCount >= 4}`);
} else {
core.setFailed(`PR must have the 'verified', 'ready', or 'ready-run-all-tests' label (the ready labels also trigger tests) or the author must have at least 4 merged PRs (found ${mergedCount}).`);
core.setFailed(`PR must have the 'verified' or 'ready' (which also triggers tests) label or the author must have at least 4 merged PRs (found ${mergedCount}).`);
}
pre-commit:
-4
View File
@@ -18,9 +18,6 @@ vllm/third_party/deep_gemm/
# fmha_sm100 vendored package built from source
vllm/third_party/fmha_sm100/
# tml-fa4 vendored package built from source
vllm/third_party/tml_fa4/
# triton jit
.triton
@@ -257,4 +254,3 @@ vllm/grpc/vllm_engine_pb2.pyi
# Ignore generated cpu headers
csrc/cpu/cpu_attn_dispatch_generated.h
rust-coverage-tools/
+2 -2
View File
@@ -30,7 +30,7 @@ repos:
- id: markdownlint-cli2
language_version: lts
args: [--fix]
exclude: (^|/)CLAUDE\.md$
exclude: ^CLAUDE\.md$
- repo: https://github.com/rhysd/actionlint
rev: v1.7.7
hooks:
@@ -210,7 +210,7 @@ repos:
name: Check SPDX headers
entry: python tools/pre_commit/check_spdx_header.py
language: python
types_or: [python, rust, proto]
types: [python]
- id: check-root-lazy-imports
name: Check root lazy imports
entry: python tools/pre_commit/check_init_lazy_imports.py
+2
View File
@@ -0,0 +1,2 @@
collect_env.py
vllm/model_executor/layers/fla/ops/*.py
+18 -29
View File
@@ -29,7 +29,6 @@ Do not open one-off PRs for tiny edits (single typo, isolated style change, one
- PR descriptions for AI-assisted work **must** include:
- Why this is not duplicating an existing PR.
- Test commands run and results.
- Model evaluation results when the change affects output, accuracy, or serving.
- Clear statement that AI assistance was used.
### Fail-closed behavior
@@ -67,38 +66,23 @@ VLLM_USE_PRECOMPILED=1 uv pip install -e . --torch-backend=auto
uv pip install -e . --torch-backend=auto
```
### Tests
### Running tests
> Requires [Environment setup](#environment-setup) and [Installing dependencies](#installing-dependencies).
```bash
# Install test dependencies (use cuda.in on non-x86_64):
uv pip install -r requirements/test/cuda.in
# Install test dependencies.
# requirements/test/cuda.txt is pinned to x86_64; on other platforms, use the
# unpinned source file instead:
uv pip install -r requirements/test/cuda.in # resolves for current platform
# Or on x86_64:
uv pip install -r requirements/test/cuda.txt
# Run a specific test file:
# Run a specific test file (use .venv/bin/python directly;
# `source activate` does not persist in non-interactive shells):
.venv/bin/python -m pytest tests/path/to/test_file.py -v
```
When adding tests:
- **Design before you write.** Answer four questions first: what is the module
for, what is its I/O contract, what failure am I guarding against, and what is
the cheapest level that catches it (unit over integration over e2e)?
- **Reuse before create.** Extend existing test files, `conftest.py` fixtures, and
helpers; add a new file only when no nearby suite fits.
- **Test behavior with intent.** Assert observable outcomes through public APIs;
state why in the name or docstring. Skip trivial wiring; flaky tests are worse
than no tests.
- **Keep it minimal.** One behavior per test and the smallest setup that
triggers it; if the test diff dwarfs the code change, cut scope.
- **No one-off kernel benchmarks in `tests/`.** Put kernel perf work in
`benchmarks/kernels/`; prove correctness in existing pytest suites.
- **Run model evals for model-affecting changes.** Search `tests/evals/` or use
`vllm bench` and include results in the PR — do not wait for reviewers to ask.
For model-specific requirements, see
[`docs/contributing/model/tests.md`](docs/contributing/model/tests.md).
### Running linters
> Requires [Environment setup](#environment-setup).
@@ -123,18 +107,23 @@ Use [Google-style docstrings](https://google.github.io/styleguide/pyguide.html#3
### Coding style guidelines
- Match existing code style
- Minimize use of comments. Eliminate comments which are redundant, preferring legible and self-documenting code. When used, keep docstrings and comments brief and direct.
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:`):
Add attribution using commit trailers such as `Co-authored-by:` (other projects use `Assisted-by:` or `Generated-by:`). For example:
```text
Your commit message here
Co-authored-by: Agent Name Here
Co-authored-by: GitHub Copilot
Co-authored-by: Claude
Co-authored-by: gemini-code-assist
Signed-off-by: Your Name <your.email@example.com>
```
+34 -31
View File
@@ -70,15 +70,6 @@ endif()
#
set(TORCH_SUPPORTED_VERSION_CUDA "2.11.0")
set(TORCH_SUPPORTED_VERSION_ROCM "2.11.0")
# TORCH_NIGHTLY=1 builds run against unpinned nightly wheels, so the supported-
# version check would always warn. Only treat it as a nightly build when the
# value is exactly "1" (the bootstrap exports TORCH_NIGHTLY=0 by default, which
# must NOT suppress the warning for normal builds).
if (DEFINED ENV{TORCH_NIGHTLY} AND "$ENV{TORCH_NIGHTLY}" STREQUAL "1")
set(TORCH_NIGHTLY_BUILD TRUE)
else()
set(TORCH_NIGHTLY_BUILD FALSE)
endif()
#
# Try to find python package with an executable that exactly matches
@@ -184,7 +175,7 @@ endif()
if (NOT HIP_FOUND AND NOT PYTORCH_FOUND_HIP AND CUDA_FOUND)
set(VLLM_GPU_LANG "CUDA")
if (NOT TORCH_NIGHTLY_BUILD AND NOT Torch_VERSION VERSION_EQUAL ${TORCH_SUPPORTED_VERSION_CUDA})
if (NOT Torch_VERSION VERSION_EQUAL ${TORCH_SUPPORTED_VERSION_CUDA})
message(WARNING "Pytorch version ${TORCH_SUPPORTED_VERSION_CUDA} "
"expected for CUDA build, saw ${Torch_VERSION} instead.")
endif()
@@ -197,7 +188,7 @@ elseif(HIP_FOUND OR PYTORCH_FOUND_HIP)
enable_language(HIP)
# ROCm 5.X and 6.X
if (NOT TORCH_NIGHTLY_BUILD AND ROCM_VERSION_DEV_MAJOR GREATER_EQUAL 5 AND
if (ROCM_VERSION_DEV_MAJOR GREATER_EQUAL 5 AND
Torch_VERSION VERSION_LESS ${TORCH_SUPPORTED_VERSION_ROCM})
message(WARNING "Pytorch version >= ${TORCH_SUPPORTED_VERSION_ROCM} "
"expected for ROCm build, saw ${Torch_VERSION} instead.")
@@ -359,7 +350,9 @@ endif()
if(VLLM_GPU_LANG STREQUAL "HIP")
set(VLLM_EXT_SRC
"csrc/torch_bindings.cpp"
"csrc/custom_quickreduce.cu")
"csrc/custom_quickreduce.cu"
"csrc/cuda_view.cu"
"csrc/libtorch_stable/cuda_utils_kernels.cu")
message(STATUS "Enabling C extension.")
define_extension_target(
@@ -387,10 +380,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
#
set(VLLM_STABLE_EXT_SRC
"csrc/libtorch_stable/torch_bindings.cpp"
"csrc/libtorch_stable/cuda_view.cu"
"csrc/libtorch_stable/cuda_utils_kernels.cu"
"csrc/libtorch_stable/activation_kernels.cu"
"csrc/libtorch_stable/ngram_embedding_kernels.cu"
"csrc/libtorch_stable/quantization/activation_kernels.cu"
"csrc/libtorch_stable/quantization/w8a8/int8/scaled_quant.cu"
"csrc/libtorch_stable/quantization/w8a8/fp8/common.cu"
@@ -409,6 +399,9 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
"csrc/libtorch_stable/sampler.cu"
"csrc/libtorch_stable/topk.cu"
"csrc/libtorch_stable/mamba/selective_scan_fwd.cu"
"csrc/libtorch_stable/attention/paged_attention_v1.cu"
"csrc/libtorch_stable/attention/paged_attention_v2.cu"
"csrc/libtorch_stable/cache_kernels.cu"
"csrc/libtorch_stable/cache_kernels.cu"
"csrc/libtorch_stable/cache_kernels_fused.cu"
"csrc/libtorch_stable/custom_all_reduce.cu"
@@ -466,6 +459,8 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
FetchContent_MakeAvailable(cutlass)
list(APPEND VLLM_STABLE_EXT_SRC
"csrc/libtorch_stable/cuda_view.cu"
"csrc/libtorch_stable/cuda_utils_kernels.cu"
"csrc/libtorch_stable/cutlass_extensions/common.cpp"
"csrc/libtorch_stable/quantization/w8a8/cutlass/scaled_mm_entry.cu"
"csrc/libtorch_stable/quantization/fp4/nvfp4_quant_entry.cu"
@@ -1096,15 +1091,14 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
USE_SABI 3
WITH_SOABI)
# Set TORCH_TARGET_VERSION for stable ABI compatibility.
# This ensures we only use C-shim APIs available in PyTorch 2.11.
# _C_stable_libtorch is abi compatible with PyTorch >= TORCH_TARGET_VERSION
# which is currently set to 2.11.
target_compile_definitions(_C_stable_libtorch PRIVATE
TORCH_TARGET_VERSION=0x020B000000000000ULL)
# Needed to use cuda/hip APIs from C-shim
if(VLLM_GPU_LANG STREQUAL "CUDA")
# Set TORCH_TARGET_VERSION for stable ABI compatibility.
# This ensures we only use C-shim APIs available in PyTorch 2.11.
# _C_stable_libtorch is abi compatible with PyTorch >= TORCH_TARGET_VERSION
# which is currently set to 2.11.
target_compile_definitions(_C_stable_libtorch PRIVATE
TORCH_TARGET_VERSION=0x020B000000000000ULL)
target_compile_definitions(_C_stable_libtorch PRIVATE USE_CUDA)
if(COOPERATIVE_TOPK_ARCHS)
target_compile_definitions(_C_stable_libtorch PRIVATE
@@ -1114,6 +1108,12 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
target_compile_definitions(_C_stable_libtorch PRIVATE
CUTLASS_ENABLE_DIRECT_CUDA_DRIVER_CALL=1)
elseif(VLLM_GPU_LANG STREQUAL "HIP")
# Set TORCH_TARGET_VERSION for stable ABI compatibility.
# This ensures we only use C-shim APIs available in PyTorch 2.10.
# _C_stable_libtorch is abi compatible with PyTorch >= TORCH_TARGET_VERSION
# which is currently set to 2.10.
target_compile_definitions(_C_stable_libtorch PRIVATE
TORCH_TARGET_VERSION=0x020A000000000000ULL)
target_compile_definitions(_C_stable_libtorch PRIVATE USE_ROCM)
endif()
@@ -1321,20 +1321,25 @@ define_extension_target(
USE_SABI 3
WITH_SOABI)
# Set TORCH_TARGET_VERSION for stable ABI compatibility.
# This ensures we only use C-shim APIs available in PyTorch 2.11.
# _moe_C_stable_libtorch is abi compatible with PyTorch >= TORCH_TARGET_VERSION
# which is currently set to 2.11.
target_compile_definitions(_moe_C_stable_libtorch PRIVATE
TORCH_TARGET_VERSION=0x020B000000000000ULL)
# Needed to use cuda/hip APIs from C-shim
if(VLLM_GPU_LANG STREQUAL "CUDA")
# Set TORCH_TARGET_VERSION for stable ABI compatibility.
# This ensures we only use C-shim APIs available in PyTorch 2.11.
# _moe_C_stable_libtorch is abi compatible with PyTorch >= TORCH_TARGET_VERSION
# which is currently set to 2.11.
target_compile_definitions(_moe_C_stable_libtorch PRIVATE
TORCH_TARGET_VERSION=0x020B000000000000ULL)
target_compile_definitions(_moe_C_stable_libtorch PRIVATE USE_CUDA)
# Needed by CUTLASS kernels
target_compile_definitions(_moe_C_stable_libtorch PRIVATE
CUTLASS_ENABLE_DIRECT_CUDA_DRIVER_CALL=1)
elseif(VLLM_GPU_LANG STREQUAL "HIP")
# Set TORCH_TARGET_VERSION for stable ABI compatibility.
# This ensures we only use C-shim APIs available in PyTorch 2.10.
# _moe_C_stable_libtorch is abi compatible with PyTorch >= TORCH_TARGET_VERSION
# which is currently set to 2.10.
target_compile_definitions(_moe_C_stable_libtorch PRIVATE
TORCH_TARGET_VERSION=0x020A000000000000ULL)
target_compile_definitions(_moe_C_stable_libtorch PRIVATE USE_ROCM)
endif()
@@ -1364,7 +1369,6 @@ if(VLLM_GPU_LANG STREQUAL "HIP")
set(VLLM_ROCM_EXT_SRC
"csrc/rocm/torch_bindings.cpp"
"csrc/rocm/skinny_gemms.cu"
"csrc/rocm/skinny_gemms_int4.cu"
"csrc/rocm/attention.cu")
set(VLLM_ROCM_HAS_GFX1100 OFF)
@@ -1408,7 +1412,6 @@ if (VLLM_GPU_LANG STREQUAL "CUDA")
include(cmake/external_projects/fmha_sm100.cmake)
include(cmake/external_projects/flashmla.cmake)
include(cmake/external_projects/qutlass.cmake)
include(cmake/external_projects/tml_fa4.cmake)
# vllm-flash-attn should be last as it overwrites some CMake functions
include(cmake/external_projects/vllm_flash_attn.cmake)
+1 -1
View File
@@ -48,7 +48,7 @@ vLLM is flexible and easy to use with:
- Tool calling and reasoning parsers
- OpenAI-compatible API server, plus Anthropic Messages API and gRPC support
- Efficient multi-LoRA support for dense and MoE layers
- Support for NVIDIA GPUs, AMD GPUs, Intel GPUs, and x86/ARM/PowerPC CPUs. Additionally, diverse hardware plugins such as Google TPUs, Intel Gaudi, IBM Spyre, Huawei Ascend, Rebellions NPU, Apple Silicon, MetaX GPU, and more.
- Support for NVIDIA GPUs, AMD GPUs, and x86/ARM/PowerPC CPUs. Additionally, diverse hardware plugins such as Google TPUs, Intel Gaudi, IBM Spyre, Huawei Ascend, Rebellions NPU, Apple Silicon, MetaX GPU, and more.
vLLM seamlessly supports 200+ model architectures on Hugging Face, including:
+3 -175
View File
@@ -75,11 +75,7 @@ def run_mla_benchmark(config: BenchmarkConfig, **kwargs) -> BenchmarkResult:
from mla_runner import run_mla_benchmark as run_mla
return run_mla(
config.backend,
config,
prefill_backend=config.prefill_backend,
sparse_mla_force_mqa=config.sparse_mla_force_mqa,
**kwargs,
config.backend, config, prefill_backend=config.prefill_backend, **kwargs
)
@@ -596,30 +592,6 @@ def main():
default="profile",
help="Output file name for ncu profile (default: 'profile').",
)
parser.add_argument(
"--torch-profile",
action="store_true",
default=False,
help="Collect a PyTorch profiler Chrome trace for each benchmark run.",
)
parser.add_argument(
"--torch-profile-dir",
default=None,
help="Directory for PyTorch profiler traces.",
)
parser.add_argument(
"--torch-profile-iters",
type=int,
default=3,
help="Number of forward passes to record per PyTorch profiler trace.",
)
parser.add_argument(
"--sparse-mla-mha-variants",
nargs="+",
default=None,
choices=["dense_mha", "mqa"],
help="Sparse MLA variants to run in mha_vs_mqa mode. Defaults to both.",
)
# Parameter sweep (use YAML config for advanced sweeps)
parser.add_argument(
@@ -669,7 +641,6 @@ def main():
# Prefill backends (e.g., ["fa3", "fa4"])
args.prefill_backends = yaml_config.get("prefill_backends", None)
args.prefill_backend = yaml_config.get("prefill_backend", None)
# FP8 output benchmark knobs; CLI wins.
if args.fp8_output_scale is None:
@@ -712,9 +683,6 @@ def main():
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)
args.max_model_len = model.get(
"max_model_len", getattr(args, "max_model_len", None)
)
# 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)
@@ -733,21 +701,6 @@ def main():
args.cuda_graphs = yaml_config["cuda_graphs"]
if "ncu_profile" in yaml_config:
args.ncu_profile = yaml_config["ncu_profile"]
if "torch_profile" in yaml_config:
args.torch_profile = yaml_config["torch_profile"]
if "torch_profile_dir" in yaml_config:
args.torch_profile_dir = yaml_config["torch_profile_dir"]
if "torch_profile_iters" in yaml_config:
args.torch_profile_iters = yaml_config["torch_profile_iters"]
args.sparse_mla_topk_pattern = yaml_config.get(
"sparse_mla_topk_pattern", "random"
)
args.sparse_mla_dense_mha_max_seq_len = yaml_config.get(
"sparse_mla_dense_mha_max_seq_len", None
)
args.sparse_mla_mha_variants = yaml_config.get(
"sparse_mla_mha_variants", args.sparse_mla_mha_variants
)
# Parameter sweep configuration
if "parameter_sweep" in yaml_config:
@@ -889,6 +842,8 @@ 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,
@@ -1108,133 +1063,6 @@ def main():
f"\n [yellow]Prefill always faster for batch_size={bs}[/]"
)
# Handle MHA vs MQA comparison mode for sparse MLA
elif hasattr(args, "mode") and args.mode == "mha_vs_mqa":
console.print("[yellow]Mode: MHA vs MQA comparison for sparse MLA[/]")
sparse_mla_topk_pattern = getattr(args, "sparse_mla_topk_pattern", "random")
dense_mha_max_seq_len = getattr(args, "sparse_mla_dense_mha_max_seq_len", None)
prefill_backend = getattr(args, "prefill_backend", None)
if prefill_backend:
console.print(f"Prefill backend: {prefill_backend}")
available_variants = [
("dense_mha", False, "dense"),
("mqa", True, "auto"),
]
requested_variants = getattr(args, "sparse_mla_mha_variants", None)
if requested_variants is not None:
valid_variants = {label for label, _, _ in available_variants}
invalid_variants = sorted(set(requested_variants) - valid_variants)
if invalid_variants:
raise ValueError(
"Invalid sparse_mla_mha_variants entries: "
f"{invalid_variants}. Valid variants are: "
f"{sorted(valid_variants)}"
)
requested_variant_set = set(requested_variants)
variants = [
variant
for variant in available_variants
if variant[0] in requested_variant_set
]
else:
variants = available_variants
formatter = ResultsFormatter(console)
total = 0
for spec in args.batch_specs:
q_len = max(request.q_len for request in parse_batch_spec(spec))
for variant_label, _, _ in variants:
if (
variant_label == "dense_mha"
and dense_mha_max_seq_len is not None
and q_len > dense_mha_max_seq_len
):
continue
total += len(backends)
with tqdm(total=total, desc="Benchmarking") as pbar:
for spec in args.batch_specs:
q_len = max(request.q_len for request in parse_batch_spec(spec))
for backend in backends:
for variant_label, force_mqa, mha_mode in variants:
if (
variant_label == "dense_mha"
and dense_mha_max_seq_len is not None
and q_len > dense_mha_max_seq_len
):
continue
config = BenchmarkConfig(
backend=f"{backend}_{variant_label}",
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,
max_model_len=getattr(args, "max_model_len", None),
kv_cache_dtype=args.kv_cache_dtype,
profile_memory=args.profile_memory,
use_cuda_graphs=args.cuda_graphs,
ncu_profile=args.ncu_profile,
torch_profile=args.torch_profile,
torch_profile_dir=args.torch_profile_dir,
torch_profile_iters=args.torch_profile_iters,
warmup_ms=args.warmup_ms,
kv_lora_rank=getattr(args, "kv_lora_rank", None),
qk_nope_head_dim=getattr(args, "qk_nope_head_dim", None),
qk_rope_head_dim=getattr(args, "qk_rope_head_dim", None),
v_head_dim=getattr(args, "v_head_dim", None),
sparse_mla_force_mqa=force_mqa,
sparse_mla_mha_mode=mha_mode,
sparse_mla_dense_mha_max_seq_len=dense_mha_max_seq_len,
sparse_mla_topk_pattern=sparse_mla_topk_pattern,
prefill_backend=prefill_backend,
)
# run_mla_benchmark needs the real backend name
from mla_runner import run_mla_benchmark as run_mla
run_label = f"{backend}_{variant_label} {spec}"
pbar.set_postfix_str(run_label)
try:
result = run_mla(
backend,
config,
prefill_backend=prefill_backend,
sparse_mla_force_mqa=force_mqa,
)
except Exception as e:
result = BenchmarkResult(
config=config,
mean_time=float("inf"),
median_time=float("inf"),
std_time=0,
min_time=float("inf"),
max_time=float("inf"),
error=str(e),
)
all_results.append(result)
if args.output_csv:
formatter.save_csv(all_results, args.output_csv)
if args.output_json:
formatter.save_json(all_results, args.output_json)
if not result.success:
console.print(
f"[red]Error {backend}_{variant_label} "
f"{spec}: {result.error}[/]"
)
pbar.update(1)
# Display results with variant labels as separate "backends"
console.print("\n[bold green]MHA vs MQA Results:[/]")
variant_backends = [f"{b}_{v}" for b in backends for v, _, _ in variants]
formatter.print_table(all_results, variant_backends)
# Handle model parameter sweep mode
elif hasattr(args, "model_parameter_sweep") and args.model_parameter_sweep:
# Model parameter sweep
+37 -68
View File
@@ -4,10 +4,8 @@
"""Common utilities for attention benchmarking."""
import csv
import gc
import json
import math
from collections.abc import Sequence
from dataclasses import asdict, dataclass
from pathlib import Path
from typing import Any
@@ -46,13 +44,10 @@ def run_do_bench(
kwargs: dict[str, Any] = {"return_mode": "all"}
if use_cuda_graphs:
result = triton.testing.do_bench_cudagraph(benchmark_fn, **kwargs)
gc.collect()
torch.accelerator.empty_cache()
else:
if warmup_ms is not None:
kwargs["warmup"] = warmup_ms
result = triton.testing.do_bench(benchmark_fn, **kwargs)
torch.accelerator.synchronize()
return result
@@ -96,6 +91,42 @@ except ImportError:
AttentionLayerBase = object # Fallback
class MockKVBProj:
"""Mock KV projection layer for MLA prefill mode.
Mimics ColumnParallelLinear behavior for kv_b_proj in MLA backends.
Projects kv_c_normed to [qk_nope_head_dim + v_head_dim] per head.
"""
def __init__(self, num_heads: int, qk_nope_head_dim: int, v_head_dim: int):
self.num_heads = num_heads
self.qk_nope_head_dim = qk_nope_head_dim
self.v_head_dim = v_head_dim
self.out_dim = qk_nope_head_dim + v_head_dim
self.weight = torch.empty(0, dtype=torch.bfloat16)
def __call__(self, x: torch.Tensor) -> tuple[torch.Tensor]:
"""
Project kv_c_normed to output space.
Args:
x: Input tensor [num_tokens, kv_lora_rank]
Returns:
Tuple containing output tensor
[num_tokens, num_heads, qk_nope_head_dim + v_head_dim]
"""
num_tokens = x.shape[0]
result = torch.randn(
num_tokens,
self.num_heads,
self.out_dim,
device=x.device,
dtype=x.dtype,
)
return (result,) # Return as tuple to match ColumnParallelLinear API
class MockIndexer:
"""Mock Indexer for sparse MLA backends.
@@ -127,60 +158,6 @@ class MockIndexer:
)
self.topk_indices_buffer[:num_tokens] = indices
def fill_indices(
self,
num_tokens: int,
max_kv_len: int,
pattern: str = "random",
requests: Sequence[Any] | None = None,
):
if pattern == "random":
self.fill_random_indices(num_tokens, max_kv_len)
return
if pattern == "prefix":
indices = torch.arange(
self.topk_tokens,
dtype=torch.int32,
device=self.topk_indices_buffer.device,
)
indices = (indices % max_kv_len).expand(num_tokens, -1)
self.topk_indices_buffer[:num_tokens] = indices
return
if pattern == "sliding_window":
if requests is None:
start = max(max_kv_len - self.topk_tokens, 0)
indices = torch.arange(
start,
start + self.topk_tokens,
dtype=torch.int32,
device=self.topk_indices_buffer.device,
)
indices = indices.clamp(max=max_kv_len - 1).expand(num_tokens, -1)
self.topk_indices_buffer[:num_tokens] = indices
return
rows = []
offsets = torch.arange(
self.topk_tokens,
dtype=torch.int32,
device=self.topk_indices_buffer.device,
) - (self.topk_tokens - 1)
for request in requests:
q_len = request.q_len
kv_len = request.kv_len
context_len = kv_len - q_len
positions = torch.arange(
context_len,
kv_len,
dtype=torch.int32,
device=self.topk_indices_buffer.device,
)
row_indices = positions[:, None] + offsets[None, :]
rows.append(row_indices.clamp(min=0, max=kv_len - 1))
self.topk_indices_buffer[:num_tokens] = torch.cat(rows, dim=0)
return
raise ValueError(f"Unknown sparse MLA topk pattern: {pattern}")
class MockLayer(AttentionLayerBase):
"""Mock attention layer with scale parameters and impl.
@@ -275,14 +252,10 @@ class BenchmarkConfig:
num_kv_heads: int
block_size: int
device: str
max_model_len: int | None = None
dtype: torch.dtype = torch.float16
profile_memory: bool = False
use_cuda_graphs: bool = True
use_cuda_graphs: bool = False
ncu_profile: bool = False
torch_profile: bool = False
torch_profile_dir: str | None = None
torch_profile_iters: int = 3
warmup_ms: int | None = None
# "auto" or "fp8"
@@ -298,10 +271,6 @@ class BenchmarkConfig:
# Backend-specific tuning
num_kv_splits: int | None = None # CUTLASS MLA
reorder_batch_threshold: int | None = None # FlashAttn MLA, FlashMLA
sparse_mla_force_mqa: bool = False # Force MQA path for sparse MLA
sparse_mla_mha_mode: str = "auto" # "auto" or "dense"
sparse_mla_dense_mha_max_seq_len: int | None = None
sparse_mla_topk_pattern: str = "random" # "random", "prefix", "sliding_window"
num_splits: int | None = None # FlashAttention split-K (0=auto, 1=disabled)
@@ -1,474 +0,0 @@
# Sparse MLA benchmark: forward_mha vs forward_mqa
#
# Usage:
# python benchmark.py --config configs/mla_sparse_mha_vs_mqa.yaml
#
# Heatmap grid:
# - batch_size: 1, 2, 4, 8, 16, 32
# - seq_len: 32, 64, 128, 256, 512, 1024, 2048
# - q_len: powers of two through seq_len
#
# Specs with q_len < seq_len include context; the q_len == seq_len diagonal
# covers pure prefill.
# The model shape below is the DP case. For the TP8 run, manually change
# model.num_q_heads from 128 to 16 before rerunning this benchmark.
mode: mha_vs_mqa
model:
name: "deepseek-v3"
num_layers: 60
num_q_heads: 128
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
max_model_len: 2048
batch_specs:
# Batch size 1
# seq_len = 32
- "1q1s32"
- "1q2s32"
- "1q4s32"
- "1q8s32"
- "1q16s32"
- "1q32"
# seq_len = 64
- "1q1s64"
- "1q2s64"
- "1q4s64"
- "1q8s64"
- "1q16s64"
- "1q32s64"
- "1q64"
# seq_len = 128
- "1q1s128"
- "1q2s128"
- "1q4s128"
- "1q8s128"
- "1q16s128"
- "1q32s128"
- "1q64s128"
- "1q128"
# seq_len = 256
- "1q1s256"
- "1q2s256"
- "1q4s256"
- "1q8s256"
- "1q16s256"
- "1q32s256"
- "1q64s256"
- "1q128s256"
- "1q256"
# seq_len = 512
- "1q1s512"
- "1q2s512"
- "1q4s512"
- "1q8s512"
- "1q16s512"
- "1q32s512"
- "1q64s512"
- "1q128s512"
- "1q256s512"
- "1q512"
# seq_len = 1024
- "1q1s1024"
- "1q2s1024"
- "1q4s1024"
- "1q8s1024"
- "1q16s1024"
- "1q32s1024"
- "1q64s1024"
- "1q128s1024"
- "1q256s1024"
- "1q512s1024"
- "1q1024"
# seq_len = 2048
- "1q1s2048"
- "1q2s2048"
- "1q4s2048"
- "1q8s2048"
- "1q16s2048"
- "1q32s2048"
- "1q64s2048"
- "1q128s2048"
- "1q256s2048"
- "1q512s2048"
- "1q1024s2048"
- "1q2048"
# Batch size 2
# seq_len = 32
- "2q1s32"
- "2q2s32"
- "2q4s32"
- "2q8s32"
- "2q16s32"
- "2q32"
# seq_len = 64
- "2q1s64"
- "2q2s64"
- "2q4s64"
- "2q8s64"
- "2q16s64"
- "2q32s64"
- "2q64"
# seq_len = 128
- "2q1s128"
- "2q2s128"
- "2q4s128"
- "2q8s128"
- "2q16s128"
- "2q32s128"
- "2q64s128"
- "2q128"
# seq_len = 256
- "2q1s256"
- "2q2s256"
- "2q4s256"
- "2q8s256"
- "2q16s256"
- "2q32s256"
- "2q64s256"
- "2q128s256"
- "2q256"
# seq_len = 512
- "2q1s512"
- "2q2s512"
- "2q4s512"
- "2q8s512"
- "2q16s512"
- "2q32s512"
- "2q64s512"
- "2q128s512"
- "2q256s512"
- "2q512"
# seq_len = 1024
- "2q1s1024"
- "2q2s1024"
- "2q4s1024"
- "2q8s1024"
- "2q16s1024"
- "2q32s1024"
- "2q64s1024"
- "2q128s1024"
- "2q256s1024"
- "2q512s1024"
- "2q1024"
# seq_len = 2048
- "2q1s2048"
- "2q2s2048"
- "2q4s2048"
- "2q8s2048"
- "2q16s2048"
- "2q32s2048"
- "2q64s2048"
- "2q128s2048"
- "2q256s2048"
- "2q512s2048"
- "2q1024s2048"
- "2q2048"
# Batch size 4
# seq_len = 32
- "4q1s32"
- "4q2s32"
- "4q4s32"
- "4q8s32"
- "4q16s32"
- "4q32"
# seq_len = 64
- "4q1s64"
- "4q2s64"
- "4q4s64"
- "4q8s64"
- "4q16s64"
- "4q32s64"
- "4q64"
# seq_len = 128
- "4q1s128"
- "4q2s128"
- "4q4s128"
- "4q8s128"
- "4q16s128"
- "4q32s128"
- "4q64s128"
- "4q128"
# seq_len = 256
- "4q1s256"
- "4q2s256"
- "4q4s256"
- "4q8s256"
- "4q16s256"
- "4q32s256"
- "4q64s256"
- "4q128s256"
- "4q256"
# seq_len = 512
- "4q1s512"
- "4q2s512"
- "4q4s512"
- "4q8s512"
- "4q16s512"
- "4q32s512"
- "4q64s512"
- "4q128s512"
- "4q256s512"
- "4q512"
# seq_len = 1024
- "4q1s1024"
- "4q2s1024"
- "4q4s1024"
- "4q8s1024"
- "4q16s1024"
- "4q32s1024"
- "4q64s1024"
- "4q128s1024"
- "4q256s1024"
- "4q512s1024"
- "4q1024"
# seq_len = 2048
- "4q1s2048"
- "4q2s2048"
- "4q4s2048"
- "4q8s2048"
- "4q16s2048"
- "4q32s2048"
- "4q64s2048"
- "4q128s2048"
- "4q256s2048"
- "4q512s2048"
- "4q1024s2048"
- "4q2048"
# Batch size 8
# seq_len = 32
- "8q1s32"
- "8q2s32"
- "8q4s32"
- "8q8s32"
- "8q16s32"
- "8q32"
# seq_len = 64
- "8q1s64"
- "8q2s64"
- "8q4s64"
- "8q8s64"
- "8q16s64"
- "8q32s64"
- "8q64"
# seq_len = 128
- "8q1s128"
- "8q2s128"
- "8q4s128"
- "8q8s128"
- "8q16s128"
- "8q32s128"
- "8q64s128"
- "8q128"
# seq_len = 256
- "8q1s256"
- "8q2s256"
- "8q4s256"
- "8q8s256"
- "8q16s256"
- "8q32s256"
- "8q64s256"
- "8q128s256"
- "8q256"
# seq_len = 512
- "8q1s512"
- "8q2s512"
- "8q4s512"
- "8q8s512"
- "8q16s512"
- "8q32s512"
- "8q64s512"
- "8q128s512"
- "8q256s512"
- "8q512"
# seq_len = 1024
- "8q1s1024"
- "8q2s1024"
- "8q4s1024"
- "8q8s1024"
- "8q16s1024"
- "8q32s1024"
- "8q64s1024"
- "8q128s1024"
- "8q256s1024"
- "8q512s1024"
- "8q1024"
# seq_len = 2048
- "8q1s2048"
- "8q2s2048"
- "8q4s2048"
- "8q8s2048"
- "8q16s2048"
- "8q32s2048"
- "8q64s2048"
- "8q128s2048"
- "8q256s2048"
- "8q512s2048"
- "8q1024s2048"
- "8q2048"
# Batch size 16
# seq_len = 32
- "16q1s32"
- "16q2s32"
- "16q4s32"
- "16q8s32"
- "16q16s32"
- "16q32"
# seq_len = 64
- "16q1s64"
- "16q2s64"
- "16q4s64"
- "16q8s64"
- "16q16s64"
- "16q32s64"
- "16q64"
# seq_len = 128
- "16q1s128"
- "16q2s128"
- "16q4s128"
- "16q8s128"
- "16q16s128"
- "16q32s128"
- "16q64s128"
- "16q128"
# seq_len = 256
- "16q1s256"
- "16q2s256"
- "16q4s256"
- "16q8s256"
- "16q16s256"
- "16q32s256"
- "16q64s256"
- "16q128s256"
- "16q256"
# seq_len = 512
- "16q1s512"
- "16q2s512"
- "16q4s512"
- "16q8s512"
- "16q16s512"
- "16q32s512"
- "16q64s512"
- "16q128s512"
- "16q256s512"
- "16q512"
# seq_len = 1024
- "16q1s1024"
- "16q2s1024"
- "16q4s1024"
- "16q8s1024"
- "16q16s1024"
- "16q32s1024"
- "16q64s1024"
- "16q128s1024"
- "16q256s1024"
- "16q512s1024"
- "16q1024"
# seq_len = 2048
- "16q1s2048"
- "16q2s2048"
- "16q4s2048"
- "16q8s2048"
- "16q16s2048"
- "16q32s2048"
- "16q64s2048"
- "16q128s2048"
- "16q256s2048"
- "16q512s2048"
- "16q1024s2048"
- "16q2048"
# Batch size 32
# seq_len = 32
- "32q1s32"
- "32q2s32"
- "32q4s32"
- "32q8s32"
- "32q16s32"
- "32q32"
# seq_len = 64
- "32q1s64"
- "32q2s64"
- "32q4s64"
- "32q8s64"
- "32q16s64"
- "32q32s64"
- "32q64"
# seq_len = 128
- "32q1s128"
- "32q2s128"
- "32q4s128"
- "32q8s128"
- "32q16s128"
- "32q32s128"
- "32q64s128"
- "32q128"
# seq_len = 256
- "32q1s256"
- "32q2s256"
- "32q4s256"
- "32q8s256"
- "32q16s256"
- "32q32s256"
- "32q64s256"
- "32q128s256"
- "32q256"
# seq_len = 512
- "32q1s512"
- "32q2s512"
- "32q4s512"
- "32q8s512"
- "32q16s512"
- "32q32s512"
- "32q64s512"
- "32q128s512"
- "32q256s512"
- "32q512"
# seq_len = 1024
- "32q1s1024"
- "32q2s1024"
- "32q4s1024"
- "32q8s1024"
- "32q16s1024"
- "32q32s1024"
- "32q64s1024"
- "32q128s1024"
- "32q256s1024"
- "32q512s1024"
- "32q1024"
# seq_len = 2048
- "32q1s2048"
- "32q2s2048"
- "32q4s2048"
- "32q8s2048"
- "32q16s2048"
- "32q32s2048"
- "32q64s2048"
- "32q128s2048"
- "32q256s2048"
- "32q512s2048"
- "32q1024s2048"
- "32q2048"
backends:
- FLASHMLA_SPARSE
device: "cuda:0"
profile_memory: false
sparse_mla_dense_mha_max_seq_len: 2048
sparse_mla_topk_pattern: "random"
output:
csv: "benchmark_output/mla_sparse_mha_vs_mqa.csv"
json: "benchmark_output/mla_sparse_mha_vs_mqa.json"
+17 -171
View File
@@ -9,8 +9,6 @@ needing full VllmConfig integration.
"""
import statistics
import tempfile
from pathlib import Path
import numpy as np
import torch
@@ -19,6 +17,7 @@ from common import (
BenchmarkResult,
MockHfConfig,
MockIndexer,
MockKVBProj,
MockLayer,
run_do_bench,
run_ncu_profile,
@@ -34,59 +33,8 @@ from vllm.config import (
VllmConfig,
set_current_vllm_config,
)
from vllm.model_executor.layers.linear import ColumnParallelLinear
from vllm.v1.attention.backends.mla.prefill.registry import MLAPrefillBackendEnum
def _safe_profile_name(value: str) -> str:
return "".join(c if c.isalnum() or c in "._-" else "_" for c in value)
def _create_kv_b_proj(
mla_dims: dict,
device: torch.device,
):
kv_b_proj = ColumnParallelLinear(
mla_dims["kv_lora_rank"],
mla_dims["num_q_heads"]
* (mla_dims["qk_nope_head_dim"] + mla_dims["v_head_dim"]),
bias=False,
params_dtype=torch.bfloat16,
quant_config=None,
prefix="benchmark.kv_b_proj",
).to(device)
with torch.no_grad():
kv_b_proj.weight.copy_(torch.randn_like(kv_b_proj.weight))
return kv_b_proj
def _ensure_single_rank_model_parallel() -> None:
import torch.distributed as dist
from vllm.distributed import (
ensure_model_parallel_initialized,
init_distributed_environment,
model_parallel_is_initialized,
)
if not dist.is_available():
return
if not dist.is_initialized():
with tempfile.NamedTemporaryFile(
prefix="vllm_bench_dist_", delete=False
) as init_file:
distributed_init_method = f"file://{init_file.name}"
init_distributed_environment(
world_size=1,
rank=0,
distributed_init_method=distributed_init_method,
local_rank=0,
backend="nccl",
)
if not model_parallel_is_initialized():
ensure_model_parallel_initialized(1, 1)
# ============================================================================
# VllmConfig Creation
# ============================================================================
@@ -118,12 +66,10 @@ def create_minimal_vllm_config(
block_size: int = 128,
max_num_seqs: int = 256,
max_num_batched_tokens: int = 8192,
max_model_len: int = 32768,
mla_dims: dict | None = None,
index_topk: int | None = None,
prefill_backend: str | None = None,
kv_cache_dtype: str = "auto",
sparse_mla_force_mqa: bool = False,
) -> VllmConfig:
"""
Create minimal VllmConfig for MLA benchmarks.
@@ -140,8 +86,6 @@ def create_minimal_vllm_config(
prefill_backend: Prefill backend name (e.g., "fa3", "fa4", "flashinfer",
"trtllm"). Configures the attention config to force
the specified prefill backend.
sparse_mla_force_mqa: If True, forces all sparse MLA tokens through
forward_mqa (even prefill tokens).
Returns:
VllmConfig for benchmarking
@@ -187,7 +131,7 @@ def create_minimal_vllm_config(
trust_remote_code=True,
dtype="bfloat16",
seed=0,
max_model_len=max_model_len,
max_model_len=32768,
quantization=None,
enforce_eager=False,
max_logprobs=20,
@@ -219,7 +163,7 @@ def create_minimal_vllm_config(
scheduler_config = SchedulerConfig(
max_num_seqs=max_num_seqs,
max_num_batched_tokens=max(max_num_batched_tokens, max_num_seqs),
max_model_len=max_model_len,
max_model_len=32768,
is_encoder_decoder=False,
enable_chunked_prefill=True,
)
@@ -248,9 +192,6 @@ def create_minimal_vllm_config(
"flash_attn_version"
]
if sparse_mla_force_mqa:
vllm_config.attention_config.sparse_mla_force_mqa = True
return vllm_config
@@ -607,7 +548,12 @@ def _create_backend_impl(
# Calculate scale
scale = 1.0 / np.sqrt(mla_dims["qk_nope_head_dim"] + mla_dims["qk_rope_head_dim"])
kv_b_proj = _create_kv_b_proj(mla_dims, device)
# Create mock kv_b_proj layer for prefill mode
mock_kv_b_proj = MockKVBProj(
num_heads=mla_dims["num_q_heads"],
qk_nope_head_dim=mla_dims["qk_nope_head_dim"],
v_head_dim=mla_dims["v_head_dim"],
)
# Create indexer for sparse backends
indexer = None
@@ -638,7 +584,7 @@ def _create_backend_impl(
"qk_rope_head_dim": mla_dims["qk_rope_head_dim"],
"qk_head_dim": mla_dims["qk_nope_head_dim"] + mla_dims["qk_rope_head_dim"],
"v_head_dim": mla_dims["v_head_dim"],
"kv_b_proj": kv_b_proj,
"kv_b_proj": mock_kv_b_proj,
}
# Add indexer for sparse backends
@@ -839,35 +785,14 @@ def _run_single_benchmark(
# Fill indexer with random indices for sparse backends
is_sparse = backend_cfg.get("is_sparse", False)
if is_sparse and indexer is not None:
indexer.fill_indices(
total_q,
max_kv_len,
getattr(config, "sparse_mla_topk_pattern", "random"),
)
indexer.fill_random_indices(total_q, max_kv_len)
# Determine which forward methods to use based on metadata.
# Non-sparse backends use .decode/.prefill sub-objects.
# Sparse backends use num_decode_tokens/num_prefills directly.
#
# sparse_mla_force_mqa overrides: even for prefill metadata, use MQA.
force_mqa = getattr(config, "sparse_mla_force_mqa", False)
force_dense_mha = getattr(config, "sparse_mla_mha_mode", "auto") == "dense"
if force_mqa:
has_decode = True
has_prefill = False
elif is_sparse:
has_decode = metadata.num_decode_tokens > 0
has_prefill = metadata.num_prefills > 0
else:
has_decode = metadata.decode is not None
has_prefill = metadata.prefill is not None
# Sparse MLA backends always use forward_mqa
has_decode = is_sparse or getattr(metadata, "decode", None) is not None
has_prefill = not is_sparse and getattr(metadata, "prefill", None) is not None
if not has_decode and not has_prefill:
raise RuntimeError("Metadata has neither decode nor prefill metadata")
if is_sparse and force_dense_mha and not has_prefill:
raise RuntimeError(
"Sparse MLA dense_mha benchmark did not produce prefill metadata. "
"Check reorder_batch_threshold/path forcing."
)
num_decode = (
metadata.num_decode_tokens
@@ -946,6 +871,7 @@ def _run_single_benchmark(
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
@@ -972,48 +898,6 @@ def _run_single_benchmark(
throughput_tokens_per_sec=0.0,
)
if config.torch_profile:
profile_dir = Path(
config.torch_profile_dir or "benchmark_outputs/torch_profiles"
)
profile_dir.mkdir(parents=True, exist_ok=True)
trace_name = _safe_profile_name(f"{config.backend}_{config.batch_spec}")
trace_path = profile_dir / f"{trace_name}.json"
iters = max(config.torch_profile_iters, 1)
forward_fn()
torch.accelerator.synchronize()
with torch.profiler.profile(
activities=[
torch.profiler.ProfilerActivity.CPU,
torch.profiler.ProfilerActivity.CUDA,
],
record_shapes=True,
profile_memory=True,
with_stack=False,
) as prof:
for _ in range(iters):
forward_fn()
torch.accelerator.synchronize()
prof.step()
prof.export_chrome_trace(str(trace_path))
print(f"Saved PyTorch profiler trace to {trace_path}")
print(
prof.key_averages().table(
sort_by="cuda_time_total",
row_limit=25,
)
)
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,
)
all_ms = run_do_bench(benchmark_fn, config.use_cuda_graphs, config.warmup_ms)
# Convert ms to seconds per layer
@@ -1036,7 +920,6 @@ def _run_mla_benchmark_batched(
configs_with_params: list[tuple], # [(config, threshold, num_splits), ...]
index_topk: int = 2048,
prefill_backend: str | None = None,
sparse_mla_force_mqa: bool = False,
output_scale: float | None = None,
fuse_quant_op: bool = False,
) -> list[BenchmarkResult]:
@@ -1057,8 +940,6 @@ def _run_mla_benchmark_batched(
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.
sparse_mla_force_mqa: If True, forces all sparse MLA tokens through
forward_mqa (even prefill tokens).
Returns:
List of BenchmarkResult objects
@@ -1099,41 +980,21 @@ def _run_mla_benchmark_batched(
sum(r.q_len for r in parse_batch_spec(cfg.batch_spec))
for cfg, *_ in configs_with_params
)
max_model_len = max(
max_total_q,
max(
getattr(cfg, "max_model_len", None) or 32768
for cfg, *_ in configs_with_params
),
)
# Create and set vLLM config for MLA (reused across all benchmarks)
vllm_config = create_minimal_vllm_config(
model_name="deepseek-v3", # Used only for model path
block_size=block_size,
max_num_batched_tokens=max_total_q,
max_model_len=max_model_len,
mla_dims=mla_dims, # Use custom dims from config or default
index_topk=index_topk if is_sparse else None,
prefill_backend=prefill_backend,
kv_cache_dtype=kv_cache_dtype,
sparse_mla_force_mqa=sparse_mla_force_mqa,
)
results = []
# Initialize workspace manager (needed by metadata builders)
from vllm.v1.worker.workspace import (
init_workspace_manager,
is_workspace_manager_initialized,
)
if not is_workspace_manager_initialized():
init_workspace_manager(device)
with set_current_vllm_config(vllm_config):
_ensure_single_rank_model_parallel()
# Create backend impl, layer, builder, and indexer (reused across benchmarks)
impl, layer, builder_instance, indexer = _create_backend_impl(
backend_cfg,
@@ -1179,20 +1040,9 @@ def _run_mla_benchmark_batched(
for config, threshold, num_splits in configs_with_params:
# Set threshold for this benchmark (FlashAttn/FlashMLA only)
original_threshold = None
effective_threshold = threshold
force_dense_mha = (
is_sparse
and getattr(config, "sparse_mla_mha_mode", "auto") == "dense"
and not getattr(config, "sparse_mla_force_mqa", False)
)
if force_dense_mha:
# Sparse MLA normally treats q_len <= 1 as decode. Use an
# impossible threshold so dense_mha benchmarks actually run
# the prefill/MHA path, including q_len=1 short extends.
effective_threshold = -1
if effective_threshold is not None and builder_instance:
if threshold is not None and builder_instance:
original_threshold = builder_instance.reorder_batch_threshold
builder_instance.reorder_batch_threshold = effective_threshold
builder_instance.reorder_batch_threshold = threshold
# Set num_splits for CUTLASS
original_num_splits = None
@@ -1240,7 +1090,6 @@ def run_mla_benchmark(
num_kv_splits: int | None = None,
index_topk: int = 2048,
prefill_backend: str | None = None,
sparse_mla_force_mqa: bool = False,
output_scale: float | None = None,
fuse_quant_op: bool = False,
) -> BenchmarkResult | list[BenchmarkResult]:
@@ -1262,8 +1111,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.
sparse_mla_force_mqa: If True, forces all sparse MLA tokens through
forward_mqa (even prefill tokens).
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.
@@ -1295,7 +1142,6 @@ def run_mla_benchmark(
configs_with_params,
index_topk,
prefill_backend=prefill_backend,
sparse_mla_force_mqa=sparse_mla_force_mqa,
output_scale=output_scale,
fuse_quant_op=fuse_quant_op,
)
+2 -2
View File
@@ -12,7 +12,7 @@ from dataclasses import dataclass, field
import aiohttp
import huggingface_hub.constants
from tqdm.asyncio import tqdm
from transformers import AutoTokenizer, PythonBackend, TokenizersBackend
from transformers import AutoTokenizer, PreTrainedTokenizer, PreTrainedTokenizerFast
# NOTE(simon): do not import vLLM here so the benchmark script
# can run without vLLM installed.
@@ -609,7 +609,7 @@ def get_tokenizer(
tokenizer_mode: str = "auto",
trust_remote_code: bool = False,
**kwargs,
) -> PythonBackend | TokenizersBackend:
) -> PreTrainedTokenizer | PreTrainedTokenizerFast:
if pretrained_model_name_or_path is not None and not os.path.exists(
pretrained_model_name_or_path
):
+2 -4
View File
@@ -132,10 +132,8 @@ def benchmark_function(
reset_memory_stats()
# Benchmark
start_events = [
torch.cuda.Event(enable_timing=True) for _ in range(benchmark_iters)
]
end_events = [torch.cuda.Event(enable_timing=True) for _ in range(benchmark_iters)]
start_events = [torch.Event(enable_timing=True) for _ in range(benchmark_iters)]
end_events = [torch.Event(enable_timing=True) for _ in range(benchmark_iters)]
for i in range(benchmark_iters):
logits_copy = logits.clone()
+4 -3
View File
@@ -69,11 +69,12 @@ def make_inputs(total_tokens, num_reqs, block_size):
# Output workspace
dst = torch.zeros(total_tokens, HEAD_DIM, dtype=torch.bfloat16, device="cuda")
seq_lens_t = torch.tensor(seq_lens, dtype=torch.int32, device="cuda")
workspace_starts_t = torch.tensor(
workspace_starts, dtype=torch.int32, device="cuda"
)
return cache, dst, block_table, workspace_starts_t
return cache, dst, block_table, seq_lens_t, workspace_starts_t
def bench_scenario(label, num_reqs, total_tokens_list, save_path):
@@ -93,7 +94,7 @@ def bench_scenario(label, num_reqs, total_tokens_list, save_path):
)
)
def bench_fn(total_tokens, provider, num_reqs):
cache, dst, block_table, ws_starts = make_inputs(
cache, dst, block_table, seq_lens_t, ws_starts = make_inputs(
total_tokens, num_reqs, BLOCK_SIZE
)
@@ -101,7 +102,7 @@ def bench_scenario(label, num_reqs, total_tokens_list, save_path):
ms, min_ms, max_ms = triton.testing.do_bench_cudagraph(
lambda: ops.cp_gather_and_upconvert_fp8_kv_cache(
cache, dst, block_table, ws_starts, num_reqs
cache, dst, block_table, seq_lens_t, ws_starts, num_reqs
),
quantiles=quantiles,
rep=500,
@@ -17,7 +17,7 @@ from vllm.model_executor.layers.fused_moe.fused_flydsl_moe import fused_flydsl_m
from vllm.model_executor.layers.quantization.compressed_tensors.compressed_tensors_moe import ( # noqa: E501
compressed_tensors_moe_w4a16_flydsl,
)
from vllm.utils.platform_utils import get_device_name_as_file_name
from vllm.platforms import current_platform
RoutingBuffers = tuple[
torch.Tensor, # sorted_token_ids
@@ -259,7 +259,7 @@ def tune_flydsl_moe_w4a16(
)
us_best = us
tuned_config[str(num_tokens)] = tile_config
device_name = get_device_name_as_file_name()
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"
+2 -2
View File
@@ -134,8 +134,8 @@ def benchmark_config(
torch.accelerator.synchronize()
# Benchmark
start = torch.cuda.Event(enable_timing=True)
end = torch.cuda.Event(enable_timing=True)
start = torch.Event(enable_timing=True)
end = torch.Event(enable_timing=True)
start.record()
for _ in range(num_iters):
with override_config(config):
+93 -38
View File
@@ -19,11 +19,13 @@ from vllm.utils.torch_utils import (
logger = init_logger(__name__)
NUM_BLOCKS = 128 * 1024
PARTITION_SIZE = 512
PARTITION_SIZE_ROCM = 256
@torch.inference_mode()
def main(
version: str,
num_seqs: int,
seq_len: int,
num_query_heads: int,
@@ -80,20 +82,27 @@ def main(
# Prepare for the paged attention kernel.
output = torch.empty_like(query)
num_partitions = (max_seq_len + PARTITION_SIZE_ROCM - 1) // PARTITION_SIZE_ROCM
tmp_output = torch.empty(
size=(num_seqs, num_query_heads, num_partitions, head_size),
dtype=output.dtype,
device=output.device,
)
exp_sums = torch.empty(
size=(num_seqs, num_query_heads, num_partitions),
dtype=torch.float32,
device=output.device,
)
max_logits = torch.empty_like(exp_sums)
if version == "v2":
if current_platform.is_rocm():
global PARTITION_SIZE
if not args.custom_paged_attn and not current_platform.is_navi():
PARTITION_SIZE = 1024
else:
PARTITION_SIZE = PARTITION_SIZE_ROCM
num_partitions = (max_seq_len + PARTITION_SIZE - 1) // PARTITION_SIZE
tmp_output = torch.empty(
size=(num_seqs, num_query_heads, num_partitions, head_size),
dtype=output.dtype,
device=output.device,
)
exp_sums = torch.empty(
size=(num_seqs, num_query_heads, num_partitions),
dtype=torch.float32,
device=output.device,
)
max_logits = torch.empty_like(exp_sums)
def run_benchmark(num_iters: int, profile: bool = False) -> float:
def run_cuda_benchmark(num_iters: int, profile: bool = False) -> float:
torch.accelerator.synchronize()
if profile:
torch.cuda.cudart().cudaProfilerStart()
@@ -103,26 +112,67 @@ def main(
k_scale = v_scale = torch.tensor(1.0, dtype=torch.float32, device=device)
for _ in range(num_iters):
ops.paged_attention_rocm(
output,
exp_sums,
max_logits,
tmp_output,
query,
key_cache,
value_cache,
num_kv_heads,
scale,
block_tables,
seq_lens,
None,
block_size,
max_seq_len,
alibi_slopes,
kv_cache_dtype,
k_scale,
v_scale,
)
if version == "v1":
ops.paged_attention_v1(
output,
query,
key_cache,
value_cache,
num_kv_heads,
scale,
block_tables,
seq_lens,
block_size,
max_seq_len,
alibi_slopes,
kv_cache_dtype,
k_scale,
v_scale,
)
elif version == "v2":
if not args.custom_paged_attn:
ops.paged_attention_v2(
output,
exp_sums,
max_logits,
tmp_output,
query,
key_cache,
value_cache,
num_kv_heads,
scale,
block_tables,
seq_lens,
block_size,
max_seq_len,
alibi_slopes,
kv_cache_dtype,
k_scale,
v_scale,
)
else:
ops.paged_attention_rocm(
output,
exp_sums,
max_logits,
tmp_output,
query,
key_cache,
value_cache,
num_kv_heads,
scale,
block_tables,
seq_lens,
None,
block_size,
max_seq_len,
alibi_slopes,
kv_cache_dtype,
k_scale,
v_scale,
)
else:
raise ValueError(f"Invalid version: {version}")
torch.accelerator.synchronize()
end_time = time.perf_counter()
@@ -132,6 +182,7 @@ def main(
# Warmup.
print("Warming up...")
run_benchmark = run_cuda_benchmark
run_benchmark(num_iters=3, profile=False)
# Benchmark.
@@ -144,13 +195,12 @@ def main(
if __name__ == "__main__":
logger.warning(
"This script benchmarks the ROCm paged attention kernel. "
"This script benchmarks the paged attention kernel. "
"By default this is no longer used in vLLM inference."
)
if not current_platform.is_rocm():
raise RuntimeError("This benchmark requires the ROCm platform.")
parser = FlexibleArgumentParser(description="Benchmark the paged attention kernel.")
parser.add_argument("--version", type=str, choices=["v1", "v2"], default="v2")
parser.add_argument("--batch-size", type=int, default=8)
parser.add_argument("--seq-len", type=int, default=4096)
parser.add_argument("--num-query-heads", type=int, default=64)
@@ -158,7 +208,7 @@ if __name__ == "__main__":
parser.add_argument(
"--head-size",
type=int,
choices=[64, 128],
choices=[64, 80, 96, 112, 120, 128, 192, 256],
default=128,
)
parser.add_argument("--block-size", type=int, choices=[16, 32], default=16)
@@ -174,7 +224,11 @@ if __name__ == "__main__":
choices=["auto", "fp8", "fp8_e5m2", "fp8_e4m3"],
default="auto",
help="Data type for kv cache storage. If 'auto', will use model "
"data type. ROCm (AMD GPU) supports fp8 (=fp8_e4m3)",
"data type. CUDA 11.8+ supports fp8 (=fp8_e4m3) and fp8_e5m2. "
"ROCm (AMD GPU) supports fp8 (=fp8_e4m3)",
)
parser.add_argument(
"--custom-paged-attn", action="store_true", help="Use custom paged attention"
)
args = parser.parse_args()
print(args)
@@ -182,6 +236,7 @@ if __name__ == "__main__":
if args.num_query_heads % args.num_kv_heads != 0:
raise ValueError("num_query_heads must be divisible by num_kv_heads")
main(
version=args.version,
num_seqs=args.batch_size,
seq_len=args.seq_len,
num_query_heads=args.num_query_heads,
@@ -1,201 +0,0 @@
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
"""
Benchmark the RDNAHybridW4A16LinearKernel across decode and prefill shapes.
Usage:
python benchmark_int4_gemm.py
python benchmark_int4_gemm.py --models Qwen/Qwen3-4B
python benchmark_int4_gemm.py --group-size 128
"""
import argparse
import copy
import itertools
import os
import torch
from vllm.triton_utils import triton
# ---------------------------------------------------------------------------
# Weight shapes: [K, N], TP_SPLIT_DIM
# ---------------------------------------------------------------------------
WEIGHT_SHAPES = {
"Qwen/Qwen3-4B": [
([2560, 3840], 1), # qkv_proj
([2560, 2560], 0), # o_proj
([2560, 19456], 1), # gate_up_proj
([9728, 2560], 0), # down_proj
],
"Qwen/Qwen2.5-7B-Instruct": [
([3584, 4608], 1),
([3584, 3584], 0),
([3584, 37888], 1),
([18944, 3584], 0),
],
"trymirai/SmolLM2-1.7B-Instruct-AWQ": [
([2048, 6144], 1), # qkv_proj
([2048, 2048], 0), # o_proj
([2048, 16384], 1), # gate_up_proj
([8192, 2048], 0), # down_proj
],
"RedHatAI/Qwen3-8B-quantized.w4a16": [
([4096, 6144], 1), # qkv_proj
([4096, 4096], 0), # o_proj
([4096, 24576], 1), # gate_up_proj
([12288, 4096], 0), # down_proj
],
}
# ---------------------------------------------------------------------------
# Weight packing
# ---------------------------------------------------------------------------
def prepare_hybrid_weights(K, N, group_size, device="cuda"):
"""Create random weights for benchmarking.
Returns (w_q_skinny, w_s_skinny, w_fp16, w_zp). The triton path derives
its int32 view from w_q_skinny, so no separate int32 buffer is returned.
"""
num_groups = K // group_size
# Random packed weights — actual values don't matter for throughput
w_q_skinny_i32 = torch.randint(
0, 2**31, (N, K // 8), dtype=torch.int32, device=device
)
w_q_skinny = w_q_skinny_i32.view(torch.int8).contiguous()
w_s_skinny = torch.randn(N, num_groups, dtype=torch.float16, device=device) * 0.01
# Raw per-group zero-points for asymmetric benchmarks
w_zp = torch.randint(0, 16, (N, num_groups), dtype=torch.int32, device=device).to(
torch.float16
)
# FP16 baseline for F.linear
w_fp16 = torch.randn(N, K, dtype=torch.float16, device=device) * 0.01
return w_q_skinny, w_s_skinny, w_fp16, w_zp
# ---------------------------------------------------------------------------
# Benchmark
# ---------------------------------------------------------------------------
PROVIDERS = ["torch-fp16", "hybrid-w4a16", "hybrid-w4a16-zp"]
@triton.testing.perf_report(
triton.testing.Benchmark(
x_names=["batch_size"],
x_vals=[1, 2, 4, 8, 16, 32, 64, 128, 256, 512, 1024, 2048, 4096],
x_log=False,
line_arg="provider",
line_vals=PROVIDERS,
line_names=PROVIDERS,
ylabel="TFLOP/s (larger is better)",
plot_name="FP16 vs Hybrid W4A16",
args={},
)
)
def benchmark(batch_size, provider, N, K, group_size, weights):
M = batch_size
device = "cuda"
dtype = torch.float16
a = torch.randn((M, K), device=device, dtype=dtype)
quantiles = [0.5, 0.2, 0.8]
if provider == "torch-fp16":
w_fp16 = weights["w_fp16"]
ms, min_ms, max_ms = triton.testing.do_bench_cudagraph(
lambda: torch.nn.functional.linear(a, w_fp16),
quantiles=quantiles,
)
elif provider in ("hybrid-w4a16", "hybrid-w4a16-zp"):
from vllm.model_executor.kernels.linear.mixed_precision import (
rdna_hybrid_w4a16 as _k,
)
_rdna_hybrid_w4a16_apply_impl = _k._rdna_hybrid_w4a16_apply_impl
from vllm.utils.platform_utils import num_compute_units
w = weights
cu_count = num_compute_units()
use_zp = provider == "hybrid-w4a16-zp"
def run():
return _rdna_hybrid_w4a16_apply_impl(
a,
w["w_q_skinny"],
w["w_s_skinny"],
w["w_zp"] if use_zp else None,
None, # bias
cu_count,
group_size,
)
ms, min_ms, max_ms = triton.testing.do_bench_cudagraph(
run,
quantiles=quantiles,
)
else:
return 0.0, 0.0, 0.0
to_tflops = lambda t_ms: (2 * M * N * K) * 1e-12 / (t_ms * 1e-3)
return to_tflops(ms), to_tflops(max_ms), to_tflops(min_ms)
def prepare_shapes(args):
KN_model_names = []
for model, tp_size in itertools.product(args.models, args.tp_sizes):
for KN, tp_dim in copy.deepcopy(WEIGHT_SHAPES[model]):
KN[tp_dim] //= tp_size
KN.append(model)
KN_model_names.append(KN)
return KN_model_names
if __name__ == "__main__":
parser = argparse.ArgumentParser(
description="Benchmark RDNAHybridW4A16LinearKernel"
)
parser.add_argument(
"--models",
nargs="+",
type=str,
default=["Qwen/Qwen3-4B"],
choices=list(WEIGHT_SHAPES.keys()),
)
parser.add_argument("--tp-sizes", nargs="+", type=int, default=[1])
parser.add_argument("--group-size", type=int, default=128)
parser.add_argument("--save-path", type=str, default=None)
args = parser.parse_args()
for K, N, model in prepare_shapes(args):
group_size = args.group_size
print(f"\n{'=' * 70}")
print(f"{model}, N={N} K={K}, group_size={group_size}")
print(f"{'=' * 70}")
w_q_skinny, w_s_skinny, w_fp16, w_zp = prepare_hybrid_weights(K, N, group_size)
weights = {
"w_q_skinny": w_q_skinny,
"w_s_skinny": w_s_skinny,
"w_fp16": w_fp16,
"w_zp": w_zp,
}
save_path = args.save_path or f"bench_int4_res_n{N}_k{K}"
os.makedirs(save_path, exist_ok=True)
benchmark.run(
print_data=True,
show_plots=False,
save_path=save_path,
N=N,
K=K,
group_size=group_size,
weights=weights,
)
print("\nBenchmark finished!")
@@ -1,108 +0,0 @@
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
# Benchmark ReLUSquaredActivation: custom CUDA kernel vs forward_native, both
# eager and under torch.compile (Inductor fuses relu+square into one kernel).
import itertools
import torch
import torch.nn.functional as F
import vllm.model_executor.layers.activation # noqa: F401
from vllm.benchmarks.lib.utils import default_vllm_config
from vllm.triton_utils import triton
from vllm.utils.argparse_utils import FlexibleArgumentParser
from vllm.utils.torch_utils import STR_DTYPE_TO_TORCH_DTYPE, set_random_seed
# Capped so the largest tensor stays under 2**31 elements: the shared activation
# kernel computes the per-token pointer offset (blockIdx.x * d) in 32-bit, which
# overflows for tensors with >2**32 elements. Realistic token counts are well
# below this; the kernel-vs-native gap is already clear at these sizes.
batch_size_range = [1, 16, 128]
seq_len_range = [1, 16, 64, 1024]
intermediate_size = [3072, 9728, 12288]
configs = list(itertools.product(batch_size_range, seq_len_range, intermediate_size))
@default_vllm_config()
def benchmark_relu_squared(
batch_size: int,
seq_len: int,
intermediate_size: int,
provider: str,
dtype: torch.dtype,
):
device = "cuda"
num_tokens = batch_size * seq_len
set_random_seed(42)
torch.set_default_device(device)
x = torch.randn(num_tokens, intermediate_size, dtype=dtype, device=device)
out = torch.empty_like(x)
def native(x: torch.Tensor) -> torch.Tensor:
return torch.square(F.relu(x))
# Verify the custom kernel matches the native implementation before timing.
ref = native(x)
torch.ops._C.relu_squared(out, x)
torch.testing.assert_close(out, ref)
if provider == "custom":
# Custom CUDA kernel — single fused kernel.
fn = lambda: torch.ops._C.relu_squared(out, x)
elif provider == "native":
# forward_native, eager — relu and square as separate ops.
fn = lambda: native(x)
elif provider == "native_compiled":
# forward_native under torch.compile — Inductor fuses relu+square.
# This is the real production baseline (custom ops are off when
# Inductor is enabled), so it is the comparison reviewers care about.
compiled = torch.compile(native)
compiled(x) # warm up / trigger compilation before timing
fn = lambda: compiled(x)
ms, min_ms, max_ms = triton.testing.do_bench_cudagraph(
fn, quantiles=[0.5, 0.2, 0.8]
)
return ms, max_ms, min_ms
if __name__ == "__main__":
parser = FlexibleArgumentParser(
description="Benchmark ReLUSquaredActivation: custom kernel vs native."
)
parser.add_argument(
"--dtype",
type=str,
choices=["half", "bfloat16", "float"],
default="bfloat16",
)
args = parser.parse_args()
dtype = STR_DTYPE_TO_TORCH_DTYPE[args.dtype]
perf_report = triton.testing.perf_report(
triton.testing.Benchmark(
x_names=["batch_size", "seq_len", "intermediate_size"],
x_vals=configs,
line_arg="provider",
line_vals=["custom", "native_compiled", "native"],
line_names=[
"Custom Kernel",
"Native (torch.compile)",
"Native (eager)",
],
styles=[("blue", "-"), ("green", "-"), ("red", "-")],
ylabel="ms",
plot_name="relu_squared-eager-performance",
args={},
)
)
perf_report(
lambda batch_size, seq_len, intermediate_size, provider: benchmark_relu_squared(
batch_size, seq_len, intermediate_size, provider, dtype
)
).run(print_data=True)
@@ -170,8 +170,8 @@ def benchmark_config(
graph.replay()
torch.accelerator.synchronize()
start = torch.cuda.Event(enable_timing=True)
end = torch.cuda.Event(enable_timing=True)
start = torch.Event(enable_timing=True)
end = torch.Event(enable_timing=True)
latencies: list[float] = []
for _ in range(num_iters):
start.record()

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