diff --git a/.buildkite/image_build/image_build.sh b/.buildkite/image_build/image_build.sh index b81cd4ec0c..6d22acd5b1 100755 --- a/.buildkite/image_build/image_build.sh +++ b/.buildkite/image_build/image_build.sh @@ -79,6 +79,54 @@ setup_buildx_builder() { docker buildx ls | grep -E '^\*|^NAME' || docker buildx ls } +download_ci_hcl() { + # Returns non-zero instead of exiting: the optional export path must be + # able to give up without failing a build whose image already exists. + echo "--- :arrow_down: Downloading ci.hcl" + curl -sSfL -o "${CI_HCL_PATH}" "${CI_HCL_URL}" || return $? + echo "Downloaded to ${CI_HCL_PATH}" + + if [[ ! -f "${CI_HCL_PATH}" ]]; then + echo "Error: ci.hcl not found at ${CI_HCL_PATH}" + return 1 + fi +} + +export_kernel_symbol_map_from_cache() { + # The existing-image path: everything here is best effort. Called as + # `... || echo`, so errexit is suspended inside and any failure just + # returns; the build keeps the image it already has. + local rc=0 + download_ci_hcl \ + && setup_buildx_builder \ + && resolve_parent_commit \ + && export PARENT_COMMIT \ + && BUILD_TMP_DIR="$(mktemp -d)" \ + && export_kernel_symbol_map || rc=$? + rm -rf -- "${BUILD_TMP_DIR:-}" + return "${rc}" +} + +export_kernel_symbol_map() { + # Only when the build asked for it (nightly/post-merge). The map was + # produced inside the csrc-build stage during the main bake; this second + # bake of the scratch stage is a cache hit that just writes the file out. + if [[ "${VLLM_KERNEL_SYMBOL_MAP:-0}" != "1" ]]; then + return 0 + fi + echo "--- :world_map: Exporting kernel symbol map" + local out_dir="${BUILD_TMP_DIR}/kernel-symbol-map" + mkdir -p "${out_dir}" + if docker buildx bake -f "${VLLM_BAKE_FILE_PATH}" -f "${CI_HCL_PATH}" --progress plain \ + --set "kernel-symbol-map.output=type=local,dest=${out_dir}" kernel-symbol-map \ + && [[ -s "${out_dir}/kernel_symbol_map.json.gz" ]]; then + ls -la "${out_dir}" + (cd "${out_dir}" && buildkite-agent artifact upload "kernel_symbol_map.json.gz") + else + echo "kernel symbol map export failed; continuing without it" >&2 + fi +} + annotate_image_tags() { .buildkite/scripts/annotate-image-build.sh \ "${IMAGE_TAG:-}" "${IMAGE_TAG_LATEST:-}" @@ -90,6 +138,15 @@ check_and_skip_if_image_exists() { if docker manifest inspect "${IMAGE_TAG}" >/dev/null 2>&1; then echo "Image already exists: ${IMAGE_TAG}" echo "Skipping build" + if [[ "${VLLM_KERNEL_SYMBOL_MAP:-0}" == "1" ]]; then + # The image is reused, but the symbol map is an artifact of + # this build. Bake just the export target: its layers come + # from the registry cache, and the map step itself reruns + # if the cached image was built without the arg. Never + # fails the build: the image is already there. + export_kernel_symbol_map_from_cache \ + || echo "kernel symbol map: export from the cached image failed; continuing" >&2 + fi annotate_image_tags exit 0 fi @@ -154,7 +211,7 @@ print_bake_config() { local bake_tmp bake_tmp="$(mktemp -d)" BAKE_CONFIG_FILE="${bake_tmp}/bake-config-build-${BUILDKITE_BUILD_NUMBER:-local}.json" - docker buildx bake -f "${VLLM_BAKE_FILE_PATH}" -f "${CI_HCL_PATH}" --print "${TARGET}" | tee "${BAKE_CONFIG_FILE}" || true + docker buildx bake "${BAKE_FILES[@]}" --print "${TARGET}" | tee "${BAKE_CONFIG_FILE}" || true echo "Saved bake config to ${BAKE_CONFIG_FILE}" echo "--- :arrow_down: Uploading bake config to Buildkite" (cd "$(dirname "${BAKE_CONFIG_FILE}")" && buildkite-agent artifact upload "$(basename "${BAKE_CONFIG_FILE}")") @@ -230,8 +287,10 @@ TARGET="test-ci" VLLM_BAKE_FILE_PATH="${VLLM_BAKE_FILE_PATH:-docker/docker-bake.hcl}" BUILDER_NAME="${BUILDER_NAME:-vllm-builder}" CI_HCL_URL="${CI_HCL_URL:-https://raw.githubusercontent.com/vllm-project/ci-infra/main/docker/ci.hcl}" -CI_HCL_PATH="/tmp/ci.hcl" +CI_HCL_PATH="${CI_HCL_PATH:-/tmp/ci.hcl}" +ZSTD_HCL_PATH="${ZSTD_HCL_PATH:-.buildkite/image_build/zstd.hcl}" BUILDKIT_SOCKET="/run/buildkit/buildkitd.sock" +BAKE_FILES=(-f "${VLLM_BAKE_FILE_PATH}" -f "${CI_HCL_PATH}" -f "${ZSTD_HCL_PATH}") prepare_cache_tags ecr_login @@ -288,14 +347,7 @@ if [[ ! -f "${VLLM_BAKE_FILE_PATH}" ]]; then exit 1 fi -echo "--- :arrow_down: Downloading ci.hcl" -curl -sSfL -o "${CI_HCL_PATH}" "${CI_HCL_URL}" -echo "Downloaded to ${CI_HCL_PATH}" - -if [[ ! -f "${CI_HCL_PATH}" ]]; then - echo "Error: ci.hcl not found at ${CI_HCL_PATH}" - exit 1 -fi +download_ci_hcl || exit 1 setup_buildx_builder @@ -309,7 +361,7 @@ BUILD_TMP_DIR="$(mktemp -d)" trap 'rm -rf -- "${BUILD_TMP_DIR}"' EXIT BUILD_METADATA_FILE="${BUILD_TMP_DIR}/build-metadata.json" BUILD_STATUS=0 -docker --debug buildx bake -f "${VLLM_BAKE_FILE_PATH}" -f "${CI_HCL_PATH}" \ +docker --debug buildx bake "${BAKE_FILES[@]}" \ --progress plain --metadata-file "${BUILD_METADATA_FILE}" "${TARGET}" || BUILD_STATUS=$? record_buildkit_trace "${BUILD_METADATA_FILE}" || true @@ -319,5 +371,6 @@ if [[ "${BUILD_STATUS}" -ne 0 ]]; then fi echo "--- :white_check_mark: Build complete" +export_kernel_symbol_map || true annotate_image_tags diff --git a/.buildkite/image_build/zstd.hcl b/.buildkite/image_build/zstd.hcl new file mode 100644 index 0000000000..1f795c5f91 --- /dev/null +++ b/.buildkite/image_build/zstd.hcl @@ -0,0 +1,3 @@ +target "test-ci" { + output = ["type=registry,compression=zstd,compression-level=3,force-compression=true,oci-mediatypes=true"] +} diff --git a/.buildkite/intel_jobs/lm_eval_intel.yaml b/.buildkite/intel_jobs/lm_eval_intel.yaml index 1ef00a4794..e1f649ad91 100644 --- a/.buildkite/intel_jobs/lm_eval_intel.yaml +++ b/.buildkite/intel_jobs/lm_eval_intel.yaml @@ -26,3 +26,27 @@ steps: 'uv pip install "gpt-oss[eval]==0.0.5" && cd tests && pytest -s -v evals/gpt_oss/test_gpqa_correctness.py --config-list-file=configs/models-xpu.txt' +- label: LM Eval Large Models(2 GPUs) + key: lm-eval-large-models-2-gpus + timeout_in_minutes: 45 + device: intel_gpu + agent_tags: + label: production + gpu: 2+ + mem: 24+ + no_plugin: true + optional: true + working_dir: "." + env: + REGISTRY: "public.ecr.aws/q9t5s3a7" + REPO: "vllm-ci-test-repo" + VLLM_TEST_DEVICE: "xpu" + source_file_dependencies: + - vllm/ + - .buildkite/lm-eval-harness/ + commands: + - >- + bash .buildkite/scripts/hardware_ci/run-intel-test.sh + 'cd .buildkite/lm-eval-harness && + export VLLM_USE_DEEP_GEMM=0 && + pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-large-xpu-fp8.txt --tp-size=2' diff --git a/.buildkite/intel_jobs/misc_intel.yaml b/.buildkite/intel_jobs/misc_intel.yaml index 39b61b3f7d..5fa4d833d7 100644 --- a/.buildkite/intel_jobs/misc_intel.yaml +++ b/.buildkite/intel_jobs/misc_intel.yaml @@ -2,6 +2,34 @@ group: Miscellaneous Intel depends_on: - image-build-xpu steps: +- label: XPU Batch Invariance + key: xpu-batch-invariance + timeout_in_minutes: 60 + num_devices: 2 + device: intel_gpu + agent_tags: + label: production + gpu: 2+ + mem: 24+ + no_plugin: true + working_dir: "." + env: + REGISTRY: "public.ecr.aws/q9t5s3a7" + REPO: "vllm-ci-test-repo" + VLLM_TEST_DEVICE: "xpu" + source_file_dependencies: + - vllm/ + - tests/v1/determinism/test_xpu_batch_invariant_ut.py + - .buildkite/scripts/hardware_ci/run-xpu-batch-invariance.sh + - .buildkite/intel_jobs/misc_intel.yaml + commands: + - >- + bash .buildkite/scripts/hardware_ci/run-intel-test.sh + 'export VLLM_BATCH_INVARIANT=1 && + cd tests && + pytest -v -s v1/determinism/test_xpu_batch_invariant_ut.py && + bash ../.buildkite/scripts/hardware_ci/run-xpu-batch-invariance.sh' + - label: V1 Core + KV + Metrics key: v1-core-kv-metrics-intel timeout_in_minutes: 45 @@ -290,4 +318,4 @@ steps: - >- bash .buildkite/scripts/hardware_ci/run-intel-test.sh 'cd tests && - pytest -v -s compile/passes/test_qk_norm_rope_fusion.py' \ No newline at end of file + pytest -v -s compile/passes/test_qk_norm_rope_fusion.py' diff --git a/.buildkite/intel_jobs/model_runner_v2_intel.yaml b/.buildkite/intel_jobs/model_runner_v2_intel.yaml deleted file mode 100644 index 7f6767eba2..0000000000 --- a/.buildkite/intel_jobs/model_runner_v2_intel.yaml +++ /dev/null @@ -1,120 +0,0 @@ -group: Model Runner V2 Intel -depends_on: - - image-build-xpu -steps: -- label: Model Runner V2 Core Tests (Intel) - key: model-runner-v2-core-tests-intel - timeout_in_minutes: 45 - device: intel_gpu - agent_tags: - label: production - gpu: 2+ - mem: 24+ - no_plugin: true - working_dir: "." - env: - REGISTRY: "public.ecr.aws/q9t5s3a7" - REPO: "vllm-ci-test-repo" - VLLM_TEST_DEVICE: "xpu" - source_file_dependencies: - - vllm/v1/worker/gpu/ - - vllm/v1/worker/gpu_worker.py - - vllm/v1/core/sched/ - - vllm/v1/attention/ - - tests/v1/engine/test_llm_engine.py - - tests/v1/e2e/ - commands: - - >- - bash .buildkite/scripts/hardware_ci/run-intel-test.sh - 'export VLLM_USE_V2_MODEL_RUNNER=1 && - cd tests && - pytest -v -s v1/engine/test_llm_engine.py -k "not test_engine_metrics" && - pytest -v -s v1/e2e/general/test_context_length.py && - ENFORCE_EAGER=1 pytest -v -s v1/e2e/general/test_async_scheduling.py -k "not ngram" && - pytest -v -s entrypoints/llm/test_struct_output_generate.py -k "xgrammar and not speculative_config6 and not speculative_config7 and not speculative_config8 and not speculative_config0" && - pytest -v -s v1/e2e/general/test_min_tokens.py' - -- label: Model Runner V2 Examples (Intel) - key: model-runner-v2-examples-intel - timeout_in_minutes: 45 - device: intel_gpu - agent_tags: - label: production - gpu: 1+ - mem: 24+ - no_plugin: true - working_dir: "." - env: - REGISTRY: "public.ecr.aws/q9t5s3a7" - REPO: "vllm-ci-test-repo" - VLLM_TEST_DEVICE: "xpu" - source_file_dependencies: - - vllm/v1/worker/gpu/ - - vllm/v1/core/sched/ - - vllm/v1/worker/gpu_worker.py - - examples/basic/offline_inference/ - - examples/generate/multimodal/ - - examples/features/ - commands: - - >- - bash .buildkite/scripts/hardware_ci/run-intel-test.sh - 'export VLLM_USE_V2_MODEL_RUNNER=1 && - cd examples && - python3 basic/offline_inference/chat.py && - python3 basic/offline_inference/generate.py --model facebook/opt-125m && - python3 generate/multimodal/vision_language_offline.py --seed 0 && - python3 features/automatic_prefix_caching/prefix_caching_offline.py' - -- label: Model Runner V2 Distributed (2 GPUs) - key: model-runner-v2-distributed-2-gpus - timeout_in_minutes: 50 - device: intel_gpu - agent_tags: - label: production - gpu: 2+ - mem: 16+ - no_plugin: true - working_dir: "." - env: - REGISTRY: "public.ecr.aws/q9t5s3a7" - REPO: "vllm-ci-test-repo" - VLLM_TEST_DEVICE: "xpu" - source_file_dependencies: - - vllm/v1/worker/gpu/ - - vllm/v1/worker/gpu_worker.py - - tests/basic_correctness/test_basic_correctness.py - - tests/v1/distributed/test_async_llm_dp.py - - tests/v1/distributed/test_eagle_dp.py - commands: - - >- - bash .buildkite/scripts/hardware_ci/run-intel-test.sh - 'export VLLM_USE_V2_MODEL_RUNNER=1 && - cd tests && - TARGET_TEST_SUITE=L4 pytest -v -s basic_correctness/test_basic_correctness.py -m "distributed\(num_gpus=2\)" -k "not ray and not True"' - -- label: Model Runner V2 Spec Decode - key: model-runner-v2-spec-decode - timeout_in_minutes: 50 - device: intel_gpu - agent_tags: - label: production - gpu: 1+ - mem: 24+ - no_plugin: true - working_dir: "." - env: - REGISTRY: "public.ecr.aws/q9t5s3a7" - REPO: "vllm-ci-test-repo" - VLLM_TEST_DEVICE: "xpu" - source_file_dependencies: - - vllm/v1/worker/gpu/ - - vllm/v1/worker/gpu_worker.py - - tests/v1/spec_decode/test_max_len.py - - tests/v1/spec_decode/test_rejection_sampler_utils.py - - tests/v1/spec_decode/test_synthetic_rejection_sampler_utils.py - commands: - - >- - bash .buildkite/scripts/hardware_ci/run-intel-test.sh - 'export VLLM_USE_V2_MODEL_RUNNER=1 && - cd tests && - pytest -v -s v1/spec_decode/test_synthetic_rejection_sampler_utils.py' diff --git a/.buildkite/intel_jobs/test-intel.yaml b/.buildkite/intel_jobs/test-intel.yaml index ef2b4ba073..d5ae40ec6e 100644 --- a/.buildkite/intel_jobs/test-intel.yaml +++ b/.buildkite/intel_jobs/test-intel.yaml @@ -106,7 +106,7 @@ steps: pytest -v -s v1/test_serial_utils.py && pytest -v -s v1/e2e/general/test_correctness_sliding_window.py && pytest -v -s v1/spec_decode --ignore=v1/spec_decode/test_max_len.py --ignore=v1/spec_decode/test_speculators_eagle3.py --ignore=v1/spec_decode/test_acceptance_length.py --ignore=v1/spec_decode/test_speculators_correctness.py && - pytest -v -s v1/kv_connector/unit --ignore=v1/kv_connector/unit/test_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' + pytest -v -s v1/kv_connector/unit --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 --deselect "v1/kv_connector/unit/test_nixl_connector.py::test_abort_timeout_on_prefiller[ray]"' parallelism: 2 - label: "XPU server test" key: xpu-server-test diff --git a/.buildkite/lm-eval-harness/configs/NVIDIA-Nemotron-3-Nano-30B-A3B-FP8-xpu.yaml b/.buildkite/lm-eval-harness/configs/NVIDIA-Nemotron-3-Nano-30B-A3B-FP8-xpu.yaml new file mode 100644 index 0000000000..72ba861195 --- /dev/null +++ b/.buildkite/lm-eval-harness/configs/NVIDIA-Nemotron-3-Nano-30B-A3B-FP8-xpu.yaml @@ -0,0 +1,16 @@ +model_name: "nvidia/NVIDIA-Nemotron-3-Nano-30B-A3B-FP8" +tasks: +- name: "gsm8k" + metrics: + - name: "exact_match,strict-match" + value: 0.7142 + - name: "exact_match,flexible-extract" + value: 0.4579 +limit: 256 +num_fewshot: 5 +max_model_len: 262144 +kv_cache_dtype: fp8 +enforce_eager: true +apply_chat_template: true +fewshot_as_multiturn: true +trust_remote_code: true diff --git a/.buildkite/lm-eval-harness/configs/models-large-xpu-fp8.txt b/.buildkite/lm-eval-harness/configs/models-large-xpu-fp8.txt new file mode 100644 index 0000000000..73b28f062c --- /dev/null +++ b/.buildkite/lm-eval-harness/configs/models-large-xpu-fp8.txt @@ -0,0 +1 @@ +NVIDIA-Nemotron-3-Nano-30B-A3B-FP8-xpu.yaml \ No newline at end of file diff --git a/.buildkite/scripts/hardware_ci/run-xpu-batch-invariance.sh b/.buildkite/scripts/hardware_ci/run-xpu-batch-invariance.sh new file mode 100755 index 0000000000..4a8e41d580 --- /dev/null +++ b/.buildkite/scripts/hardware_ci/run-xpu-batch-invariance.sh @@ -0,0 +1,243 @@ +#!/usr/bin/env bash +set -euo pipefail + +repo_root="$(cd -- "$(dirname -- "${BASH_SOURCE[0]}")/../../.." && pwd)" + +model="ibm-research/PowerMoE-3b" +port=8000 +num_decode_tokens=128 +server_pid= +server_log= + +# XPU MoE kernels may not safely handle the -1 routes emitted for padding. +export VLLM_MOE_SKIP_PADDING=0 + +cd "$repo_root" + +cleanup_server() { + local attempt + + if [[ -n "$server_pid" ]] && kill -0 "$server_pid" 2>/dev/null; then + kill "$server_pid" 2>/dev/null || true + + for ((attempt = 0; attempt < 30; attempt++)); do + if ! kill -0 "$server_pid" 2>/dev/null; then + break + fi + sleep 1 + done + + if kill -0 "$server_pid" 2>/dev/null; then + echo "vLLM did not stop after 30 seconds; sending SIGKILL" >&2 + kill -KILL "$server_pid" 2>/dev/null || true + fi + + wait "$server_pid" 2>/dev/null || true + fi + + server_pid= +} + +on_exit() { + local status=$? + + cleanup_server + if [[ "$status" -ne 0 && -n "$server_log" && -f "$server_log" ]]; then + if command -v buildkite-agent >/dev/null 2>&1; then + buildkite-agent artifact upload "$server_log" || + echo "Failed to upload $server_log as a Buildkite artifact" >&2 + fi + fi + [[ -z "$server_log" ]] || rm -f "$server_log" + exit "$status" +} +trap on_exit EXIT + +wait_for_server() { + for _ in $(seq 1 600); do + if curl -sf "http://127.0.0.1:$port/health" >/dev/null; then + return + fi + if ! kill -0 "$server_pid" 2>/dev/null; then + echo "vLLM exited before becoming healthy; see the server log above." >&2 + return 1 + fi + sleep 1 + done + echo "Timed out waiting for vLLM server health check" >&2 + return 1 +} + +request_outputs() { + local prompts_json=$1 + local expected_choices=$2 + local response + + response=$(printf \ + '{"model":"%s","prompt":%s,"temperature":0.7,"top_p":0.95,"seed":0,"max_tokens":%s,"logprobs":64,"ignore_eos":true,"return_token_ids":true}' \ + "$model" "$prompts_json" "$num_decode_tokens" | + curl --fail --silent --show-error \ + --max-time 1200 \ + "http://127.0.0.1:$port/v1/completions" \ + -H "Content-Type: application/json" \ + --data-binary @-) + + local -a token_id_arrays logprob_arrays + mapfile -t token_id_arrays < <( + grep -o '"token_ids":\[[^]]*\]' <<<"$response" | + sed 's/^"token_ids":\[//; s/\]$//' + ) + mapfile -t logprob_arrays < <( + grep -o '"token_logprobs":\[[^]]*\]' <<<"$response" | + sed 's/^"token_logprobs":\[//; s/\]$//' + ) + if [[ "${#token_id_arrays[@]}" -ne "$expected_choices" ]] || + [[ "${#logprob_arrays[@]}" -ne "$expected_choices" ]]; then + echo "Expected $expected_choices choices with token IDs and logits" >&2 + return 1 + fi + local index + local array + local value + local -a values + for index in "${!logprob_arrays[@]}"; do + array="${logprob_arrays[index]}" + IFS=, read -r -a values <<<"$array" + if [[ "${#values[@]}" -ne "$num_decode_tokens" ]]; then + echo "Completion returned an unexpected number of decode steps" >&2 + return 1 + fi + IFS=, read -r -a values <<<"${token_id_arrays[index]}" + if [[ "${#values[@]}" -ne "$num_decode_tokens" ]]; then + echo "Completion returned an unexpected number of generated tokens" >&2 + return 1 + fi + for value in "${values[@]}"; do + if [[ ! "$value" =~ ^[0-9]+$ ]]; then + echo "Completion returned an invalid token ID" >&2 + return 1 + fi + done + IFS=, read -r -a values <<<"$array" + for value in "${values[@]}"; do + if [[ ! "$value" =~ ^-?([0-9]+([.][0-9]*)?|[.][0-9]+)([eE][+-]?[0-9]+)?$ ]]; then + echo "Completion returned non-finite raw logits" >&2 + return 1 + fi + done + printf '%s|%s\n' "${token_id_arrays[index]}" "${logprob_arrays[index]}" + done +} + +make_prompt() { + local length=$1 + local offset=${2:-0} + local i + local prompt="Use the following system-design context to answer the question. Scenario $offset:" + for ((i = 0; i < length; i++)); do + prompt+=" context" + done + prompt+="\\n\\nQuestion: How should batch-invariant inference behave when this same request is decoded alongside unrelated requests?\\nAnswer:" + printf '"%s"' "$prompt" +} + +first_difference() { + local actual=$1 + local expected=$2 + local position=$3 + local actual_tokens=${actual%%|*} + local expected_tokens=${expected%%|*} + local actual_logits=${actual#*|} + local expected_logits=${expected#*|} + local -a actual_token_values + local -a expected_token_values + local -a actual_logit_values + local -a expected_logit_values + local step + + IFS=, read -r -a actual_token_values <<<"$actual_tokens" + IFS=, read -r -a expected_token_values <<<"$expected_tokens" + IFS=, read -r -a actual_logit_values <<<"$actual_logits" + IFS=, read -r -a expected_logit_values <<<"$expected_logits" + for ((step = 0; step < num_decode_tokens; step++)); do + if [[ "${actual_token_values[step]}" != "${expected_token_values[step]}" ]]; then + printf 'batch position %s: step=%s, baseline token=%s, actual token=%s' \ + "$position" \ + "$step" \ + "${expected_token_values[step]}" \ + "${actual_token_values[step]}" + return + fi + if [[ "${actual_logit_values[step]}" != "${expected_logit_values[step]}" ]]; then + printf 'batch position %s: step=%s, token=%s, baseline logit=%s, actual logit=%s' \ + "$position" \ + "$step" \ + "${expected_token_values[step]}" \ + "${expected_logit_values[step]}" \ + "${actual_logit_values[step]}" + return + fi + done +} + +run_e2e() { + local needle + local batch + local baseline_output + local batch_response + local difference + local position + local -a fillers + local -a batch_outputs + local -a differences=() + + echo "Testing PowerMoE-3b: TP=2, VLLM_BATCH_INVARIANT=1, VLLM_MOE_SKIP_PADDING=0" + server_log=$(mktemp "${TMPDIR:-/tmp}/xpu-batch-invariance.XXXXXX.log") + VLLM_BATCH_INVARIANT=1 vllm serve "$model" \ + --tensor-parallel-size 2 \ + --max-model-len 1024 \ + --max-num-seqs 16 \ + --max-num-batched-tokens 512 \ + --enable-chunked-prefill \ + --no-enable-prefix-caching \ + --enforce-eager \ + --seed 0 \ + --max-logprobs 64 \ + --logprobs-mode raw_logits \ + --port "$port" > >(tee "$server_log") 2>&1 & + server_pid=$! + wait_for_server + + needle=$(make_prompt 257) + local filler_lengths=(7 31 63 95 127 159 191 223 255 287 319 383 447) + for position in "${!filler_lengths[@]}"; do + fillers+=( "$(make_prompt "${filler_lengths[position]}" "$((position + 1))")" ) + done + batch="[$needle,${fillers[0]},${fillers[1]},${fillers[2]},${fillers[3]},${fillers[4]},${fillers[5]},$needle,${fillers[6]},${fillers[7]},${fillers[8]},${fillers[9]},${fillers[10]},${fillers[11]},${fillers[12]},$needle]" + + echo "Requesting BS=1 baseline ($num_decode_tokens decode tokens)..." + baseline_output=$(request_outputs "$needle" 1) + echo "Received BS=1 baseline." + echo "Requesting BS=16 batch ($((16 * num_decode_tokens)) decode tokens)..." + batch_response=$(request_outputs "$batch" 16) + mapfile -t batch_outputs <<<"$batch_response" + echo "Received BS=16 batch; comparing generated tokens and logits..." + + for position in 0 7 15; do + difference=$(first_difference "${batch_outputs[position]}" "$baseline_output" "$position") + if [[ -n "$difference" ]]; then + differences+=("$difference") + fi + done + + if [[ "${#differences[@]}" -ne 0 ]]; then + echo "Batch-invariant PowerMoE-3b output changed: ${differences[*]}" >&2 + return 1 + fi + echo "PowerMoE-3b batch invariance check passed" + cleanup_server + rm -f "$server_log" + server_log= +} + +run_e2e diff --git a/.buildkite/scripts/initialized-snapshot-e2e.sh b/.buildkite/scripts/initialized-snapshot-e2e.sh new file mode 100755 index 0000000000..cd0511ffdd --- /dev/null +++ b/.buildkite/scripts/initialized-snapshot-e2e.sh @@ -0,0 +1,311 @@ +#!/usr/bin/env bash +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +set -Eeuo pipefail +readonly MODEL="Qwen/Qwen3-0.6B" +readonly MODEL_REVISION="c1899de289a04d12100db370d81485cdf75e47ca" +readonly EXPECTED_TOKEN_ID="12095" +readonly EXPECTED_TEXT=" Paris" +readonly CLEANUP_RESERVE_S=300 +RUN_LIMIT_S="${SNAPSHOT_E2E_RUN_LIMIT_S:-3300}" +[[ "$RUN_LIMIT_S" =~ ^[1-9][0-9]{2,3}$ ]] \ + || { echo "SNAPSHOT_E2E_RUN_LIMIT_S must be a bounded decimal integer" >&2; exit 1; } +(( RUN_LIMIT_S >= CLEANUP_RESERVE_S + 60 && RUN_LIMIT_S <= 3300 )) \ + || { echo "SNAPSHOT_E2E_RUN_LIMIT_S must be between 360 and 3300" >&2; exit 1; } +readonly RUN_LIMIT_S +readonly HOST_HF_CACHE="${SNAPSHOT_E2E_HOST_HF_CACHE:-}" +HF_CACHE_MOUNT=() +if [[ -n "$HOST_HF_CACHE" ]]; then + [[ "$HOST_HF_CACHE" == /* && -d "$HOST_HF_CACHE" ]] \ + || { echo "SNAPSHOT_E2E_HOST_HF_CACHE must be an existing absolute directory" >&2; exit 1; } + HF_CACHE_MOUNT=(--mount "type=bind,source=$HOST_HF_CACHE,target=/e2e/hf,readonly") + PREPARE_ARTIFACT_ROOT='install -d -m 0700 /e2e' +else + PREPARE_ARTIFACT_ROOT='mkdir -m 0700 /e2e && mkdir -m 0700 /e2e/hf' +fi +readonly PREPARE_ARTIFACT_ROOT +if (( $# != 2 )); then + echo "usage: $0 IMAGE_TAG BUILDKITE_COMMIT" >&2 + exit 2 +fi +readonly IMAGE_TAG="$1" +readonly EXPECTED_COMMIT="$2" +readonly DEADLINE_EPOCH="$(( $(date +%s) + RUN_LIMIT_S ))" +RUN_ID="" CONTAINER_NAME="" CONTAINER_ID="" +GPU_UUID="" PORT_ONE="" PORT_TWO="" +LINK_REMAP_BASELINE="" +SNAPSHOT_PIDS=() +die() { + echo "initialized snapshot E2E: $*" >&2 + exit 1 +} + +run() { + local label="$1" maximum_s="$2" available limit + shift 2 + available="$((DEADLINE_EPOCH - $(date +%s) - CLEANUP_RESERVE_S))" + (( available > 0 )) || die "no time remains before the cleanup reserve" + (( available < maximum_s )) && limit="$available" || limit="$maximum_s" + echo "--- ${label} (timeout ${limit}s)" >&2 + timeout --signal=TERM --kill-after=30s "${limit}s" "$@" +} + +link_remaps() { + find /dev/shm -maxdepth 1 -name 'link_remap.*' \ + -printf '%p|%D|%i|%m|%s\n' | sort +} + +gpu_pids() { + nvidia-smi --id="$GPU_UUID" --query-compute-apps=pid \ + --format=csv,noheader,nounits | sed '/^[[:space:]]*$/d' +} + +port_listeners() { + ss -H -ltn "sport = :$1" +} + +state_clean() { + local pid port output + for pid in "${SNAPSHOT_PIDS[@]}"; do + [[ ! -e "/proc/$pid" ]] || return 1 + done + output="$(gpu_pids)" || return 1 + [[ -z "$output" ]] || return 1 + output="$(link_remaps)" || return 1 + [[ "$output" == "$LINK_REMAP_BASELINE" ]] || return 1 + for port in "$PORT_ONE" "$PORT_TWO"; do + [[ -n "$port" ]] || continue + output="$(port_listeners "$port")" || return 1 + [[ -z "$output" ]] || return 1 + done +} + +wait_clean() { + local _attempt + for ((_attempt = 0; _attempt < 60; _attempt++)); do + state_clean && return 0 + sleep 1 + done + return 1 +} + +container_names() { + timeout 15s docker container ls --all --format '{{.Names}}' +} + +cleanup() { + local failed=0 names="" actual="" expected="" + if [[ -n "$CONTAINER_NAME" ]]; then + names="$(container_names 2>/dev/null)" || return 1 + if grep -Fxq "$CONTAINER_NAME" <<< "$names"; then + actual="$(docker container inspect \ + --format '{{.Id}}|{{index .Config.Labels "ai.vllm.snapshot.e2e.run"}}' \ + "$CONTAINER_NAME" 2>/dev/null)" || return 1 + expected="${CONTAINER_ID:-${actual%%|*}}|$RUN_ID" + if [[ "$actual" != "$expected" ]]; then + echo "refusing to remove a container with mismatched identity" >&2 + failed=1 + else + timeout 60s docker rm --force "$CONTAINER_NAME" >/dev/null \ + || failed=1 + fi + fi + names="$(container_names 2>/dev/null)" || return 1 + grep -Fxq "$CONTAINER_NAME" <<< "$names" && failed=1 + fi + if [[ -n "$PORT_TWO" ]] && ! wait_clean; then + echo "process, GPU, port, or link-remap residue remained" >&2 + failed=1 + fi + return "$failed" +} + +on_exit() { + local status=$? + trap - EXIT INT TERM + set +e + cleanup || status=1 + exit "$status" +} + +trap on_exit EXIT +trap 'exit 130' INT +trap 'exit 143' TERM + +parse_inspect() { + timeout 30s python3 -c ' +import json, math, sys +data = json.load(sys.stdin) +model, revision, gpu_uuid, token_id, text = sys.argv[1:] +expected = { + "boundary": "post-engine-init-reloadable-state-released", + "model": model, "model_revision": revision, + "gpu_uuid": gpu_uuid, "oracle_token_ids": [int(token_id)], + "oracle_text": text, +} +bad = next((field for field, value in expected.items() if data.get(field) != value), None) +if bad: + raise SystemExit(f"inspect field {bad} mismatch: {data.get(bad)!r}") +logprob = data.get("oracle_sampled_token_logprob") +if type(logprob) is not float or not math.isfinite(logprob): + raise SystemExit("inspect logprob is not finite") +pids = data.get("process_tree") +if not pids: + raise SystemExit("inspect process_tree is empty") +for pid in pids: + print(f"pid={pid}") +print(f"oracle_logprob={logprob}") +' "$MODEL" "$MODEL_REVISION" "$GPU_UUID" "$EXPECTED_TOKEN_ID" \ + "$EXPECTED_TEXT" +} + +assert_oracle() { + local port="$1" ordinal="$2" response + response="$(run "restore ${ordinal} public completion" 120 curl \ + --fail --silent --show-error --max-time 120 \ + --header 'Content-Type: application/json' \ + --data "{\"model\":\"${MODEL}\",\"prompt\":\"The capital of France is\",\"temperature\":0,\"seed\":0,\"min_tokens\":1,\"max_tokens\":1,\"logprobs\":0,\"return_token_ids\":true}" \ + "http://127.0.0.1:${port}/v1/completions")" + printf '%s' "$response" | timeout 30s python3 -c ' +import json, math, sys +choice = json.load(sys.stdin)["choices"][0] +token_id, text, expected, ordinal = sys.argv[1:] +values = choice["logprobs"]["token_logprobs"] +actual = float(values[0]) if len(values) == 1 else math.nan +if (choice["token_ids"] != [int(token_id)] or choice["text"] != text + or not math.isfinite(actual) or abs(actual - float(expected)) > 1e-3): + raise SystemExit(f"public oracle mismatch: {choice!r}") +print(f"restore={ordinal} token_id={token_id} text={text!r} logprob={actual}") +' "$EXPECTED_TOKEN_ID" "$EXPECTED_TEXT" "$INSPECT_LOGPROB" "$ordinal" +} + +[[ "$(uname -s)" == "Linux" && "$(uname -m)" == "x86_64" ]] \ + || die "host must be Linux x86_64" +[[ "$EXPECTED_COMMIT" =~ ^[0-9a-f]{40}$ ]] \ + || die "BUILDKITE_COMMIT must be a full lowercase commit hash" +[[ "$(< /proc/sys/kernel/io_uring_disabled)" == "2" ]] \ + || die "kernel.io_uring_disabled must already equal 2" +run "Docker preflight" 60 docker info >/dev/null + +VISIBLE_GPU="${NVIDIA_VISIBLE_DEVICES:-}" +if [[ -z "$VISIBLE_GPU" || "$VISIBLE_GPU" == "all" || "$VISIBLE_GPU" == "void" ]]; then + VISIBLE_GPU="${CUDA_VISIBLE_DEVICES:-}" +fi +if [[ -z "$VISIBLE_GPU" || "$VISIBLE_GPU" == "all" || "$VISIBLE_GPU" == "void" ]]; then + CDI_GPU="${BUILDKITE_PLUGIN_DOCKER_DEVICES_0:-}" + if [[ "$CDI_GPU" == nvidia.com/gpu=* ]]; then + VISIBLE_GPU="${CDI_GPU#nvidia.com/gpu=}" + else + VISIBLE_GPU="" + fi +fi +[[ -n "$VISIBLE_GPU" && "$VISIBLE_GPU" != all && "$VISIBLE_GPU" != void \ + && "$VISIBLE_GPU" != *","* ]] \ + || die "the runner must expose exactly one selected GPU" +GPU_ROW="$(nvidia-smi --id="$VISIBLE_GPU" \ + --query-gpu=name,uuid,mig.mode.current,memory.total \ + --format=csv,noheader,nounits)" +IFS=',' read -r GPU_NAME GPU_UUID MIG_MODE GPU_MEMORY_MIB <<< "$GPU_ROW" +for variable in GPU_NAME GPU_UUID MIG_MODE GPU_MEMORY_MIB; do + printf -v "$variable" '%s' "${!variable#"${!variable%%[![:space:]]*}"}" +done +[[ "$GPU_NAME" == *H200* && "$GPU_UUID" == GPU-* && "$MIG_MODE" == Disabled ]] \ + || die "selected device is not a full non-MIG H200: $GPU_ROW" +(( GPU_MEMORY_MIB >= 130000 )) || die "selected H200 does not expose full memory" +GPU_PIDS="$(gpu_pids)" || die "failed to inspect the selected GPU" +[[ -z "$GPU_PIDS" ]] || die "selected H200 is not idle" + +RUN_ID="snapshot-e2e-$(date +%s)-$$-$RANDOM" +CONTAINER_NAME="vllm-$RUN_ID" +read -r PORT_ONE PORT_TWO < <(python3 -c 'import socket; s=[socket.socket(),socket.socket()]; [x.bind(("127.0.0.1",0)) for x in s]; print(*(x.getsockname()[1] for x in s))') +LINK_REMAP_BASELINE="$(link_remaps)" + +run "pull exact candidate image" 900 docker pull --platform linux/amd64 \ + "$IMAGE_TAG" >/dev/null +IFS='|' read -r IMAGE_ID IMAGE_PLATFORM OCI_COMMIT VLLM_COMMIT < <( + docker image inspect \ + --format '{{.Id}}|{{.Os}}/{{.Architecture}}|{{index .Config.Labels "org.opencontainers.image.revision"}}|{{index .Config.Labels "ai.vllm.build.commit"}}' \ + "$IMAGE_TAG" +) +[[ "$IMAGE_ID" =~ ^sha256:[0-9a-f]{64}$ && "$IMAGE_PLATFORM" == linux/amd64 ]] \ + || die "candidate image identity is invalid" +[[ "$OCI_COMMIT" == "$EXPECTED_COMMIT" && "$VLLM_COMMIT" == "$EXPECTED_COMMIT" ]] \ + || die "candidate image commit labels do not match" +CONTAINER_ID="$(run "start exact candidate container" 90 docker run --detach \ + --name "$CONTAINER_NAME" --label "ai.vllm.snapshot.e2e.run=$RUN_ID" \ + --gpus "device=$GPU_UUID" --user 0 --privileged --pid=host --ipc=host \ + --network=host --env "CUDA_VISIBLE_DEVICES=$GPU_UUID" --env HF_HUB_DISABLE_TELEMETRY=1 \ + --env VLLM_NO_USAGE_STATS=1 --env VLLM_USE_V2_MODEL_RUNNER=1 \ + ${HF_CACHE_MOUNT[@]+"${HF_CACHE_MOUNT[@]}"} \ + --entrypoint sleep "$IMAGE_ID" infinity)" +[[ "$CONTAINER_ID" =~ ^[0-9a-f]{64}$ ]] || die "invalid container ID" +run "prepare private artifact root" 60 docker exec "$CONTAINER_NAME" sh -c \ + "$PREPARE_ARTIFACT_ROOT && df -Pk / | awk 'NR == 2 && \$4 >= 10485760 {ok=1} END {exit !ok}'" +run "verify snapshot runtime" 60 docker exec "$CONTAINER_NAME" sh -c \ + 'criu --version; test -x /usr/local/sbin/cuda-checkpoint; test -f /usr/local/lib/criu/cuda_plugin.so' +run "verify selected CUDA device" 60 docker exec "$CONTAINER_NAME" python3 -c ' +import sys, torch +if torch.cuda.device_count() != 1: + raise SystemExit("expected exactly one visible CUDA device") +actual = str(torch.cuda.get_device_properties(0).uuid).removeprefix("GPU-") +if actual != sys.argv[1].removeprefix("GPU-"): + raise SystemExit(f"CUDA device {actual} does not match {sys.argv[1]}") +print(f"selected_cuda_uuid=GPU-{actual}") +' "$GPU_UUID" +OFFLINE_EXEC=(docker exec --env HF_HOME=/e2e/hf --env HF_HUB_OFFLINE=1 \ + --env TRANSFORMERS_OFFLINE=1 --env VLLM_SNAPSHOT_TIMEOUT_S=900 \ + "$CONTAINER_NAME") +if [[ -n "$HOST_HF_CACHE" ]]; then + run "validate pinned model in read-only host cache" 120 \ + "${OFFLINE_EXEC[@]}" python3 -c \ + 'from pathlib import Path; from huggingface_hub import snapshot_download; import sys; p=snapshot_download(repo_id=sys.argv[1], revision=sys.argv[2], local_files_only=True); assert Path(p).name == sys.argv[2]' \ + "$MODEL" "$MODEL_REVISION" +else + run "prefetch pinned public model" 900 docker exec --env HF_HOME=/e2e/hf \ + "$CONTAINER_NAME" python3 -c \ + 'from pathlib import Path; from huggingface_hub import snapshot_download; import sys; p=snapshot_download(repo_id=sys.argv[1], revision=sys.argv[2]); assert Path(p).name == sys.argv[2]' \ + "$MODEL" "$MODEL_REVISION" +fi + +run "create compact initialized snapshot" 1200 "${OFFLINE_EXEC[@]}" \ + vllm snapshot create "$MODEL" --revision "$MODEL_REVISION" \ + --tokenizer-revision "$MODEL_REVISION" --snapshot-dir /e2e/artifact \ + --dtype float16 --max-model-len 512 --gpu-memory-utilization 0.50 +wait_clean || die "create left process, GPU, port, or link-remap residue" + +INSPECT_JSON="$(run "inspect compact initialized snapshot" 120 \ + "${OFFLINE_EXEC[@]}" vllm snapshot inspect /e2e/artifact)" +INSPECT_RECEIPT="$(printf '%s' "$INSPECT_JSON" | parse_inspect)" +mapfile -t SNAPSHOT_PIDS < <(sed -n 's/^pid=//p' <<< "$INSPECT_RECEIPT") +INSPECT_LOGPROB="$(sed -n 's/^oracle_logprob=//p' <<< "$INSPECT_RECEIPT")" +(( ${#SNAPSHOT_PIDS[@]} > 0 )) || die "inspect returned no PIDs" +state_clean || die "create left recorded state alive" +echo "$INSPECT_RECEIPT" + +run "restore compact snapshot 1" 900 "${OFFLINE_EXEC[@]}" \ + vllm snapshot restore /e2e/artifact --host 127.0.0.1 --port "$PORT_ONE" +PORT_LISTENERS="$(port_listeners "$PORT_ONE")" || die "failed to inspect restore 1 port" +[[ -n "$PORT_LISTENERS" ]] || die "restore 1 did not bind its port" +GPU_PIDS="$(gpu_pids)" || die "failed to inspect the selected GPU" +[[ -n "$GPU_PIDS" ]] || die "restore 1 did not engage the selected GPU" +assert_oracle "$PORT_ONE" 1 +run "stop container after restore 1" 60 docker stop --time 20 "$CONTAINER_NAME" >/dev/null +[[ "$(docker container inspect --format '{{.Id}}' "$CONTAINER_NAME")" == "$CONTAINER_ID" ]] \ + || die "docker stop replaced the container" +wait_clean || die "restore 1 teardown left residue" + +run "start the same container" 60 docker start "$CONTAINER_NAME" >/dev/null +[[ "$(docker container inspect --format '{{.Id}}' "$CONTAINER_NAME")" == "$CONTAINER_ID" ]] \ + || die "docker start replaced the container" +run "restore compact snapshot 2" 900 "${OFFLINE_EXEC[@]}" \ + vllm snapshot restore /e2e/artifact --host 127.0.0.1 --port "$PORT_TWO" +PORT_LISTENERS="$(port_listeners "$PORT_TWO")" || die "failed to inspect restore 2 port" +[[ -n "$PORT_LISTENERS" ]] || die "restore 2 did not bind its port" +GPU_PIDS="$(gpu_pids)" || die "failed to inspect the selected GPU" +[[ -n "$GPU_PIDS" ]] || die "restore 2 did not engage the selected GPU" +assert_oracle "$PORT_TWO" 2 +run "stop container after restore 2" 60 docker stop --time 20 "$CONTAINER_NAME" >/dev/null +wait_clean || die "restore 2 teardown left residue" +cleanup || die "exact cleanup failed" +echo "container_id=$CONTAINER_ID restore_1_port=$PORT_ONE restore_2_port=$PORT_TWO" +echo "result=PASS" diff --git a/.buildkite/scripts/install-kv-connectors.sh b/.buildkite/scripts/install-kv-connectors.sh index ef47c3625f..8733a21347 100755 --- a/.buildkite/scripts/install-kv-connectors.sh +++ b/.buildkite/scripts/install-kv-connectors.sh @@ -9,9 +9,20 @@ if python3 -c "import torch; raise SystemExit(0 if torch.version.hip is not None exit 0 fi -REQUIREMENTS_FILE="${KV_CONNECTORS_REQUIREMENTS:-/vllm-workspace/requirements/kv_connectors.txt}" +# Select requirements before optional packages can change the installed Torch. +CUDA_MAJOR=$(python3 -c 'import torch; print(torch.version.cuda.split(".", 1)[0])') +DEFAULT_REQUIREMENTS=/vllm-workspace/requirements/kv_connectors.txt +if [ "${CUDA_MAJOR}" = "12" ]; then + DEFAULT_REQUIREMENTS=/vllm-workspace/requirements/kv_connectors_cu12.txt +fi +REQUIREMENTS_FILE="${KV_CONNECTORS_REQUIREMENTS:-${DEFAULT_REQUIREMENTS}}" -uv pip install --system -r "${REQUIREMENTS_FILE}" +CONSTRAINTS_DIR=$(mktemp -d) +trap 'rm -rf "${CONSTRAINTS_DIR}"' EXIT +uv pip freeze --system > "${CONSTRAINTS_DIR}/installed.txt" +grep -E '^(torch|torchaudio|torchvision|triton|cuda-bindings|cuda-python|cuda-toolkit)==' \ + "${CONSTRAINTS_DIR}/installed.txt" > "${CONSTRAINTS_DIR}/constraints.txt" +uv pip install --system -c "${CONSTRAINTS_DIR}/constraints.txt" -r "${REQUIREMENTS_FILE}" KV_METADATA=$(python3 - <<'PY' import importlib.metadata as metadata @@ -56,5 +67,5 @@ PY # build first to avoid a clash. if [ "${CUDA_MAJOR}" = "13" ] && [ -n "${MOONCAKE_VERSION}" ]; then uv pip uninstall --system mooncake-transfer-engine 2>/dev/null || true - uv pip install --system "mooncake-transfer-engine-cuda13==${MOONCAKE_VERSION}" + uv pip install --system -c "${CONSTRAINTS_DIR}/constraints.txt" "mooncake-transfer-engine-cuda13==${MOONCAKE_VERSION}" fi diff --git a/.buildkite/scripts/publish-release-images.sh b/.buildkite/scripts/publish-release-images.sh index 2fa334241f..dec366390b 100755 --- a/.buildkite/scripts/publish-release-images.sh +++ b/.buildkite/scripts/publish-release-images.sh @@ -61,6 +61,13 @@ if target_enabled cuda-13-0; then docker manifest create "vllm/vllm-openai:v${RELEASE_VERSION}" "vllm/vllm-openai:v${RELEASE_VERSION}-x86_64" "vllm/vllm-openai:v${RELEASE_VERSION}-aarch64" docker manifest push vllm/vllm-openai:latest docker manifest push "vllm/vllm-openai:v${RELEASE_VERSION}" + + ZSTD_DIGEST=$(docker buildx imagetools inspect \ + "public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64" \ + --format '{{json .Manifest.Digest}}' | tr -d '"') + .buildkite/scripts/publish-zstd-image.sh \ + "public.ecr.aws/q9t5s3a7/vllm-release-repo@${ZSTD_DIGEST}" \ + "vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-zstd" fi # ---- CUDA 12.9 ---- diff --git a/.buildkite/scripts/publish-zstd-image.sh b/.buildkite/scripts/publish-zstd-image.sh new file mode 100755 index 0000000000..bd3d2f99d6 --- /dev/null +++ b/.buildkite/scripts/publish-zstd-image.sh @@ -0,0 +1,45 @@ +#!/bin/bash +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project + +set -euo pipefail + +if [ "$#" -ne 2 ]; then + echo "Usage: $0 SOURCE@sha256:DIGEST DESTINATION-zstd" >&2 + exit 2 +fi + +SOURCE=$1 +DESTINATION=$2 +if [[ ! "$SOURCE" =~ ^[a-zA-Z0-9][a-zA-Z0-9._:/-]*@sha256:[a-f0-9]{64}$ ]] || + [[ ! "$DESTINATION" =~ ^[a-zA-Z0-9][a-zA-Z0-9._:/-]*:[a-zA-Z0-9_][a-zA-Z0-9_.-]*-zstd$ ]]; then + echo "ERROR: source must use an immutable sha256 digest and destination must end in -zstd" >&2 + exit 2 +fi + +BUILDER= +cleanup() { + status=$? + trap - EXIT INT TERM + if [ -n "$BUILDER" ] && ! docker buildx rm --force "$BUILDER"; then + echo "ERROR: failed to remove Buildx builder $BUILDER" >&2 + [ "$status" -ne 0 ] || status=1 + fi + exit "$status" +} +trap cleanup EXIT +trap 'exit 130' INT +trap 'exit 143' TERM + +BUILDER=$(docker buildx create \ + --driver docker-container \ + --driver-opt network=host \ + --driver-opt 'image=moby/buildkit:v0.32.2@sha256:28a898719c18a33f4e8000685287fa36fd0dd9560c6440227d3a732d79bb41d8') + +printf '# check=error=true\n\nFROM %s\n' "$SOURCE" | docker buildx build \ + --builder "$BUILDER" \ + --platform linux/amd64 \ + --progress plain \ + --provenance=mode=min \ + --output "type=image,name=${DESTINATION},push=true,compression=zstd,compression-level=3,force-compression=true,oci-mediatypes=true" \ + - diff --git a/.buildkite/test-amd.yaml b/.buildkite/test-amd.yaml index fe2cfd18c8..9a2116a16d 100644 --- a/.buildkite/test-amd.yaml +++ b/.buildkite/test-amd.yaml @@ -436,7 +436,7 @@ steps: - tests/v1/e2e/general/ - vllm/platforms/rocm.py commands: - - pytest -v -s v1/e2e/general --ignore v1/e2e/general/test_async_scheduling.py + - pytest -v -s v1/e2e/general --ignore v1/e2e/general/test_async_scheduling.py --ignore v1/e2e/general/test_kv_sharing_fast_prefill.py -k "not test_mamba_prefix_cache_mrv2" - label: ":amd: (MI250) E2E Scheduling" # TBD timeout_in_minutes: 55 @@ -1361,6 +1361,23 @@ steps: - bash ../tools/install_torchcodec_rocm.sh || exit 1 - pytest -s entrypoints/openai/correctness/ +- label: ":amd: (MI300) Scale-out EC Connector E2E" + timeout_in_minutes: 60 + mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300] + dind: false + agent_pool: mi300_2 + num_gpus: 2 + optional: true + working_dir: "/vllm-workspace/tests" + source_file_dependencies: + - vllm/distributed/ec_transfer/ + - vllm/entrypoints/scale_out/ + - tests/entrypoints/scale_out/ec_integration/ + - vllm/platforms/rocm.py + commands: + - export VLLM_WORKER_MULTIPROC_METHOD=spawn + - bash entrypoints/scale_out/ec_integration/run_scale_out_ec_e2e_test.sh + #----------------------------------------------------------- mi300 · evals -----------------------------------------------------------# - label: ":amd: (MI300) DeepSeek V2-Lite Prefetch Offload Accuracy" # TBD @@ -1737,29 +1754,6 @@ steps: - pytest -v -s tests/ir - pytest -v -s tests/kernels/ir -- label: ":amd: (MI300) MLA Kernels" # TBD - timeout_in_minutes: 25 - mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300] - dind: false - agent_pool: mi300_1 - working_dir: "/vllm-workspace/" - source_file_dependencies: - - vllm/v1/attention/backends/mla/rocm_aiter_mla.py - - vllm/v1/attention/selector.py - - vllm/platforms/rocm.py - - vllm/_aiter_ops.py - - tests/kernels/attention/test_rocm_aiter_mla_decode.py - - tests/kernels/attention/test_rocm_aiter_mla_decode_metadata.py - - tests/kernels/attention/test_rocm_aiter_mla_fp8_support.py - - tests/kernels/attention/test_rocm_aiter_mla_op_registration.py - optional: true - commands: - - amd-smi - - pytest -v -s tests/kernels/attention/test_rocm_aiter_mla_decode_metadata.py - - pytest -v -s tests/kernels/attention/test_rocm_aiter_mla_decode.py - - pytest -v -s tests/kernels/attention/test_rocm_aiter_mla_fp8_support.py - - pytest -v -s tests/kernels/attention/test_rocm_aiter_mla_op_registration.py - - label: ":amd: (MI300) Attention Kernels Shard %N" # TBD timeout_in_minutes: 90 mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300] @@ -1778,7 +1772,7 @@ steps: - vllm/platforms/rocm.py - vllm/utils/flashinfer.py commands: - - pytest -v -s kernels/attention --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT + - pytest -v -s kernels/attention --ignore=kernels/attention/test_triton_unified_attention_diffkv.py --ignore=kernels/attention/test_rocm_aiter_mla_decode.py --ignore=kernels/attention/test_rocm_aiter_mla_decode_metadata.py --ignore=kernels/attention/test_rocm_aiter_mla_fp8_support.py --ignore=kernels/attention/test_rocm_aiter_mla_op_registration.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT - label: ":amd: (MI300) MiniMax Reduce RMS Kernels" # TBD timeout_in_minutes: 45 @@ -2566,7 +2560,7 @@ steps: - vllm/platforms/rocm.py - vllm/entrypoints/generate/beam_search commands: - - pytest -v -s samplers + - pytest -v -s samplers -k "not test_beam_search_passes_multimodal_data" #------------------------------------------------------------ mi300 · misc ------------------------------------------------------------# @@ -3213,6 +3207,58 @@ steps: v1/e2e/spec_decode/draft_model/test_draft_model.py::test_draft_model_tensor_parallelism v1/e2e/spec_decode/draft_model/test_draft_model.py::test_draft_model_engine_args_tensor_parallelism +- label: ":amd: (MI300) Mooncake EC TCP E2E" + timeout_in_minutes: 60 + mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300] + dind: false + agent_pool: mi300_2 + num_gpus: 2 + optional: true + working_dir: "/vllm-workspace" + source_file_dependencies: + - vllm/distributed/ec_transfer/ + - vllm/config/ec_transfer.py + - vllm/config/multimodal.py + - vllm/config/vllm.py + - vllm/multimodal/ + - vllm/v1/core/encoder_cache_manager.py + - vllm/v1/core/sched/ + - vllm/v1/engine/ + - vllm/v1/worker/ec_connector_model_runner_mixin.py + - vllm/v1/worker/gpu_model_runner.py + - vllm/v1/worker/gpu_worker.py + - vllm/v1/worker/gpu/ec_connector.py + - vllm/v1/worker/gpu/model_runner.py + - vllm/v1/worker/mm_encoder_model_runner.py + - vllm/v1/worker/encoder_cudagraph.py + - vllm/v1/worker/encoder_cudagraph_defs.py + - vllm/v1/attention/ops/vit_attn_wrappers.py + - vllm/v1/worker/gpu/model_states/ + - vllm/v1/worker/gpu/mm/ + - tests/v1/cudagraph/test_encoder_cudagraph.py + - tests/models/multimodal/generation/test_vit_cudagraph.py + - examples/disaggregated/disaggregated_encoder/ + - tests/v1/ec_connector/ + - requirements/kv_connectors.txt + - vllm/platforms/rocm.py + - requirements/rocm.txt + commands: + - export VLLM_WORKER_MULTIPROC_METHOD=spawn + - export PYTHON_BIN='python' + - export PYTHONUNBUFFERED='1' + - export VLLM_HOST_IP='127.0.0.1' + - export VLLM_USE_V2_MODEL_RUNNER='1' + - export E_CUDAGRAPH_MM_ENCODER='1' + - export E_MM_ENCODER_ATTN_BACKEND='TRITON_ATTN' + - export MOONCAKE_EC_PROTOCOL='tcp' + - export USE_MM_PROMPTS='1' + - export SKIP_BASELINE='0' + - export CONCURRENCY='3' + - export REPEAT='2' + - export LOG_PATH='/tmp/mooncake-ec-e2e' + - export BASELINE_FILE='/tmp/mooncake-ec-e2e/baseline.json' + - bash tests/v1/ec_connector/integration/run_epd_mooncake_ec_full_pipeline.sh + - label: ":amd: (MI300) CrossLayer KV layout Distributed NixlConnector PD accuracy" # TBD timeout_in_minutes: 60 mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300] @@ -3230,7 +3276,7 @@ steps: - CROSS_LAYERS_BLOCKS=True ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh - label: ":amd: (MI300) Distributed DP Extended" # TBD - timeout_in_minutes: 60 + timeout_in_minutes: 75 mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300] dind: false agent_pool: mi300_4 @@ -3364,8 +3410,71 @@ steps: commands: - python3 benchmarks/attention_benchmarks/benchmark.py --backends ROCM_ATTN ROCM_AITER_FA ROCM_AITER_UNIFIED_ATTN --batch-specs "8q1s1k" +#---------------------------------------------------------- mi355 · compile ----------------------------------------------------------# + +- label: ":amd: (MI355) PyTorch Compilation H100 Cases" + timeout_in_minutes: 45 + mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355] + dind: false + agent_pool: mi355_1 + num_gpus: 1 + optional: true + working_dir: "/vllm-workspace/tests" + source_file_dependencies: + - vllm/__init__.py + - vllm/_aiter_ops.py + - vllm/_custom_ops.py + - vllm/compilation/ + - vllm/config/ + - vllm/distributed/ + - "!vllm/distributed/kv_transfer/" + - 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/h100/ + - tests/utils.py + - csrc/ + commands: + - pytest -v -s compile/h100/test_startup.py + #-------------------------------------------------------- mi355 · distributed --------------------------------------------------------# +- label: ":amd: (MI355) Sharded RDT Weight Transfer" + timeout_in_minutes: 45 + mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355] + dind: false + agent_pool: mi355_1 + num_gpus: 1 + optional: true + working_dir: "/vllm-workspace/tests" + source_file_dependencies: + - vllm/distributed/weight_transfer/ + - tests/distributed/test_sharded_rdt_plan.py + - tests/distributed/test_sharded_rdt_producer.py + - tests/distributed/test_sharded_rdt_trainer.py + - vllm/platforms/rocm.py + commands: + - export VLLM_WORKER_MULTIPROC_METHOD=spawn + - uv pip install --system --no-deps 'ray==2.56.1' + - pytest -v -s distributed/test_sharded_rdt_plan.py + - pytest -v -s distributed/test_sharded_rdt_producer.py + - pytest -v -s distributed/test_sharded_rdt_trainer.py + - label: ":amd: (MI355) Distributed DP Basic" # TBD timeout_in_minutes: 45 mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355] @@ -3445,6 +3554,35 @@ steps: - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt - DP_EP=1 ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh +- label: ":amd: (MI355) Kimi-Linear-48B-A3B Disaggregated DP EP" + timeout_in_minutes: 105 + mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355] + dind: false + agent_pool: mi355_4 + num_gpus: 4 + optional: true + working_dir: "/vllm-workspace/tests" + source_file_dependencies: + - vllm/platforms/rocm.py + - vllm/models/kimi_k3/ + - vllm/distributed/kv_transfer/kv_connector/v1/nixl/ + - tests/v1/kv_connector/nixl_integration/ + - requirements/kv_connectors_rocm.txt + commands: + - export VLLM_WORKER_MULTIPROC_METHOD=spawn + - export ENABLE_HMA_FLAG='1' + - export DP_EP='1' + - export GPU_MEMORY_UTILIZATION='0.9' + - export VLLM_SSM_CONV_STATE_LAYOUT='DS' + - export PREFILLER_TP_SIZE='2' + - export DECODER_TP_SIZE='2' + - export MODEL_NAMES='moonshotai/Kimi-Linear-48B-A3B-Instruct' + - export PREFILL_BLOCK_SIZE='2048' + - export DECODE_BLOCK_SIZE='2048' + - export VLLM_SERVE_EXTRA_ARGS='--trust-remote-code,--mamba-cache-mode align,--max-model-len 32768' + - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt + - bash v1/kv_connector/nixl_integration/run_accuracy_test.sh + - label: ":amd: (MI355) Distributed Features" # TBD timeout_in_minutes: 90 mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355] @@ -3642,6 +3780,51 @@ 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-gfx950.txt +- label: ":amd: (MI355) LM Eval Watermarking" + timeout_in_minutes: 45 + mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355] + dind: false + agent_pool: mi355_1 + num_gpus: 1 + optional: true + working_dir: "/vllm-workspace/tests" + source_file_dependencies: + - .buildkite/test_areas/lm_eval.yaml + - tests/evals/gsm8k/ + - vllm/config/watermarking.py + - vllm/sampling_params.py + - vllm/v1/watermarking/ + - vllm/v1/worker/gpu/ + - vllm/platforms/rocm.py + commands: + - export VLLM_WORKER_MULTIPROC_METHOD=spawn + - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-watermark.txt + +- label: ":amd: (MI355) LM Eval Watermark Feature Combination" + timeout_in_minutes: 60 + mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355] + dind: false + agent_pool: mi355_2 + num_gpus: 2 + optional: true + working_dir: "/vllm-workspace/tests" + source_file_dependencies: + - .buildkite/test_areas/lm_eval.yaml + - tests/evals/gsm8k/ + - vllm/config/speculative.py + - vllm/config/vllm.py + - vllm/config/watermarking.py + - vllm/sampling_params.py + - vllm/model_executor/models/qwen3_dspark.py + - vllm/model_executor/models/qwen3_dflash.py + - vllm/model_executor/warmup/ + - vllm/v1/watermarking/ + - vllm/v1/worker/gpu/ + - vllm/platforms/rocm.py + commands: + - export VLLM_WORKER_MULTIPROC_METHOD=spawn + - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-watermark-dspark-tp2.txt + - label: ":amd: (MI355) LM Eval Qwen3-5 Models" # TBD timeout_in_minutes: 40 mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355] @@ -3665,6 +3848,27 @@ steps: commands: - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-qwen35-mi355.txt +- label: ":amd: (MI355) Qwen3-8-Flash-Next-FP8 Accuracy Eval" + timeout_in_minutes: 105 + mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355] + dind: false + agent_pool: mi355_4 + num_gpus: 4 + optional: true + working_dir: "/vllm-workspace/tests" + source_file_dependencies: + - csrc/libtorch_stable/gdn/ + - tests/evals/qwen4_exp/ + - vllm/model_executor/layers/mamba/ + - vllm/models/qwen4_exp/ + - vllm/transformers_utils/configs/qwen4_exp.py + - vllm/v1/attention/backends/short_conv_attn.py + - vllm/platforms/rocm.py + commands: + - export VLLM_WORKER_MULTIPROC_METHOD=spawn + - uv pip install --system 'evalscope==1.10.0' + - pytest -s -v evals/qwen4_exp/test_accuracy.py --config-list-file=configs/models-rocm.txt + - label: ":amd: (MI355) LM Eval Small Models Distributed" timeout_in_minutes: 55 mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355] @@ -3686,6 +3890,25 @@ steps: commands: - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-mi3xx-fp8-and-mixed.txt +- label: ":amd: (MI355) MoE Refactor Integration Shard %N" + timeout_in_minutes: 105 + mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355] + dind: false + agent_pool: mi355_2 + num_gpus: 2 + parallelism: 4 + optional: true + working_dir: "/vllm-workspace/tests" + source_file_dependencies: + - vllm/platforms/rocm.py + - tests/evals/gsm8k/ + - vllm/model_executor/layers/fused_moe/ + - vllm/model_executor/layers/quantization/ + - vllm/_aiter_ops.py + commands: + - export VLLM_WORKER_MULTIPROC_METHOD=spawn + - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/moe-refactor/config-rocm-shard-$$BUILDKITE_PARALLEL_JOB.txt + - label: ":amd: (MI355) Qwen3-30B-A3B-FP8-block Sync EPLB Accuracy" # TBD timeout_in_minutes: 25 mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355] @@ -3790,6 +4013,23 @@ steps: #---------------------------------------------------------- mi355 · kernels ----------------------------------------------------------# +- label: ":amd: (MI355) Attention DiffKV Kernels" + timeout_in_minutes: 35 + mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355] + dind: false + agent_pool: mi355_1 + num_gpus: 1 + optional: true + working_dir: "/vllm-workspace/tests" + source_file_dependencies: + - vllm/v1/attention/ops/triton_unified_attention_diffkv.py + - vllm/v1/attention/backends/triton_attn_diffkv.py + - tests/kernels/attention/test_triton_unified_attention.py + - tests/kernels/attention/test_triton_unified_attention_diffkv.py + - vllm/platforms/rocm.py + commands: + - pytest -v -s kernels/attention/test_triton_unified_attention_diffkv.py + - label: ":amd: (MI355) Core Operation Kernels Shard %N" # TBD timeout_in_minutes: 40 mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355] @@ -3859,19 +4099,19 @@ steps: - vllm/_aiter_ops.py - examples/basic/offline_inference/chat.py - tests/kernels/attention/test_attention_selector.py - - tests/kernels/attention/test_rocm_aiter_mla_decode_metadata.py optional: true commands: - amd-smi - python3 examples/basic/offline_inference/chat.py --attention-backend TRITON_ATTN - pytest -v -s tests/kernels/attention/test_attention_selector.py - - pytest -v -s tests/kernels/attention/test_rocm_aiter_mla_decode_metadata.py -- label: ":amd: (MI355) MLA Kernels" # TBD - timeout_in_minutes: 25 +- label: ":amd: (MI355) MLA Kernels" + timeout_in_minutes: 40 mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355] dind: false agent_pool: mi355_1 + num_gpus: 1 + optional: true working_dir: "/vllm-workspace/" source_file_dependencies: - vllm/v1/attention/backends/mla/rocm_aiter_mla.py @@ -3882,7 +4122,6 @@ steps: - tests/kernels/attention/test_rocm_aiter_mla_decode_metadata.py - tests/kernels/attention/test_rocm_aiter_mla_fp8_support.py - tests/kernels/attention/test_rocm_aiter_mla_op_registration.py - optional: true commands: - amd-smi - pytest -v -s tests/kernels/attention/test_rocm_aiter_mla_decode_metadata.py @@ -3890,6 +4129,22 @@ steps: - pytest -v -s tests/kernels/attention/test_rocm_aiter_mla_fp8_support.py - pytest -v -s tests/kernels/attention/test_rocm_aiter_mla_op_registration.py +- label: ":amd: (MI355) DeepSeek V4 Fused Kernels" + timeout_in_minutes: 60 + mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355] + dind: false + agent_pool: mi355_1 + num_gpus: 1 + optional: true + working_dir: "/vllm-workspace/tests" + source_file_dependencies: + - csrc/libtorch_stable/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu + - vllm/models/deepseek_v4/common/ops/ + - tests/kernels/test_fused_deepseek_v4_qnorm_rope_kv_insert.py + - vllm/platforms/rocm.py + commands: + - pytest -v -s kernels/test_fused_deepseek_v4_qnorm_rope_kv_insert.py + - label: ":amd: (MI355) Attention Kernels Shard %N" # TBD timeout_in_minutes: 90 mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355] @@ -3908,7 +4163,7 @@ steps: - vllm/utils/flashinfer.py optional: true commands: - - pytest -v -s kernels/attention --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT + - pytest -v -s kernels/attention --ignore=kernels/attention/test_triton_unified_attention_diffkv.py --ignore=kernels/attention/test_rocm_aiter_mla_decode.py --ignore=kernels/attention/test_rocm_aiter_mla_decode_metadata.py --ignore=kernels/attention/test_rocm_aiter_mla_fp8_support.py --ignore=kernels/attention/test_rocm_aiter_mla_op_registration.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT - label: ":amd: (MI355) MoE Kernels Shard %N" # TBD timeout_in_minutes: 55 @@ -4050,6 +4305,21 @@ steps: #------------------------------------------------------ mi355 · models / basic -------------------------------------------------------# +- label: ":amd: (MI355) GLM5Next" + timeout_in_minutes: 35 + mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355] + dind: false + agent_pool: mi355_1 + num_gpus: 1 + optional: true + working_dir: "/vllm-workspace/tests" + source_file_dependencies: + - vllm/models/glm5next/ + - tests/models/glm5next/ + - vllm/platforms/rocm.py + commands: + - pytest -v -s models/glm5next + - label: ":amd: (MI355) Basic Models (Extra Initialization) Shard %N" # TBD timeout_in_minutes: 60 mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355] @@ -4087,7 +4357,27 @@ steps: - export GIT_CONFIG_VALUE_0=HTTP/1.1 - MAMBA_FORCE_BUILD=TRUE uv pip install --system --no-build-isolation 'git+https://github.com/AndreasKaratzas/mamba@fix-rocm-7.0-warp-size-constexpr' - CAUSAL_CONV1D_FORCE_BUILD=TRUE 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 + - 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 + +- label: ":amd: (MI355) Granite Language Model Compatibility" + timeout_in_minutes: 80 + mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355] + dind: false + agent_pool: mi355_1 + num_gpus: 1 + optional: true + working_dir: "/vllm-workspace/tests" + source_file_dependencies: + - vllm/ + - "!vllm/distributed/kv_transfer/" + - tests/models/language/generation + commands: + - export GIT_CONFIG_COUNT=1 + - export GIT_CONFIG_KEY_0=http.version + - export GIT_CONFIG_VALUE_0=HTTP/1.1 + - MAMBA_FORCE_BUILD=TRUE uv pip install --system --no-build-isolation 'git+https://github.com/AndreasKaratzas/mamba@fix-rocm-7.0-warp-size-constexpr' + - CAUSAL_CONV1D_FORCE_BUILD=TRUE 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: ":amd: (MI355) Language Models (Extended Generation)" # TBD timeout_in_minutes: 95 @@ -4308,8 +4598,65 @@ steps: # commands: # - pytest -s -v tests/quantization/test_gfx950_moe.py +#--------------------------------------------------------- mi355 · samplers ----------------------------------------------------------# + +- label: ":amd: (MI355) Samplers Multimodal Beam Search" + timeout_in_minutes: 45 + mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355] + dind: false + agent_pool: mi355_1 + num_gpus: 1 + optional: true + working_dir: "/vllm-workspace/tests" + source_file_dependencies: + - vllm/model_executor/layers + - vllm/sampling_metadata.py + - vllm/v1/sample/ + - vllm/entrypoints/generate/beam_search/ + - tests/samplers + - tests/conftest.py + - vllm/_aiter_ops.py + - vllm/platforms/rocm.py + commands: + - pytest -v -s samplers/test_beam_search.py::test_beam_search_passes_multimodal_data + #------------------------------------------------------------ mi355 · v1 -------------------------------------------------------------# +- label: ":amd: (MI355) E2E Core Large Memory" + timeout_in_minutes: 65 + mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355] + dind: false + agent_pool: mi355_1 + num_gpus: 1 + optional: true + working_dir: "/vllm-workspace/tests" + source_file_dependencies: + - vllm/v1/ + - tests/v1/e2e/general/ + - vllm/platforms/rocm.py + commands: + - pytest -v -s v1/e2e/general/test_kv_sharing_fast_prefill.py + - pytest -v -s v1/e2e/general/test_mamba_prefix_cache.py::test_mamba_prefix_cache_mrv2 v1/e2e/general/test_mamba_prefix_cache.py::test_mamba_prefix_cache_mrv2_async + +- label: ":amd: (MI355) Extract Hidden States Integration Single GPU" + timeout_in_minutes: 45 + mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355] + dind: false + agent_pool: mi355_1 + num_gpus: 1 + optional: true + working_dir: "/vllm-workspace" + source_file_dependencies: + - vllm/v1/spec_decode/extract_hidden_states.py + - vllm/model_executor/models/extract_hidden_states.py + - vllm/transformers_utils/configs/extract_hidden_states.py + - vllm/distributed/kv_transfer/kv_connector/v1/example_hidden_states_connector.py + - tests/v1/kv_connector/extract_hidden_states_integration + - vllm/platforms/rocm.py + commands: + - export VLLM_WORKER_MULTIPROC_METHOD=spawn + - pytest -v -s -m 'not distributed' tests/v1/kv_connector/extract_hidden_states_integration + - label: ":amd: (MI355) V1 Attention Shard %N" # TBD timeout_in_minutes: 125 mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355] diff --git a/.buildkite/test_areas/disaggregated.yaml b/.buildkite/test_areas/disaggregated.yaml index b187169e76..811301f58d 100644 --- a/.buildkite/test_areas/disaggregated.yaml +++ b/.buildkite/test_areas/disaggregated.yaml @@ -318,3 +318,24 @@ steps: commands: - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh - bash v1/kv_connector/nixl_integration/run_accuracy_test.sh + mirror: + amd: + label: ":amd: (MI355) Kimi-Linear-48B-A3B Disaggregated DP EP" + dind: false + device: mi355_4 + num_devices: 4 + timeout_in_minutes: 105 + working_dir: /vllm-workspace/tests + depends_on: + - image-build-amd + source_file_dependencies: + - vllm/platforms/rocm.py + - vllm/models/kimi_k3/ + - vllm/distributed/kv_transfer/kv_connector/v1/nixl/ + - tests/v1/kv_connector/nixl_integration/ + - requirements/kv_connectors_rocm.txt + env: + VLLM_WORKER_MULTIPROC_METHOD: "spawn" + commands: + - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt + - bash v1/kv_connector/nixl_integration/run_accuracy_test.sh diff --git a/.buildkite/test_areas/disaggregated_ec.yaml b/.buildkite/test_areas/disaggregated_ec.yaml index 9100ad0a43..67bfebaf8c 100644 --- a/.buildkite/test_areas/disaggregated_ec.yaml +++ b/.buildkite/test_areas/disaggregated_ec.yaml @@ -14,3 +14,22 @@ steps: - tests/entrypoints/scale_out/ec_integration/ commands: - bash entrypoints/scale_out/ec_integration/run_scale_out_ec_e2e_test.sh + mirror: + amd: + label: ":amd: (MI300) Scale-out EC Connector E2E" + dind: false + device: mi300_2 + num_devices: 2 + timeout_in_minutes: 60 + working_dir: /vllm-workspace/tests + depends_on: + - image-build-amd + source_file_dependencies: + - vllm/distributed/ec_transfer/ + - vllm/entrypoints/scale_out/ + - tests/entrypoints/scale_out/ec_integration/ + - vllm/platforms/rocm.py + env: + VLLM_WORKER_MULTIPROC_METHOD: "spawn" + commands: + - bash entrypoints/scale_out/ec_integration/run_scale_out_ec_e2e_test.sh diff --git a/.buildkite/test_areas/disaggregated_mooncake.yaml b/.buildkite/test_areas/disaggregated_mooncake.yaml index f6e944d705..d89996c9a2 100644 --- a/.buildkite/test_areas/disaggregated_mooncake.yaml +++ b/.buildkite/test_areas/disaggregated_mooncake.yaml @@ -83,3 +83,45 @@ steps: PY )" - bash tests/v1/ec_connector/integration/run_epd_mooncake_ec_full_pipeline.sh + mirror: + amd: + label: ":amd: (MI300) Mooncake EC TCP E2E" + dind: false + device: mi300_2 + num_devices: 2 + timeout_in_minutes: 60 + working_dir: /vllm-workspace + depends_on: + - image-build-amd + source_file_dependencies: + - vllm/distributed/ec_transfer/ + - vllm/config/ec_transfer.py + - vllm/config/multimodal.py + - vllm/config/vllm.py + - vllm/multimodal/ + - vllm/v1/core/encoder_cache_manager.py + - vllm/v1/core/sched/ + - vllm/v1/engine/ + - vllm/v1/worker/ec_connector_model_runner_mixin.py + - vllm/v1/worker/gpu_model_runner.py + - vllm/v1/worker/gpu_worker.py + - vllm/v1/worker/gpu/ec_connector.py + - vllm/v1/worker/gpu/model_runner.py + - vllm/v1/worker/mm_encoder_model_runner.py + - vllm/v1/worker/encoder_cudagraph.py + - vllm/v1/worker/encoder_cudagraph_defs.py + - vllm/v1/attention/ops/vit_attn_wrappers.py + - vllm/v1/worker/gpu/model_states/ + - vllm/v1/worker/gpu/mm/ + - tests/v1/cudagraph/test_encoder_cudagraph.py + - tests/models/multimodal/generation/test_vit_cudagraph.py + - examples/disaggregated/disaggregated_encoder/ + - tests/v1/ec_connector/ + - requirements/kv_connectors.txt + - vllm/platforms/rocm.py + - requirements/rocm.txt + env: + VLLM_WORKER_MULTIPROC_METHOD: "spawn" + PYTHON_BIN: "python" + commands: + - bash tests/v1/ec_connector/integration/run_epd_mooncake_ec_full_pipeline.sh diff --git a/.buildkite/test_areas/distributed.yaml b/.buildkite/test_areas/distributed.yaml index a108246563..1f2d572922 100644 --- a/.buildkite/test_areas/distributed.yaml +++ b/.buildkite/test_areas/distributed.yaml @@ -33,6 +33,29 @@ steps: - pytest -v -s distributed/test_sharded_rdt_plan.py - pytest -v -s distributed/test_sharded_rdt_producer.py - pytest -v -s distributed/test_sharded_rdt_trainer.py + mirror: + amd: + label: ":amd: (MI355) Sharded RDT Weight Transfer" + dind: false + device: mi355_1 + num_devices: 1 + timeout_in_minutes: 45 + working_dir: /vllm-workspace/tests + depends_on: + - image-build-amd + source_file_dependencies: + - vllm/distributed/weight_transfer/ + - tests/distributed/test_sharded_rdt_plan.py + - tests/distributed/test_sharded_rdt_producer.py + - tests/distributed/test_sharded_rdt_trainer.py + - vllm/platforms/rocm.py + env: + VLLM_WORKER_MULTIPROC_METHOD: "spawn" + commands: + - uv pip install --system --no-deps 'ray==2.56.1' + - pytest -v -s distributed/test_sharded_rdt_plan.py + - pytest -v -s distributed/test_sharded_rdt_producer.py + - pytest -v -s distributed/test_sharded_rdt_trainer.py - label: ":nvidia: (L4) Distributed DP Basic" key: distributed-dp-tests-2-gpus @@ -364,17 +387,18 @@ steps: - pytest -v -s tests/v1/distributed/test_dbo.py - pytest -v -s tests/distributed/test_mnnvl_alltoall.py -- label: ":nvidia: (B200) Kimi-K3 GEMM-RS/AR" - key: kimi-k3-gemm-rs-ar-2xb200 +- label: ":nvidia: (B200) GEMM-RS/AR" + key: gemm-rs-ar-2xb200 timeout_in_minutes: 20 device: b200-k8s working_dir: "/vllm-workspace/" num_devices: 2 source_file_dependencies: - - vllm/models/kimi_k3/nvidia/ops/cute_dsl/gemm_rs_ar.py - - tests/kernels/test_kimi_k3_gemm_rs_ar.py + - vllm/model_executor/kernels/linear/cute_dsl/gemm_rs_ar.py + - vllm/cute_utils/ + - tests/kernels/test_gemm_rs_ar.py commands: - - pytest -v -s tests/kernels/test_kimi_k3_gemm_rs_ar.py + - pytest -v -s tests/kernels/test_gemm_rs_ar.py - label: ":nvidia: (L4) Distributed 2-Node" key: 2-node-test-4-gpus diff --git a/.buildkite/test_areas/engine.yaml b/.buildkite/test_areas/engine.yaml index 8e47a84e18..e38cd92573 100644 --- a/.buildkite/test_areas/engine.yaml +++ b/.buildkite/test_areas/engine.yaml @@ -94,7 +94,7 @@ steps: label: ":amd: (MI250) E2E Core" dind: false commands: - - pytest -v -s v1/e2e/general --ignore v1/e2e/general/test_async_scheduling.py + - pytest -v -s v1/e2e/general --ignore v1/e2e/general/test_async_scheduling.py --ignore v1/e2e/general/test_kv_sharing_fast_prefill.py -k "not test_mamba_prefix_cache_mrv2" device: mi250_1 timeout_in_minutes: 50 depends_on: @@ -115,6 +115,23 @@ steps: commands: - pytest -v -s v1/e2e/general/test_kv_sharing_fast_prefill.py - pytest -v -s v1/e2e/general/test_mamba_prefix_cache.py::test_mamba_prefix_cache_mrv2 v1/e2e/general/test_mamba_prefix_cache.py::test_mamba_prefix_cache_mrv2_async + mirror: + amd: + label: ":amd: (MI355) E2E Core Large Memory" + dind: false + device: mi355_1 + num_devices: 1 + timeout_in_minutes: 65 + working_dir: /vllm-workspace/tests + depends_on: + - image-build-amd + source_file_dependencies: + - vllm/v1/ + - tests/v1/e2e/general/ + - vllm/platforms/rocm.py + commands: + - pytest -v -s v1/e2e/general/test_kv_sharing_fast_prefill.py + - pytest -v -s v1/e2e/general/test_mamba_prefix_cache.py::test_mamba_prefix_cache_mrv2 v1/e2e/general/test_mamba_prefix_cache.py::test_mamba_prefix_cache_mrv2_async - label: ":nvidia: (L4) V1 E2E" key: v1-e2e-2-gpus diff --git a/.buildkite/test_areas/entrypoints.yaml b/.buildkite/test_areas/entrypoints.yaml index 60ac5d932c..9e0a8399a3 100644 --- a/.buildkite/test_areas/entrypoints.yaml +++ b/.buildkite/test_areas/entrypoints.yaml @@ -2,6 +2,26 @@ group: Entrypoints depends_on: - image-build steps: +- label: ":nvidia: (H200) Initialized Snapshot E2E" + key: initialized-snapshot-e2e + device: h200 + num_devices: 1 + no_plugin: true + optional: true + timeout_in_minutes: 60 + working_dir: "." + env: + SNAPSHOT_E2E_RUN_LIMIT_S: "1200" + source_file_dependencies: + - .buildkite/scripts/initialized-snapshot-e2e.sh + - .buildkite/test_areas/entrypoints.yaml + - docker/Dockerfile + - tools/install_snapshot_runtime.sh + - vllm/entrypoints/cli/ + - vllm/snapshot/ + commands: + - bash .buildkite/scripts/initialized-snapshot-e2e.sh "$IMAGE_TAG" "$BUILDKITE_COMMIT" + - label: ":nvidia: (H200 MIG 35GB) Entrypoints Unit" device: h200_35gb key: entrypoints-unit-tests diff --git a/.buildkite/test_areas/kernels.yaml b/.buildkite/test_areas/kernels.yaml index aab6e52b8e..78d57ca455 100644 --- a/.buildkite/test_areas/kernels.yaml +++ b/.buildkite/test_areas/kernels.yaml @@ -111,6 +111,23 @@ steps: commands: - pytest -v -s kernels/test_fused_deepseek_v4_*.py - pytest -v -s kernels/test_top_k_per_row.py + mirror: + amd: + label: ":amd: (MI355) DeepSeek V4 Fused Kernels" + dind: false + device: mi355_1 + num_devices: 1 + timeout_in_minutes: 60 + working_dir: /vllm-workspace/tests + depends_on: + - image-build-amd + source_file_dependencies: + - csrc/libtorch_stable/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu + - vllm/models/deepseek_v4/common/ops/ + - tests/kernels/test_fused_deepseek_v4_qnorm_rope_kv_insert.py + - vllm/platforms/rocm.py + commands: + - pytest -v -s kernels/test_fused_deepseek_v4_qnorm_rope_kv_insert.py - label: ":nvidia: (B200) DeepSeek V4 Kernels" key: deepseek-v4-kernel-test-b200 @@ -196,6 +213,8 @@ steps: - vllm/_aiter_ops.py - vllm/envs.py - vllm/platforms/rocm.py + commands: + - pytest -v -s kernels/attention --ignore=kernels/attention/test_triton_unified_attention_diffkv.py --ignore=kernels/attention/test_rocm_aiter_mla_decode.py --ignore=kernels/attention/test_rocm_aiter_mla_decode_metadata.py --ignore=kernels/attention/test_rocm_aiter_mla_fp8_support.py --ignore=kernels/attention/test_rocm_aiter_mla_op_registration.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT - label: ":nvidia: (H100) Attention DiffKV Kernels" key: kernels-attention-diffkv-test-h100 @@ -209,6 +228,24 @@ steps: - tests/kernels/attention/test_triton_unified_attention_diffkv.py commands: - pytest -v -s kernels/attention/test_triton_unified_attention_diffkv.py + mirror: + amd: + label: ":amd: (MI355) Attention DiffKV Kernels" + dind: false + device: mi355_1 + num_devices: 1 + timeout_in_minutes: 35 + working_dir: /vllm-workspace/tests + depends_on: + - image-build-amd + source_file_dependencies: + - vllm/v1/attention/ops/triton_unified_attention_diffkv.py + - vllm/v1/attention/backends/triton_attn_diffkv.py + - tests/kernels/attention/test_triton_unified_attention.py + - tests/kernels/attention/test_triton_unified_attention_diffkv.py + - vllm/platforms/rocm.py + commands: + - pytest -v -s kernels/attention/test_triton_unified_attention_diffkv.py - label: ":nvidia: (H100) MLA Kernel Test" key: kernels-flashmla-test-h100 @@ -231,6 +268,31 @@ steps: - pytest -v -s kernels/attention/test_flashmla_sparse.py - pytest -v -s kernels/attention/test_flashinfer_mla_decode.py - pytest -v -s kernels/attention/test_mla_cross_layer_kernel_equivalence.py + mirror: + amd: + label: ":amd: (MI355) MLA Kernels" + dind: false + device: mi355_1 + num_devices: 1 + timeout_in_minutes: 40 + working_dir: /vllm-workspace/ + depends_on: + - image-build-amd + source_file_dependencies: + - vllm/v1/attention/backends/mla/rocm_aiter_mla.py + - vllm/v1/attention/selector.py + - vllm/platforms/rocm.py + - vllm/_aiter_ops.py + - tests/kernels/attention/test_rocm_aiter_mla_decode.py + - tests/kernels/attention/test_rocm_aiter_mla_decode_metadata.py + - tests/kernels/attention/test_rocm_aiter_mla_fp8_support.py + - tests/kernels/attention/test_rocm_aiter_mla_op_registration.py + commands: + - amd-smi + - pytest -v -s tests/kernels/attention/test_rocm_aiter_mla_decode_metadata.py + - pytest -v -s tests/kernels/attention/test_rocm_aiter_mla_decode.py + - pytest -v -s tests/kernels/attention/test_rocm_aiter_mla_fp8_support.py + - pytest -v -s tests/kernels/attention/test_rocm_aiter_mla_op_registration.py - label: ":nvidia: (L4) Quantization Kernels Shard %N" key: kernels-quantization-test @@ -640,31 +702,6 @@ steps: - pytest -v -s kernels/moe/test_moe_layer.py - pytest -v -s kernels/moe/test_deepep_v2_moe.py -- label: ":amd: (MI355) MLA Kernels" - key: amd-mla-kernels-mi355 - dind: false - device: mi355_1 - num_devices: 1 - timeout_in_minutes: 25 - working_dir: /vllm-workspace/ - depends_on: - - image-build-amd - source_file_dependencies: - - vllm/v1/attention/backends/mla/rocm_aiter_mla.py - - vllm/v1/attention/selector.py - - vllm/platforms/rocm.py - - vllm/_aiter_ops.py - - tests/kernels/attention/test_rocm_aiter_mla_decode.py - - tests/kernels/attention/test_rocm_aiter_mla_decode_metadata.py - - tests/kernels/attention/test_rocm_aiter_mla_fp8_support.py - - tests/kernels/attention/test_rocm_aiter_mla_op_registration.py - commands: - - amd-smi - - pytest -v -s tests/kernels/attention/test_rocm_aiter_mla_decode_metadata.py - - pytest -v -s tests/kernels/attention/test_rocm_aiter_mla_decode.py - - pytest -v -s tests/kernels/attention/test_rocm_aiter_mla_fp8_support.py - - pytest -v -s tests/kernels/attention/test_rocm_aiter_mla_op_registration.py - - label: ":amd: (MI355) FP8 MoE Kernels" key: amd-fp8-moe-kernels-mi355 dind: false @@ -720,9 +757,7 @@ steps: - vllm/_aiter_ops.py - examples/basic/offline_inference/chat.py - tests/kernels/attention/test_attention_selector.py - - tests/kernels/attention/test_rocm_aiter_mla_decode_metadata.py commands: - amd-smi - python3 examples/basic/offline_inference/chat.py --attention-backend TRITON_ATTN - pytest -v -s tests/kernels/attention/test_attention_selector.py - - pytest -v -s tests/kernels/attention/test_rocm_aiter_mla_decode_metadata.py diff --git a/.buildkite/test_areas/lm_eval.yaml b/.buildkite/test_areas/lm_eval.yaml index 3eb7f0a444..3450b70c28 100644 --- a/.buildkite/test_areas/lm_eval.yaml +++ b/.buildkite/test_areas/lm_eval.yaml @@ -66,6 +66,28 @@ steps: autorun_on_main: true commands: - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-watermark.txt + mirror: + amd: + label: ":amd: (MI355) LM Eval Watermarking" + dind: false + device: mi355_1 + num_devices: 1 + timeout_in_minutes: 45 + working_dir: /vllm-workspace/tests + depends_on: + - image-build-amd + source_file_dependencies: + - .buildkite/test_areas/lm_eval.yaml + - tests/evals/gsm8k/ + - vllm/config/watermarking.py + - vllm/sampling_params.py + - vllm/v1/watermarking/ + - vllm/v1/worker/gpu/ + - vllm/platforms/rocm.py + env: + VLLM_WORKER_MULTIPROC_METHOD: "spawn" + commands: + - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-watermark.txt # - label: LM Eval Large Models (4xA100) # key: lm-eval-large-models-4xa100 @@ -203,6 +225,33 @@ steps: autorun_on_main: true commands: - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-watermark-dspark-tp2.txt + mirror: + amd: + label: ":amd: (MI355) LM Eval Watermark Feature Combination" + dind: false + device: mi355_2 + num_devices: 2 + timeout_in_minutes: 60 + working_dir: /vllm-workspace/tests + depends_on: + - image-build-amd + source_file_dependencies: + - .buildkite/test_areas/lm_eval.yaml + - tests/evals/gsm8k/ + - vllm/config/speculative.py + - vllm/config/vllm.py + - vllm/config/watermarking.py + - vllm/sampling_params.py + - vllm/model_executor/models/qwen3_dspark.py + - vllm/model_executor/models/qwen3_dflash.py + - vllm/model_executor/warmup/ + - vllm/v1/watermarking/ + - vllm/v1/worker/gpu/ + - vllm/platforms/rocm.py + env: + VLLM_WORKER_MULTIPROC_METHOD: "spawn" + commands: + - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-watermark-dspark-tp2.txt - label: ":nvidia: (B200) LM Eval Spec Decode" key: lm-eval-spec-decode-4xb200 @@ -294,29 +343,66 @@ steps: commands: - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-qwen35-blackwell.txt -- label: ":nvidia: (H200) LM Eval Large Models" - key: lm-eval-large-models-8xh200 - timeout_in_minutes: 120 +- label: ":nvidia: (H200) LM Eval DeepSeek R1 TP" + key: lm-eval-deepseek-r1-tp-8xh200 + timeout_in_minutes: 35 device: h200 optional: true num_devices: 8 commands: - - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-h200.txt - mirror: - amd: - label: ":amd: (MI300) LM Eval Large Models" - dind: false - device: mi300_8 - timeout_in_minutes: 65 - depends_on: - - image-build-amd - commands: - - export VLLM_WORKER_MULTIPROC_METHOD=spawn - - export PYTORCH_ROCM_ARCH=gfx942 # Limit Quark compilation to save time - - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-mi3xx.txt - source_file_dependencies: - - vllm/_aiter_ops.py - - vllm/platforms/rocm.py + - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-h200-deepseek-r1-tp.txt + +- label: ":nvidia: (H200) LM Eval DeepSeek R1 DP" + key: lm-eval-deepseek-r1-dp-8xh200 + timeout_in_minutes: 35 + device: h200 + optional: true + num_devices: 8 + commands: + - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-h200-deepseek-r1-dp.txt + +- label: ":nvidia: (H200) LM Eval DeepSeek V3.2 TP" + key: lm-eval-deepseek-v32-tp-8xh200 + timeout_in_minutes: 35 + device: h200 + optional: true + num_devices: 8 + commands: + - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-h200-deepseek-v32-tp.txt + +- label: ":nvidia: (H200) LM Eval DeepSeek V3.2 DP" + key: lm-eval-deepseek-v32-dp-8xh200 + timeout_in_minutes: 35 + device: h200 + optional: true + num_devices: 8 + commands: + - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-h200-deepseek-v32-dp.txt + +- label: ":nvidia: (H200) LM Eval Nemotron 3 Super" + key: lm-eval-nemotron-3-super-8xh200 + timeout_in_minutes: 35 + device: h200 + optional: true + num_devices: 8 + commands: + - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-h200-nemotron-3-super.txt + +- label: ":amd: (MI300) LM Eval Large Models" + key: amd-lm-eval-large-models + dind: false + device: mi300_8 + optional: true + timeout_in_minutes: 65 + depends_on: + - image-build-amd + commands: + - export VLLM_WORKER_MULTIPROC_METHOD=spawn + - export PYTORCH_ROCM_ARCH=gfx942 # Limit Quark compilation to save time + - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-mi3xx.txt + source_file_dependencies: + - vllm/_aiter_ops.py + - vllm/platforms/rocm.py - label: ":nvidia: (H100) MoE Refactor Integration TEMPORARY" key: moe-refactor-integration-test-h100-temporary @@ -329,7 +415,27 @@ steps: # build 83851): shards ~12.4/13.6/12.4/14.5m of eval time. New configs must # be added to exactly one config-h100-shard-*.txt. - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/moe-refactor/config-h100-shard-$$BUILDKITE_PARALLEL_JOB.txt - + mirror: + amd: + label: ":amd: (MI355) MoE Refactor Integration Shard %N" + dind: false + device: mi355_2 + num_devices: 2 + timeout_in_minutes: 105 + working_dir: /vllm-workspace/tests + depends_on: + - image-build-amd + source_file_dependencies: + - vllm/platforms/rocm.py + - tests/evals/gsm8k/ + - vllm/model_executor/layers/fused_moe/ + - vllm/model_executor/layers/quantization/ + - vllm/_aiter_ops.py + env: + VLLM_WORKER_MULTIPROC_METHOD: "spawn" + commands: + - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/moe-refactor/config-rocm-shard-$$BUILDKITE_PARALLEL_JOB.txt + - label: ":nvidia: (B200) MoE Refactor Integration TEMPORARY Shard %N" key: moe-refactor-integration-test-b200-temporary device: b200-k8s @@ -347,47 +453,67 @@ 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: ":nvidia: (A100) LM Eval Humming FP16 TEMPORARY Shard %N" - key: lm-eval-humming-f16-a100 +- label: ":nvidia: (L4) LM Eval Humming TEMPORARY" + key: lm-eval-humming-l4 + optional: true + timeout_in_minutes: 90 + device: l4 + num_devices: 1 + source_file_dependencies: + - .buildkite/test_areas/lm_eval.yaml + - tests/evals/gsm8k/configs/humming/ + - requirements/cuda.txt + - vllm/utils/humming.py + - vllm/model_executor/layers/quantization/humming.py + - vllm/model_executor/layers/quantization/utils/humming/ + - 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-l4.txt + +- label: ":nvidia: (A100) LM Eval Humming Default TEMPORARY Shard %N" + key: lm-eval-humming-a100-default + optional: true timeout_in_minutes: 75 device: a100 - optional: true num_devices: 1 parallelism: 3 source_file_dependencies: - vllm/model_executor/layers/quantization/humming.py - - vllm/model_executor/layers/quantization/utils/humming_utils.py + - vllm/model_executor/layers/quantization/utils/humming/ - 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-a100-shard-$$BUILDKITE_PARALLEL_JOB.txt -- label: ":nvidia: (A100) LM Eval Humming Activation INT8 TEMPORARY" - key: lm-eval-humming-act-a100 +- label: ":nvidia: (A100) LM Eval Humming Activation INT8 TEMPORARY Shard %N" + key: lm-eval-humming-a100-act + optional: true timeout_in_minutes: 45 device: a100 - optional: true num_devices: 1 + parallelism: 2 source_file_dependencies: - vllm/model_executor/layers/quantization/humming.py - - vllm/model_executor/layers/quantization/utils/humming_utils.py + - vllm/model_executor/layers/quantization/utils/humming/ - 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 + - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config-a100-act-int8-shard-$$BUILDKITE_PARALLEL_JOB.txt -- label: ":nvidia: (H100) LM Eval Humming FP16 TEMPORARY Shard %N" - key: lm-eval-humming-f16-h100 +- label: ":nvidia: (H100) LM Eval Humming Default TEMPORARY Shard %N" + key: lm-eval-humming-h100-default + optional: true timeout_in_minutes: 70 device: h100 - optional: true num_devices: 1 parallelism: 3 source_file_dependencies: - vllm/model_executor/layers/quantization/humming.py - - vllm/model_executor/layers/quantization/utils/humming_utils.py + - vllm/model_executor/layers/quantization/utils/humming/ - vllm/model_executor/layers/fused_moe/experts/fused_humming_moe.py - vllm/model_executor/layers/fused_moe/oracle/ - vllm/model_executor/kernels/linear/ @@ -395,14 +521,14 @@ steps: - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config-h100-shard-$$BUILDKITE_PARALLEL_JOB.txt - label: ":nvidia: (H100) LM Eval Humming Activation FP8/INT8 TEMPORARY" - key: lm-eval-humming-act-h100 + key: lm-eval-humming-h100-act + optional: true 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/quantization/utils/humming/ - vllm/model_executor/layers/fused_moe/experts/fused_humming_moe.py - vllm/model_executor/layers/fused_moe/oracle/ - vllm/model_executor/kernels/linear/ @@ -410,15 +536,35 @@ steps: - 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: ":nvidia: (B200) LM Eval Humming FP16 TEMPORARY" - key: lm-eval-humming-f16-b200 +- label: ":nvidia: (H100) LM Eval Humming DeepEP TEMPORARY" + key: lm-eval-humming-h100-deepep + optional: true + timeout_in_minutes: 60 + device: h100 + num_devices: 2 + source_file_dependencies: + - .buildkite/test_areas/lm_eval.yaml + - tests/evals/gsm8k/configs/humming/ + - requirements/cuda.txt + - vllm/utils/humming.py + - vllm/model_executor/layers/quantization/humming.py + - vllm/model_executor/layers/quantization/utils/humming/ + - vllm/model_executor/layers/fused_moe/ + - vllm/model_executor/kernels/linear/ + - vllm/distributed/ + - csrc/moe/ + commands: + - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config-deepep.txt + +- label: ":nvidia: (B200) LM Eval Humming Default TEMPORARY" + key: lm-eval-humming-b200-default + optional: true 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/quantization/utils/humming/ - vllm/model_executor/layers/fused_moe/experts/fused_humming_moe.py - vllm/model_executor/layers/fused_moe/oracle/ - vllm/model_executor/kernels/linear/ @@ -426,14 +572,14 @@ steps: - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config.txt - label: ":nvidia: (B200) LM Eval Humming Activation FP8/INT8 TEMPORARY" - key: lm-eval-humming-act-b200 + key: lm-eval-humming-b200-act + optional: true 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/quantization/utils/humming/ - vllm/model_executor/layers/fused_moe/experts/fused_humming_moe.py - vllm/model_executor/layers/fused_moe/oracle/ - vllm/model_executor/kernels/linear/ @@ -538,6 +684,29 @@ steps: commands: - uv pip install --system 'evalscope==1.10.0' - pytest -s -v evals/qwen4_exp/test_accuracy.py --config-list-file=configs/models-h200.txt + mirror: + amd: + label: ":amd: (MI355) Qwen3.8-Flash-Next-FP8 Accuracy Eval" + dind: false + device: mi355_4 + num_devices: 4 + timeout_in_minutes: 105 + working_dir: /vllm-workspace/tests + depends_on: + - image-build-amd + source_file_dependencies: + - csrc/libtorch_stable/gdn/ + - tests/evals/qwen4_exp/ + - vllm/model_executor/layers/mamba/ + - vllm/models/qwen4_exp/ + - vllm/transformers_utils/configs/qwen4_exp.py + - vllm/v1/attention/backends/short_conv_attn.py + - vllm/platforms/rocm.py + env: + VLLM_WORKER_MULTIPROC_METHOD: "spawn" + commands: + - uv pip install --system 'evalscope==1.10.0' + - pytest -s -v evals/qwen4_exp/test_accuracy.py --config-list-file=configs/models-rocm.txt - label: ":nvidia: (DGX) Spark GPQA Eval (GPT-OSS)" key: gpqa-eval-gpt-oss-spark diff --git a/.buildkite/test_areas/misc.yaml b/.buildkite/test_areas/misc.yaml index 1bff8e1850..6709dbb565 100644 --- a/.buildkite/test_areas/misc.yaml +++ b/.buildkite/test_areas/misc.yaml @@ -237,6 +237,27 @@ steps: commands: - export VLLM_WORKER_MULTIPROC_METHOD=spawn - pytest -v -s tests/v1/kv_connector/extract_hidden_states_integration + mirror: + amd: + label: ":amd: (MI355) Extract Hidden States Integration Single GPU" + dind: false + device: mi355_1 + num_devices: 1 + timeout_in_minutes: 45 + working_dir: /vllm-workspace + depends_on: + - image-build-amd + source_file_dependencies: + - vllm/v1/spec_decode/extract_hidden_states.py + - vllm/model_executor/models/extract_hidden_states.py + - vllm/transformers_utils/configs/extract_hidden_states.py + - vllm/distributed/kv_transfer/kv_connector/v1/example_hidden_states_connector.py + - tests/v1/kv_connector/extract_hidden_states_integration + - vllm/platforms/rocm.py + env: + VLLM_WORKER_MULTIPROC_METHOD: "spawn" + commands: + - pytest -v -s -m 'not distributed' tests/v1/kv_connector/extract_hidden_states_integration - label: ":nvidia: (L4) Extract Hidden States Integration" key: extract-hidden-states-integration-2-gpus @@ -416,7 +437,6 @@ steps: dind: false device: mi300_1 timeout_in_minutes: 55 - soft_fail: true depends_on: - image-build-amd source_file_dependencies: @@ -505,6 +525,9 @@ steps: - tests/renderers - tests/standalone_tests/lazy_imports.py - tests/reasoning + - tools/pre_commit/check_test_tethering.py + - tools/pre_commit/test_tethering_allowlist.txt + - tests/tools/test_check_test_tethering.py - tests/tool_parsers - tests/tokenizers_ - tests/parser @@ -532,6 +555,7 @@ steps: commands: - pytest -v -s -m 'cpu_test' multimodal - pytest -v -s config + - pytest -v -s tools/test_check_test_tethering.py - label: ":computer: (CPU) Reasoning + Renderers" key: cpu-reasoning-renderers diff --git a/.buildkite/test_areas/models_basic.yaml b/.buildkite/test_areas/models_basic.yaml index 6702eb6dd3..930c2dd590 100644 --- a/.buildkite/test_areas/models_basic.yaml +++ b/.buildkite/test_areas/models_basic.yaml @@ -119,6 +119,17 @@ steps: - tests/models/glm5next/ commands: - pytest -v -s models/glm5next + mirror: + amd: + label: ":amd: (MI355) GLM5Next" + dind: false + device: mi355_1 + timeout_in_minutes: 35 + working_dir: "/vllm-workspace/tests" + depends_on: + - image-build-amd + source_file_dependencies: + - vllm/platforms/rocm.py - label: ":nvidia: (H200) Qwen4 Exp" key: qwen4-exp-unit-tests diff --git a/.buildkite/test_areas/models_language.yaml b/.buildkite/test_areas/models_language.yaml index 3f80dfdfa3..c8473e7f0b 100644 --- a/.buildkite/test_areas/models_language.yaml +++ b/.buildkite/test_areas/models_language.yaml @@ -91,7 +91,7 @@ steps: commands: - MAMBA_FORCE_BUILD=TRUE uv pip install --system --no-build-isolation 'git+https://github.com/AndreasKaratzas/mamba@fix-rocm-7.0-warp-size-constexpr' - CAUSAL_CONV1D_FORCE_BUILD=TRUE 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 + - 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 # Granite 4 hybrid generation is sensitive to hardware-specific Triton SSD # autotuning (https://github.com/vllm-project/vllm/issues/25194). Keep this one @@ -111,6 +111,27 @@ steps: - rm -rf /tmp/mamba-src && git clone --depth 1 --branch v2.3.0 https://github.com/state-spaces/mamba /tmp/mamba-src && sed -i 's/-std=c++17/-std=c++20/g' /tmp/mamba-src/setup.py && MAMBA_FORCE_BUILD=TRUE uv pip install --system --no-build-isolation /tmp/mamba-src - CAUSAL_CONV1D_FORCE_BUILD=TRUE 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' + mirror: + amd: + label: ":amd: (MI355) Granite Language Model Compatibility" + dind: false + device: mi355_1 + num_devices: 1 + timeout_in_minutes: 80 + working_dir: /vllm-workspace/tests + depends_on: + - image-build-amd + source_file_dependencies: + - vllm/ + - "!vllm/distributed/kv_transfer/" + - tests/models/language/generation + commands: + - export GIT_CONFIG_COUNT=1 + - export GIT_CONFIG_KEY_0=http.version + - export GIT_CONFIG_VALUE_0=HTTP/1.1 + - MAMBA_FORCE_BUILD=TRUE uv pip install --system --no-build-isolation 'git+https://github.com/AndreasKaratzas/mamba@fix-rocm-7.0-warp-size-constexpr' + - CAUSAL_CONV1D_FORCE_BUILD=TRUE 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: ":nvidia: (H200 MIG 35GB) Language Models (Extended Generation)" device: h200_35gb diff --git a/.buildkite/test_areas/pytorch.yaml b/.buildkite/test_areas/pytorch.yaml index 459f7faeb4..8204fb9670 100644 --- a/.buildkite/test_areas/pytorch.yaml +++ b/.buildkite/test_areas/pytorch.yaml @@ -87,6 +87,20 @@ steps: - tests/compile/h100/ commands: - "find compile/h100/ -name 'test_*.py' -print0 | xargs -0 -n1 -I{} pytest -s -v '{}'" + mirror: + amd: + label: ":amd: (MI355) PyTorch Compilation H100 Cases" + dind: false + device: mi355_1 + timeout_in_minutes: 45 + working_dir: "/vllm-workspace/tests" + depends_on: + - image-build-amd + source_file_dependencies: + - tests/utils.py + - csrc/ + commands: + - pytest -v -s compile/h100/test_startup.py - label: ":nvidia: (L4) PyTorch Compilation Passes" key: pytorch-compilation-passes-unit-tests diff --git a/.buildkite/test_areas/samplers.yaml b/.buildkite/test_areas/samplers.yaml index e1bc4c08d8..084583c8a7 100644 --- a/.buildkite/test_areas/samplers.yaml +++ b/.buildkite/test_areas/samplers.yaml @@ -35,7 +35,7 @@ steps: - vllm/_aiter_ops.py - vllm/platforms/rocm.py commands: - - pytest -v -s samplers + - pytest -v -s samplers -k "not test_beam_search_passes_multimodal_data" # Qwen2-Audio-7B leaves no room for KV cache on a 16 GiB slice. - label: ":nvidia: (H200 MIG 35GB) Samplers Multimodal Beam Search" @@ -51,3 +51,24 @@ steps: commands: - VLLM_USE_FLASHINFER_SAMPLER=0 pytest -v -s samplers/test_beam_search.py::test_beam_search_passes_multimodal_data - VLLM_USE_FLASHINFER_SAMPLER=1 pytest -v -s samplers/test_beam_search.py::test_beam_search_passes_multimodal_data + mirror: + amd: + label: ":amd: (MI355) Samplers Multimodal Beam Search" + dind: false + device: mi355_1 + num_devices: 1 + timeout_in_minutes: 45 + working_dir: /vllm-workspace/tests + depends_on: + - image-build-amd + source_file_dependencies: + - vllm/model_executor/layers + - vllm/sampling_metadata.py + - vllm/v1/sample/ + - vllm/entrypoints/generate/beam_search/ + - tests/samplers + - tests/conftest.py + - vllm/_aiter_ops.py + - vllm/platforms/rocm.py + commands: + - pytest -v -s samplers/test_beam_search.py::test_beam_search_passes_multimodal_data diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml index b12c8ce508..4823f788dc 100644 --- a/.pre-commit-config.yaml +++ b/.pre-commit-config.yaml @@ -184,6 +184,14 @@ repos: files: ^\.buildkite/.*\.ya?ml$ pass_filenames: false additional_dependencies: [PyYAML] + - id: check-test-tethering + name: Check new tests are wired into Buildkite CI + language: python + entry: python tools/pre_commit/check_test_tethering.py + files: ^(tests/.*\.py|\.buildkite/(test_areas|hardware_tests|intel_jobs)/.*\.yaml|\.buildkite/(test-amd|ci_config[^/]*)\.yaml|tools/pre_commit/test_tethering_allowlist\.txt)$ + additional_dependencies: [PyYAML, regex] + verbose: true # full-scan `note:` lines are advisory, so the hook passes + require_serial: true # one full scan, not one per pre-commit file batch - id: mypy-3.10 # TODO: Use https://github.com/pre-commit/mirrors-mypy when mypy setup is less awkward name: Run mypy for Python 3.10 entry: python tools/pre_commit/mypy.py "3.10" diff --git a/CMakeLists.txt b/CMakeLists.txt index fbacfd9091..dbf4ca90e2 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -460,6 +460,8 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP") DEFINED CMAKE_CUDA_COMPILER_VERSION AND CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL 12.0) + list(APPEND VLLM_STABLE_EXT_SRC "csrc/libtorch_stable/all_reduce_mhc.cu") + if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0) cuda_archs_loose_intersection(COOPERATIVE_TOPK_ARCHS "9.0a;10.0f;10.1f;10.3f;10.7f;11.0f;12.0f;12.1f" "${CUDA_ARCHS}") diff --git a/benchmarks/kernels/benchmark_gemm_rs_ar.py b/benchmarks/kernels/benchmark_gemm_rs_ar.py new file mode 100644 index 0000000000..3e47f61c74 --- /dev/null +++ b/benchmarks/kernels/benchmark_gemm_rs_ar.py @@ -0,0 +1,407 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""Benchmark the SM100 GEMM-RS/AR kernel against the model's unfused path. + +The baseline is what a model runs without the fusion: the projection's own +linear kernel followed by vLLM's collective (``sp_reduce_scatter`` for RS, +``tensor_model_parallel_all_reduce`` for AR), with the collective backend +chosen by vLLM's dispatch exactly as in serving. The same two kernels are also +timed bare (the GEMM alone, the collective alone on a precomputed partial); +their max is the floor a perfectly overlapped fused kernel could reach. + +All ranks must belong to one NVLink domain. Run a TP8 sweep over the Kimi-K3 +projections with: + + torchrun --nproc-per-node=8 -- benchmarks/kernels/benchmark_gemm_rs_ar.py + +or the DeepSeek-V4.1 MXFP8 ``wo_b`` projection at TP4 with: + + torchrun --nproc-per-node=4 -- benchmarks/kernels/benchmark_gemm_rs_ar.py \ + --model deepseek_v41 --backend flashinfer_cutedsl +""" + +import argparse +import json +import os +import statistics +from collections.abc import Callable + +import torch +import torch.distributed as dist + +from vllm.config import ModelConfig, VllmConfig, set_current_vllm_config +from vllm.distributed import ( + cleanup_dist_env_and_memory, + tensor_model_parallel_all_reduce, +) +from vllm.distributed.parallel_state import ( + get_tp_group, + init_distributed_environment, + initialize_model_parallel, +) +from vllm.model_executor.kernels.linear.cute_dsl.gemm_rs_ar import GemmRsAr +from vllm.model_executor.layers.linear import RowParallelLinear +from vllm.models.common.ops.sequence_parallel import sp_reduce_scatter + +# Per model: (hidden size N, global input widths K). Kimi-K3 lists the +# shared-expert down-proj and attention O-proj; DeepSeek-V4.1 lists the +# ``wo_b`` output projection (o_groups * o_lora_rank). +_MODEL_PROJECTIONS = { + "kimi_k3": (7168, (6144, 12288)), + "deepseek_v41": (5120, (8192,)), +} + + +def parse_args() -> argparse.Namespace: + parser = argparse.ArgumentParser() + parser.add_argument( + "--mode", + choices=("rs", "ar"), + default="rs", + help="Collective mode to benchmark.", + ) + parser.add_argument( + "--m", + type=int, + nargs="+", + default=[128, 512, 2048, 8192, 32768], + help="Global token counts to benchmark.", + ) + parser.add_argument( + "--model", + choices=tuple(_MODEL_PROJECTIONS), + default="kimi_k3", + help="Model whose projection shapes set the default --n and --k.", + ) + parser.add_argument( + "--backend", + choices=("bf16", "flashinfer_cutlass", "flashinfer_cutedsl"), + default="bf16", + help="Weight dtype/layout: BF16 or an online-MXFP8 FlashInfer kernel.", + ) + parser.add_argument( + "--k", + type=int, + nargs="+", + help=( + "Per-rank input dimensions. By default, derive the selected " + "model's projection dimensions from the TP world size." + ), + ) + parser.add_argument( + "--n", + type=int, + help="Output width. Defaults to the selected model's hidden size.", + ) + parser.add_argument( + "--num-workspaces", + type=int, + default=10, + help="Pointer-distinct inputs and CUDA graphs to rotate.", + ) + parser.add_argument("--warmup-replays", type=int, default=5) + parser.add_argument("--samples", type=int, default=20) + parser.add_argument("--json", type=str, help="Also write the results here.") + return parser.parse_args() + + +def make_projection( + gemm_rs_ar: GemmRsAr, + weight: torch.Tensor, + world_size: int, + backend: str, +) -> RowParallelLinear: + """Wrap a BF16 [N, K] weight the way the model would hand it to GEMM-RS. + + For the MXFP8 backends this quantizes online through the FlashInfer + kernel's post-load step, so the fused kernel and the unfused baseline see + the production weight and scale layouts. + """ + N, K = weight.shape + quant_config = None + if backend != "bf16": + from vllm.config.quantization import QuantizationConfigArgs + from vllm.model_executor.layers.quantization.online.base import ( + OnlineQuantizationConfig, + ) + + quant_config = OnlineQuantizationConfig(QuantizationConfigArgs(linear="mxfp8")) + with torch.device(weight.device): + linear = RowParallelLinear( + K * world_size, + N, + bias=False, + params_dtype=weight.dtype, + quant_config=quant_config, + # The baseline reduces the local partial itself, like the models. + reduce_results=False, + return_bias=False, + ) + # Select the fused path before online quantization replaces the weights. + assert gemm_rs_ar.can_run(linear) + linear.weight = torch.nn.Parameter(weight, requires_grad=False) + linear.quant_method.process_weights_after_loading(linear) + return linear + + +def capture_graph( + op: Callable[[], torch.Tensor], + stream: torch.cuda.Stream, + cpu_group: dist.ProcessGroup, +) -> torch.cuda.CUDAGraph: + stream.wait_stream(torch.cuda.current_stream()) + dist.barrier(group=cpu_group) + with torch.cuda.stream(stream): + for _ in range(3): + op() + stream.synchronize() + dist.barrier(group=cpu_group) + + graph = torch.cuda.CUDAGraph() + with torch.cuda.graph(graph, stream=stream): + op() + torch.cuda.current_stream().wait_stream(stream) + dist.barrier(group=cpu_group) + return graph + + +def benchmark_graphs( + candidate_graphs: dict[str, list[torch.cuda.CUDAGraph]], + warmup_replays: int, + samples: int, + device_group: dist.ProcessGroup, + device_barrier: Callable[[], None], +) -> dict[str, float]: + """Median over samples of the slowest rank's replay time, in us.""" + names = list(candidate_graphs) + for round_index in range(warmup_replays): + for offset in range(len(names)): + name = names[(round_index + offset) % len(names)] + graphs = candidate_graphs[name] + device_barrier() + graphs[round_index % len(graphs)].replay() + torch.accelerator.synchronize() + + timings: dict[str, list[float]] = {name: [] for name in names} + start = torch.cuda.Event(enable_timing=True) + end = torch.cuda.Event(enable_timing=True) + for sample_index in range(samples): + for offset in range(len(names)): + name = names[(sample_index + offset) % len(names)] + graphs = candidate_graphs[name] + device_barrier() + start.record() + graphs[sample_index % len(graphs)].replay() + end.record() + end.synchronize() + + elapsed = torch.tensor( + start.elapsed_time(end) * 1000, + dtype=torch.float64, + device=torch.accelerator.current_device_index(), + ) + dist.all_reduce(elapsed, op=dist.ReduceOp.MAX, group=device_group) + timings[name].append(elapsed.item()) + return {name: statistics.median(values) for name, values in timings.items()} + + +def benchmark_shape( + gemm_rs_ar: GemmRsAr, + mode: str, + backend: str, + M: int, + N: int, + K: int, + num_workspaces: int, + warmup_replays: int, + samples: int, + device_group: dist.ProcessGroup, + cpu_group: dist.ProcessGroup, + device_barrier: Callable[[], None], +) -> dict[str, float | int | str]: + all_reduce = mode == "ar" + world_size = dist.get_world_size(device_group) + rank = dist.get_rank(device_group) + device = torch.device("cuda", torch.accelerator.current_device_index()) + local_M = (M + world_size - 1) // world_size + + rng = torch.Generator(device=device) + rng.manual_seed(1000 + rank * 10 + M + K) + inputs = [ + torch.randn(M, K, dtype=torch.bfloat16, device=device, generator=rng) + for _ in range(num_workspaces) + ] + projections = [ + make_projection( + gemm_rs_ar, + torch.randn(N, K, dtype=torch.bfloat16, device=device, generator=rng), + world_size, + backend, + ) + for _ in range(num_workspaces) + ] + + def collective(partial: torch.Tensor) -> torch.Tensor: + if all_reduce: + return tensor_model_parallel_all_reduce(partial) + return sp_reduce_scatter(partial) + + def gemm(x: torch.Tensor, linear: RowParallelLinear) -> torch.Tensor: + return linear(x) + + def unfused(x: torch.Tensor, linear: RowParallelLinear) -> torch.Tensor: + return collective(linear(x)) + + def fused(x: torch.Tensor, linear: RowParallelLinear) -> torch.Tensor: + return gemm_rs_ar.apply(x, linear) + + # The bare collective runs on the GEMM's own output so it sees the same + # message size and layout the unfused path hands to the collective. + partials = [linear(x) for x, linear in zip(inputs, projections)] + + expected = unfused(inputs[0], projections[0]) + actual = fused(inputs[0], projections[0]) + torch.accelerator.synchronize(device) + rows = M if all_reduce else min(max(M - rank * local_M, 0), local_M) + torch.testing.assert_close( + actual[:rows], expected[:rows], rtol=5e-2, atol=4.0, msg=f"M={M}, K={K}" + ) + + candidates: dict[str, list[Callable[[], torch.Tensor]]] = { + "gemm": [lambda x=x, w=w: gemm(x, w) for x, w in zip(inputs, projections)], + "collective": [lambda p=p: collective(p) for p in partials], + "vllm": [lambda x=x, w=w: unfused(x, w) for x, w in zip(inputs, projections)], + "gemm_rs_ar": [ + lambda x=x, w=w: fused(x, w) for x, w in zip(inputs, projections) + ], + } + + candidate_graphs = {} + keepalive: list[object] = [] + for name, ops in candidates.items(): + stream = torch.cuda.Stream() + candidate_graphs[name] = [capture_graph(op, stream, cpu_group) for op in ops] + keepalive.append(stream) + + times = benchmark_graphs( + candidate_graphs, warmup_replays, samples, device_group, device_barrier + ) + floor = max(times["gemm"], times["collective"]) + return { + "mode": mode.upper(), + "backend": backend, + "M": M, + "N": N, + "K": K, + "gemm_us": times["gemm"], + "collective_us": times["collective"], + "floor_us": floor, + "vllm_us": times["vllm"], + "gemm_rs_ar_us": times["gemm_rs_ar"], + "latency_change": (times["gemm_rs_ar"] - times["vllm"]) / times["vllm"], + "floor_gap": (times["gemm_rs_ar"] - floor) / floor, + } + + +def print_results(results: list[dict[str, float | int | str]]) -> None: + collective = results[0]["mode"] + backend = results[0]["backend"] + print(f"### GEMM-{collective} vs vLLM unfused path ({backend} weights)") + print( + f"| M | N | K | GEMM (us) | {collective} (us) | floor (us) | vLLM (us) | " + f"GEMM-{collective} (us) | vs vLLM | vs floor |" + ) + print("|---:|---:|---:|---:|---:|---:|---:|---:|---:|---:|") + for r in results: + print( + f"| {r['M']} | {r['N']} | {r['K']} | {r['gemm_us']:.1f} | " + f"{r['collective_us']:.1f} | {r['floor_us']:.1f} | {r['vllm_us']:.1f} | " + f"{r['gemm_rs_ar_us']:.1f} | {r['latency_change']:+.1%} | " + f"{r['floor_gap']:+.1%} |" + ) + + +def main() -> None: + args = parse_args() + default_n, default_k = _MODEL_PROJECTIONS[args.model] + if args.n is None: + args.n = default_n + assert args.m and min(args.m) > 0 + assert args.n % 128 == 0 + assert args.num_workspaces > 0 + assert args.warmup_replays >= 0 + assert args.samples > 0 + + local_rank = int(os.environ["LOCAL_RANK"]) + torch.accelerator.set_device_index(local_rank) + init_distributed_environment() + world_size = dist.get_world_size() + if args.k is None: + assert all(K % world_size == 0 for K in default_k) + K_values = [K // world_size for K in default_k] + else: + K_values = args.k + k_alignment = 64 if args.backend == "bf16" else 128 + assert all(K % k_alignment == 0 for K in K_values) + + config = VllmConfig() + config.model_config = ModelConfig(dtype="bfloat16") + # Pin the online-MXFP8 kernel the projections quantize through. + config.kernel_config.linear_backend = ( + "auto" if args.backend == "bf16" else args.backend + ) + with set_current_vllm_config(config): + # Builds the TP communicator as in serving, so the baseline's + # collective is whatever vLLM's dispatch selects (FlashInfer all-reduce + # first, then NCCL symmetric memory, custom all-reduce, pynccl). + initialize_model_parallel(tensor_model_parallel_size=world_size) + + tp_group = get_tp_group() + pynccl_comm = tp_group.device_communicator.pynccl_comm + assert pynccl_comm is not None + sync_input = torch.zeros(1, device=torch.accelerator.current_device_index()) + sync_output = torch.empty_like(sync_input) + + def device_barrier() -> None: + # Order the timed launch after a device-side rank rendezvous without + # including the rendezvous itself in the measured event interval. + pynccl_comm.all_reduce(sync_input, sync_output) + + gemm_rs_ar = GemmRsAr( + max_M=max(args.m), + N=args.n, + all_reduce=args.mode == "ar", + ) + with set_current_vllm_config(config): + results = [ + benchmark_shape( + gemm_rs_ar, + args.mode, + args.backend, + M, + args.n, + K, + args.num_workspaces, + args.warmup_replays, + args.samples, + tp_group.device_group, + tp_group.cpu_group, + device_barrier, + ) + for K in K_values + for M in args.m + ] + del gemm_rs_ar + + if tp_group.rank_in_group == 0: + print_results(results) + if args.json: + with open(args.json, "w") as f: + json.dump(results, f, indent=2) + + dist.barrier(group=tp_group.cpu_group) + cleanup_dist_env_and_memory() + + +if __name__ == "__main__": + main() diff --git a/benchmarks/kernels/benchmark_kimi_k3_gemm_rs_ar.py b/benchmarks/kernels/benchmark_kimi_k3_gemm_rs_ar.py deleted file mode 100644 index e4daea26bb..0000000000 --- a/benchmarks/kernels/benchmark_kimi_k3_gemm_rs_ar.py +++ /dev/null @@ -1,469 +0,0 @@ -# SPDX-License-Identifier: Apache-2.0 -# SPDX-FileCopyrightText: Copyright contributors to the vLLM project -"""Benchmark the SM100 Kimi-K3 GEMM-RS/AR kernel. - -All ranks must belong to one NVLink domain. For example, run a TP8 sweep with: - - torchrun --nproc-per-node=8 \ - benchmarks/kernels/benchmark_kimi_k3_gemm_rs_ar.py -""" - -import argparse -import os -import statistics -from collections.abc import Callable -from dataclasses import dataclass - -import pandas as pd -import torch -import torch.distributed as dist -import torch.distributed._symmetric_memory as symm_mem - -from vllm.config import VllmConfig, set_current_vllm_config -from vllm.distributed import cleanup_dist_env_and_memory -from vllm.distributed.parallel_state import ( - get_tp_group, - init_distributed_environment, - initialize_model_parallel, -) -from vllm.models.kimi_k3.nvidia.ops.cute_dsl.gemm_rs_ar import GemmRsAr - -# Shared-expert down-proj and attention O-proj. -_KIMI_K3_PROJECTION_K = (6144, 12288) - - -@dataclass -class Candidate: - name: str - runs: list[Callable[[], torch.Tensor]] - check_correctness: bool = True - - -def parse_args() -> argparse.Namespace: - parser = argparse.ArgumentParser() - parser.add_argument( - "--mode", - choices=("rs", "ar"), - default="rs", - help="Collective mode to benchmark.", - ) - parser.add_argument( - "--m", - type=int, - nargs="+", - default=[128, 512, 2048, 8192, 32768], - help="Global token counts to benchmark.", - ) - parser.add_argument( - "--k", - type=int, - nargs="+", - help=( - "Per-rank input dimensions. By default, derive the Kimi-K3 " - "shared-expert down-proj and O-proj dimensions from the TP " - "world size." - ), - ) - parser.add_argument("--n", type=int, default=7168) - parser.add_argument( - "--num-workspaces", - type=int, - default=10, - help="Pointer-distinct inputs and CUDA graphs to rotate.", - ) - parser.add_argument("--warmup-replays", type=int, default=5) - parser.add_argument("--samples", type=int, default=20) - return parser.parse_args() - - -def capture_graph( - op: Callable[[], torch.Tensor], - stream: torch.cuda.Stream, - cpu_group: dist.ProcessGroup, -) -> tuple[torch.cuda.CUDAGraph, list[torch.Tensor | None]]: - result: list[torch.Tensor | None] = [None] - stream.wait_stream(torch.cuda.current_stream()) - dist.barrier(group=cpu_group) - with torch.cuda.stream(stream): - for _ in range(3): - result[0] = op() - stream.synchronize() - dist.barrier(group=cpu_group) - - graph = torch.cuda.CUDAGraph() - with torch.cuda.graph(graph, stream=stream): - result[0] = op() - torch.cuda.current_stream().wait_stream(stream) - dist.barrier(group=cpu_group) - return graph, result - - -def benchmark_graphs( - candidate_graphs: dict[str, list[torch.cuda.CUDAGraph]], - warmup_replays: int, - samples: int, - device_group: dist.ProcessGroup, - device_barrier: Callable[[], None], -) -> dict[str, float]: - candidate_names = list(candidate_graphs) - for round_index in range(warmup_replays): - for candidate_index in range(len(candidate_names)): - candidate_id = (round_index + candidate_index) % len(candidate_names) - name = candidate_names[candidate_id] - graphs = candidate_graphs[name] - device_barrier() - graphs[round_index % len(graphs)].replay() - torch.accelerator.synchronize() - - timings: dict[str, list[float]] = {name: [] for name in candidate_names} - start = torch.cuda.Event(enable_timing=True) - end = torch.cuda.Event(enable_timing=True) - for sample_index in range(samples): - for candidate_index in range(len(candidate_names)): - candidate_id = (sample_index + candidate_index) % len(candidate_names) - name = candidate_names[candidate_id] - graphs = candidate_graphs[name] - device_barrier() - start.record() - graphs[sample_index % len(graphs)].replay() - end.record() - end.synchronize() - - elapsed = torch.tensor( - start.elapsed_time(end) * 1000, - dtype=torch.float64, - device=torch.accelerator.current_device_index(), - ) - dist.all_reduce(elapsed, op=dist.ReduceOp.MAX, group=device_group) - timings[name].append(elapsed.item()) - return {name: statistics.median(values) for name, values in timings.items()} - - -def valid_rows(M: int, local_M: int, rank: int) -> int: - return min(max(M - rank * local_M, 0), local_M) - - -def benchmark_shape( - gemm_rs_ar: GemmRsAr, - mode: str, - M: int, - N: int, - K: int, - num_workspaces: int, - warmup_replays: int, - samples: int, - device_group: dist.ProcessGroup, - cpu_group: dist.ProcessGroup, - device_barrier: Callable[[], None], -) -> dict[str, float | int | str]: - all_reduce = mode == "ar" - world_size = dist.get_world_size(device_group) - rank = dist.get_rank(device_group) - device = torch.device("cuda", torch.accelerator.current_device_index()) - padded_M = (M + world_size - 1) // world_size * world_size - local_M = padded_M // world_size - - rng = torch.Generator(device=device) - rng.manual_seed(1000 + rank * 10 + M + K) - inputs = [ - torch.randn(M, K, dtype=torch.bfloat16, device=device, generator=rng) - for _ in range(num_workspaces) - ] - weights = [ - torch.randn(N, K, dtype=torch.bfloat16, device=device, generator=rng) - for _ in range(num_workspaces) - ] - - partial = torch.empty((padded_M, N), dtype=torch.bfloat16, device=device) - symm_partial = symm_mem.empty((padded_M, N), dtype=torch.bfloat16, device=device) - symm_partial_handle = symm_mem.rendezvous(symm_partial, device_group) - collective_inputs = [torch.empty_like(partial) for _ in range(num_workspaces)] - symm_collective_inputs = [] - symm_collective_handles = [] - for _ in range(num_workspaces): - collective_input = symm_mem.empty( - (padded_M, N), - dtype=torch.bfloat16, - device=device, - ) - symm_collective_inputs.append(collective_input) - symm_collective_handles.append( - symm_mem.rendezvous(collective_input, device_group) - ) - torch_output = torch.empty((local_M, N), dtype=torch.bfloat16, device=device) - symm_output = torch.empty_like(torch_output) - gemm_output = torch.empty((M, N), dtype=torch.bfloat16, device=device) - - if padded_M > M: - partial[M:].zero_() - symm_partial[M:].zero_() - - def make_torch_ring_ll_gemm_collective( - x: torch.Tensor, weight: torch.Tensor - ) -> Callable[[], torch.Tensor]: - def run() -> torch.Tensor: - torch.mm(x, weight.T, out=partial[:M]) - if all_reduce: - dist.all_reduce(partial, group=device_group) - return partial[:M] - dist.reduce_scatter_single(torch_output, partial, group=device_group) - return torch_output - - return run - - def make_torch_ldmc_gemm_collective( - x: torch.Tensor, weight: torch.Tensor - ) -> Callable[[], torch.Tensor]: - def run() -> torch.Tensor: - torch.mm(x, weight.T, out=symm_partial[:M]) - if all_reduce: - dist.all_reduce(symm_partial, group=device_group) - return symm_partial[:M] - dist.reduce_scatter_single( - symm_output, - symm_partial, - group=device_group, - ) - return symm_output - - return run - - def make_fused_gemm_collective( - x: torch.Tensor, weight: torch.Tensor - ) -> Callable[[], torch.Tensor]: - def run() -> torch.Tensor: - return gemm_rs_ar(x, weight) - - return run - - def make_torch_gemm( - x: torch.Tensor, - weight: torch.Tensor, - ) -> Callable[[], torch.Tensor]: - def run() -> torch.Tensor: - return torch.mm(x, weight.T, out=gemm_output) - - return run - - def make_ring_ll_collective( - collective_input: torch.Tensor, - ) -> Callable[[], torch.Tensor]: - def run() -> torch.Tensor: - if all_reduce: - dist.all_reduce(collective_input, group=device_group) - return collective_input - dist.reduce_scatter_single( - torch_output, collective_input, group=device_group - ) - return torch_output - - return run - - def make_ldmc_collective( - collective_input: torch.Tensor, - ) -> Callable[[], torch.Tensor]: - def run() -> torch.Tensor: - if all_reduce: - dist.all_reduce(collective_input, group=device_group) - return collective_input - dist.reduce_scatter_single( - symm_output, collective_input, group=device_group - ) - return symm_output - - return run - - candidates = ( - Candidate( - "ring_ll_us", - [make_torch_ring_ll_gemm_collective(x, w) for x, w in zip(inputs, weights)], - ), - Candidate( - "ldmc_us", - [make_torch_ldmc_gemm_collective(x, w) for x, w in zip(inputs, weights)], - ), - Candidate( - "gemm_rs_ar_us", - [make_fused_gemm_collective(x, w) for x, w in zip(inputs, weights)], - ), - Candidate( - "torch_gemm_us", - [make_torch_gemm(x, w) for x, w in zip(inputs, weights)], - check_correctness=False, - ), - Candidate( - "ring_ll_collective_us", - [make_ring_ll_collective(x) for x in collective_inputs], - check_correctness=False, - ), - Candidate( - "ldmc_collective_us", - [make_ldmc_collective(x) for x in symm_collective_inputs], - check_correctness=False, - ), - ) - - expected = candidates[0].runs[0]() - rows = M if all_reduce else valid_rows(M, local_M, rank) - for candidate in candidates[1:]: - if not candidate.check_correctness: - continue - actual = candidate.runs[0]() - torch.accelerator.synchronize(device) - torch.testing.assert_close( - actual[:rows], - expected[:rows], - rtol=5e-2, - atol=4.0, - ) - - candidate_graphs = {} - graph_keepalive: list[object] = [symm_partial_handle, *symm_collective_handles] - for candidate in candidates: - stream = torch.cuda.Stream() - bundles = [capture_graph(run, stream, cpu_group) for run in candidate.runs] - candidate_graphs[candidate.name] = [graph for graph, _ in bundles] - graph_keepalive.extend(bundles) - graph_keepalive.append(stream) - - times = benchmark_graphs( - candidate_graphs, - warmup_replays, - samples, - device_group, - device_barrier, - ) - - best_nccl_us = min(times["ring_ll_collective_us"], times["ldmc_collective_us"]) - return { - "mode": mode.upper(), - "M": M, - "N": N, - "K": K, - **times, - "best_nccl_us": best_nccl_us, - "speedup_vs_ring_ll": times["ring_ll_us"] / times["gemm_rs_ar_us"], - "speedup_vs_ldmc": times["ldmc_us"] / times["gemm_rs_ar_us"], - } - - -def print_results(results: list[dict[str, float | int | str]]) -> None: - results_df = pd.DataFrame(results) - collective = str(results_df["mode"].iloc[0]) - end_to_end = ( - results_df[ - [ - "M", - "N", - "K", - "ring_ll_us", - "ldmc_us", - "gemm_rs_ar_us", - "speedup_vs_ring_ll", - "speedup_vs_ldmc", - ] - ] - .rename( - columns={ - "ring_ll_us": f"Torch GEMM + NCCL {collective} (RING_LL) (us)", - "ldmc_us": f"Torch GEMM + NCCL {collective} (LDMC) (us)", - "gemm_rs_ar_us": f"GEMM-{collective} (us)", - "speedup_vs_ring_ll": "Speedup vs RING_LL", - "speedup_vs_ldmc": "Speedup vs LDMC", - } - ) - .round(3) - ) - components = ( - results_df[["M", "N", "K", "torch_gemm_us", "best_nccl_us", "gemm_rs_ar_us"]] - .rename( - columns={ - "torch_gemm_us": "Torch GEMM (us)", - "best_nccl_us": f"NCCL {collective} (best) (us)", - "gemm_rs_ar_us": f"GEMM-{collective} (us)", - } - ) - .round(2) - ) - - print(f"### {collective} end-to-end latency") - print(end_to_end.to_markdown(index=False)) - print(f"\n### {collective} component latency") - print(components.to_markdown(index=False)) - print( - f"\nNCCL {collective} (best) is the faster of RING_LL and LDMC " - "for each shape.\n" - ) - - -def main() -> None: - args = parse_args() - assert args.m and min(args.m) > 0 - assert args.n % 256 == 0 - assert args.num_workspaces > 0 - assert args.warmup_replays >= 0 - assert args.samples > 0 - - local_rank = int(os.environ["LOCAL_RANK"]) - torch.accelerator.set_device_index(local_rank) - init_distributed_environment() - world_size = dist.get_world_size() - if args.k is None: - assert all(K % world_size == 0 for K in _KIMI_K3_PROJECTION_K) - K_values = [K // world_size for K in _KIMI_K3_PROJECTION_K] - else: - K_values = args.k - assert all(K % 64 == 0 for K in K_values) - # Reserve symmetric memory for the NCCL-managed benchmark allocations. - os.environ["VLLM_ALLREDUCE_USE_SYMM_MEM"] = "0" - # NCCL-managed symmetric allocations select the NVLS/LDMC collective path. - symm_mem.set_backend("NCCL") - with set_current_vllm_config(VllmConfig()): - initialize_model_parallel(tensor_model_parallel_size=world_size) - - tp_group = get_tp_group() - group_warmup = torch.zeros(1, device=torch.accelerator.current_device_index()) - dist.all_reduce(group_warmup, group=tp_group.device_group) - pynccl_comm = tp_group.device_communicator.pynccl_comm - assert pynccl_comm is not None - sync_input = torch.zeros(1, device=torch.accelerator.current_device_index()) - sync_output = torch.empty_like(sync_input) - - def device_barrier() -> None: - # Order the timed launch after a device-side rank rendezvous without - # including the rendezvous itself in the measured event interval. - pynccl_comm.all_reduce(sync_input, sync_output) - - gemm_rs_ar = GemmRsAr( - max_M=max(args.m), - N=args.n, - all_reduce=args.mode == "ar", - ) - results = [ - benchmark_shape( - gemm_rs_ar, - args.mode, - M, - args.n, - K, - args.num_workspaces, - args.warmup_replays, - args.samples, - tp_group.device_group, - tp_group.cpu_group, - device_barrier, - ) - for K in K_values - for M in args.m - ] - del gemm_rs_ar - - if tp_group.rank_in_group == 0: - print_results(results) - - dist.barrier(group=tp_group.cpu_group) - cleanup_dist_env_and_memory() - - -if __name__ == "__main__": - main() diff --git a/cmake/external_projects/deepgemm.cmake b/cmake/external_projects/deepgemm.cmake index 9ac0d05414..a17031cc5e 100644 --- a/cmake/external_projects/deepgemm.cmake +++ b/cmake/external_projects/deepgemm.cmake @@ -29,10 +29,11 @@ if(DEEPGEMM_SRC_DIR) else() # Keep in sync with tools/install_deepgemm.sh set(_DEEPGEMM_UPSTREAM_REPO "https://github.com/vllm-project/DeepGEMM.git") - # Pinned to the tip of the fork's dev branch: upstream 2.8.0 plus the SM120 + # Pinned to the fork's dev branch: upstream 2.8.0 plus the SM120 # port, the SM90 paged-MQA kv_block=32/next_n=4 port, configurable SwiGLU - # alpha/beta, and SiTU for FP8/FP4 Mega MoE. - set(_DEEPGEMM_UPSTREAM_TAG "a6bbb8000161c0dc3a85a0300a905f76898a7913") + # alpha/beta, and SiTU for FP8/FP4 Mega MoE. Also includes the CUDA 12.x + # layout header fix from vllm-project/DeepGEMM#12. + set(_DEEPGEMM_UPSTREAM_TAG "e1f418c2a4f20818221f6b0e578b4c2f634d4c3f") set(_deepgemm_fc_root "${FETCHCONTENT_BASE_DIR}") if(NOT _deepgemm_fc_root) diff --git a/cmake/external_projects/vllm_flash_attn.cmake b/cmake/external_projects/vllm_flash_attn.cmake index 1b3be244de..2a26023f74 100644 --- a/cmake/external_projects/vllm_flash_attn.cmake +++ b/cmake/external_projects/vllm_flash_attn.cmake @@ -39,7 +39,7 @@ else() FetchContent_Declare( vllm-flash-attn GIT_REPOSITORY https://github.com/vllm-project/flash-attention.git - GIT_TAG 506341a143fcabd4bb79052a7605ada727d6b3f5 + GIT_TAG 9cd61de38763d712bb6ce56e2a02cc2bf718c89f GIT_PROGRESS TRUE # Don't share the vllm-flash-attn build between build types BINARY_DIR ${CMAKE_BINARY_DIR}/vllm-flash-attn diff --git a/csrc/cpu/sgl-kernels/common.h b/csrc/cpu/sgl-kernels/common.h index de1260bb33..dc0769450e 100644 --- a/csrc/cpu/sgl-kernels/common.h +++ b/csrc/cpu/sgl-kernels/common.h @@ -280,26 +280,10 @@ inline int get_thread_num() { // balance payload across each thread template inline void balance211(T n, T nth, T ith, T& n_start, T& n_end) { -#if 0 - // onednn partition pattern - T& n_my = n_end; - if (nth <= 1 || n == 0) { - n_start = 0; - n_my = n; - } else { - T n1 = div_up(n, nth); - T n2 = n1 - 1; - T T1 = n - n2 * nth; - n_my = ith < T1 ? n1 : n2; - n_start = ith <= T1 ? ith*n1 : T1 * n1 + (ith - T1) * n2; - } - n_end += n_start; -#else // pytorch aten partition pattern T n_my = div_up(n, nth); n_start = ith * n_my; n_end = std::min(n_start + n_my, n); -#endif } template diff --git a/csrc/libtorch_stable/all_reduce_mhc.cu b/csrc/libtorch_stable/all_reduce_mhc.cu new file mode 100644 index 0000000000..dc58371156 --- /dev/null +++ b/csrc/libtorch_stable/all_reduce_mhc.cu @@ -0,0 +1,212 @@ +// SPDX-License-Identifier: Apache-2.0 +// SPDX-FileCopyrightText: Copyright contributors to the vLLM project + +#include "torch_utils.h" +#include "custom_all_gather_reduce_scatter.cuh" +#include +#include + +namespace { +using Tensor = torch::stable::Tensor; +using Pack = vllm::packed_t::P; +constexpr int kHidden = 5120; +constexpr int kRanks = 4; +constexpr int kThreads = 128; +constexpr int kCluster = 5; +constexpr int kPacks = kHidden / Pack::size; + +__global__ __launch_bounds__(kThreads) + __cluster_dims__(1, kCluster, 1) void all_reduce_mhc_kernel( + const Pack* input, const Pack* residual, const float* post, + const float* comb, const float* pre, const Pack* weight, Pack* output, + Pack* normalized, Pack* local, Pack* multicast, uint32_t* epochs, + int rank, int stage_size, float eps) { +#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 900 + cudaGridDependencySynchronize(); + const int token = blockIdx.x; + const int column = blockIdx.y * kThreads + threadIdx.x; + const int idx = token * kPacks + column; + const int size = gridDim.x * kPacks; + const int stage = epochs[0] % 3; + const int dirty = vllm::mnnvl_lamport_dirty_stage(stage); + const int dirty_size = epochs[2 + dirty]; + Pack value = input[idx]; + vllm::store_multimem_lamport_payload( + multicast + stage * stage_size + rank * size + idx, + vllm::sanitize_lamport_payload(value)); + // Multi-token forwards benefit from launching dependents before mHC finishes. + if (gridDim.x > 1) cudaTriggerProgrammaticLaunchCompletion(); + vllm::lamport_cta_arrive(epochs + 1); + for (int i = idx; i < dirty_size; i += size) { + local[dirty * stage_size + i] = vllm::lamport_sentinel(); + } + Pack peers[kRanks]; + bool pending; + do { + vllm::LamportPack packets[kRanks]; + #pragma unroll + for (int peer = 0; peer < kRanks; ++peer) { + auto ptr = local + stage * stage_size + idx + peer * size; + asm volatile("ld.relaxed.sys.global.v4.u32 {%0,%1,%2,%3}, [%4];" + : "=r"(packets[peer].words[0]), "=r"(packets[peer].words[1]), + "=r"(packets[peer].words[2]), "=r"(packets[peer].words[3]) + : "l"(ptr) + : "memory"); + } + pending = false; + #pragma unroll + for (int peer = 0; peer < kRanks; ++peer) { + pending |= vllm::is_lamport_dirty(packets[peer]); + peers[peer] = packets[peer].packed; + } + } while (pending); + auto sum = vllm::upcast(peers[0]); + #pragma unroll + for (int r = 1; r < kRanks; ++r) { + vllm::packed_assign_add(sum, vllm::upcast(peers[r])); + } + // Keep the BF16 boundary between the collective and mHC post. + auto reduced = vllm::upcast(vllm::downcast(sum)); + decltype(sum) streams[4]; + #pragma unroll + for (int r = 0; r < 4; ++r) { + streams[r] = vllm::upcast(residual[(token * 4 + r) * kPacks + column]); + } + decltype(sum) collapsed{}; + #pragma unroll + for (int c = 0; c < 4; ++c) { + decltype(sum) mixed; + #pragma unroll + for (int e = 0; e < Pack::size; ++e) { + // Match post*x followed by the four residual FMAs. + float acc = __fmul_rn(post[token * 4 + c], reduced.data[e]); + #pragma unroll + for (int r = 0; r < 4; ++r) { + acc = fmaf(comb[token * 16 + r * 4 + c], streams[r].data[e], acc); + } + mixed.data[e] = acc; + } + Pack rounded = vllm::downcast(mixed); + output[(token * 4 + c) * kPacks + column] = rounded; + auto as_float = vllm::upcast(rounded); + #pragma unroll + for (int e = 0; e < Pack::size; ++e) { + collapsed.data[e] = + fmaf(pre[token * 4 + c], as_float.data[e], collapsed.data[e]); + } + } + // RMSNorm consumes the rounded BF16 collapse in the unfused path. + collapsed = vllm::upcast(vllm::downcast(collapsed)); + float sqr = 0; + #pragma unroll + for (int e = 0; e < Pack::size; ++e) { + sqr = fmaf(collapsed.data[e], collapsed.data[e], sqr); + } + #pragma unroll + for (int delta = 16; delta; delta /= 2) { + sqr += __shfl_xor_sync(0xffffffff, sqr, delta); + } + __shared__ float warp_sums[kCluster][kThreads / 32]; + auto cluster = cooperative_groups::this_cluster(); + const int lane = threadIdx.x % 32; + if (lane < kCluster) { + *cluster.map_shared_rank(&warp_sums[blockIdx.y][threadIdx.x / 32], lane) = + sqr; + } + cluster.sync(); + float total = 0; + #pragma unroll + for (int c = 0; c < kCluster; ++c) { + #pragma unroll + for (int w = 0; w < kThreads / 32; ++w) total += warp_sums[c][w]; + } + const float scale = rsqrtf(__fadd_rn(total / kHidden, eps)); + auto weights = vllm::upcast(weight[column]); + #pragma unroll + for (int e = 0; e < Pack::size; ++e) { + collapsed.data[e] = collapsed.data[e] * scale * weights.data[e]; + } + normalized[idx] = vllm::downcast(collapsed); + if (idx == 0) { + while (*reinterpret_cast(epochs + 1) < + gridDim.x * gridDim.y) { + } + epochs[2 + stage] = kRanks * size; + epochs[0] = vllm::mnnvl_lamport_next_stage(stage); + epochs[1] = 0; + } + if (gridDim.x == 1) cudaTriggerProgrammaticLaunchCompletion(); +#else + asm volatile("trap;"); +#endif +} + +void all_reduce_mhc(Tensor input, Tensor residual, Tensor post, Tensor comb, + Tensor pre, Tensor weight, Tensor output, Tensor normalized, + int64_t local, int64_t multicast, Tensor epochs, + int64_t rank, int64_t stage_bytes, double eps) { + using torch::headeronly::ScalarType; + const torch::stable::accelerator::DeviceGuard guard(input.get_device_index()); + const int64_t n = input.size(0); + STD_TORCH_CHECK(input.dim() == 2 && input.size(1) == kHidden && n > 0 && + n <= 16); + for (const auto& t : {input, residual, weight, output, normalized}) { + STD_TORCH_CHECK(t.is_contiguous() && + t.scalar_type() == ScalarType::BFloat16); + STD_TORCH_CHECK(t.get_device_index() == input.get_device_index()); + STD_TORCH_CHECK(reinterpret_cast(t.const_data_ptr()) % 16 == 0); + } + for (const auto& t : {post, comb, pre}) { + STD_TORCH_CHECK(t.is_contiguous() && t.scalar_type() == ScalarType::Float); + STD_TORCH_CHECK(t.get_device_index() == input.get_device_index()); + } + STD_TORCH_CHECK(residual.numel() == n * 4 * kHidden && + output.numel() == residual.numel() && + normalized.numel() == input.numel() && + weight.numel() == kHidden); + STD_TORCH_CHECK(post.numel() == n * 4 && pre.numel() == n * 4 && + comb.numel() == n * 16); + STD_TORCH_CHECK(epochs.is_contiguous() && + epochs.scalar_type() == ScalarType::Int && + epochs.numel() >= 5 && + epochs.get_device_index() == input.get_device_index()); + STD_TORCH_CHECK(local && multicast && rank >= 0 && rank < kRanks && + stage_bytes % 16 == 0 && + stage_bytes >= n * kHidden * 2 * kRanks); + STD_TORCH_CHECK(get_device_prop()->major == 10); + cudaLaunchAttribute attr{}; + attr.id = cudaLaunchAttributeProgrammaticStreamSerialization; + attr.val.programmaticStreamSerializationAllowed = 1; + cudaLaunchConfig_t config{ + .gridDim = dim3(n, kCluster), + .blockDim = dim3(kThreads), + .dynamicSmemBytes = 0, + .stream = get_current_cuda_stream(input.get_device_index()), + .attrs = &attr, + .numAttrs = 1}; + STD_CUDA_CHECK(cudaLaunchKernelEx( + &config, all_reduce_mhc_kernel, + static_cast(input.const_data_ptr()), + static_cast(residual.const_data_ptr()), + static_cast(post.const_data_ptr()), + static_cast(comb.const_data_ptr()), + static_cast(pre.const_data_ptr()), + static_cast(weight.const_data_ptr()), + static_cast(output.mutable_data_ptr()), + static_cast(normalized.mutable_data_ptr()), + reinterpret_cast(local), reinterpret_cast(multicast), + static_cast(epochs.mutable_data_ptr()), static_cast(rank), + static_cast(stage_bytes / sizeof(Pack)), static_cast(eps))); +} +} // namespace + +STABLE_TORCH_LIBRARY_FRAGMENT(_C_custom_ar, mhc) { + mhc.def( + "all_reduce_mhc(Tensor input, Tensor residual, Tensor post, Tensor comb, " + "Tensor pre, Tensor weight, Tensor! output, Tensor! normalized, " + "int local, int multicast, Tensor! epochs, int rank, int stage_bytes, " + "float eps) -> ()"); +} +STABLE_TORCH_LIBRARY_IMPL(_C_custom_ar, CUDA, mhc) { + mhc.impl("all_reduce_mhc", TORCH_BOX(&all_reduce_mhc)); +} diff --git a/csrc/libtorch_stable/cuda_vec_utils.cuh b/csrc/libtorch_stable/cuda_vec_utils.cuh index ec6e60724e..c962c99ec2 100644 --- a/csrc/libtorch_stable/cuda_vec_utils.cuh +++ b/csrc/libtorch_stable/cuda_vec_utils.cuh @@ -320,10 +320,6 @@ __host__ __device__ __forceinline__ bool is_16byte_aligned(const void* ptr) { return (reinterpret_cast(ptr) & 15) == 0; } -__host__ __device__ __forceinline__ bool is_32byte_aligned(const void* ptr) { - return (reinterpret_cast(ptr) & 31) == 0; -} - // ============================================================ // Packed type conversion and arithmetic // ============================================================ diff --git a/csrc/libtorch_stable/moe/moe_align_sum_kernels.cu b/csrc/libtorch_stable/moe/moe_align_sum_kernels.cu index d66b78fe83..5c9a522feb 100644 --- a/csrc/libtorch_stable/moe/moe_align_sum_kernels.cu +++ b/csrc/libtorch_stable/moe/moe_align_sum_kernels.cu @@ -236,7 +236,7 @@ __device__ void _moe_align_block_size_small_batch_expert( int32_t* __restrict__ expert_map, int32_t num_experts, int32_t block_size, size_t numel, int32_t max_num_tokens_padded, int32_t max_num_m_blocks, int32_t inactive_expert_id, int32_t model_offset, int32_t topk_num, - int32_t* token_mask, bool has_expert_map) { + int32_t* token_mask, bool has_expert_map, int32_t* scatter_idx = nullptr) { // Compute input buffer offsets. Typically these will all be 0, except when // using Multi LoRA. int sorted_token_ids_offset = max_num_tokens_padded * model_offset; @@ -318,9 +318,14 @@ __device__ void _moe_align_block_size_small_batch_expert( } for (size_t i = tid; i < numel; i += stride) { - if (int expert_id = get_local_expert_id(i, topk_ids, expert_map, - num_experts, has_expert_map); - expert_id != -1) { + int expert_id = get_local_expert_id(i, topk_ids, expert_map, num_experts, + has_expert_map); + bool valid = + expert_id != -1 && (token_mask == nullptr || token_mask[i / topk_num]); + if (scatter_idx != nullptr) { + scatter_idx[i] = valid ? static_cast(i) : -1; + } + if (expert_id != -1) { int32_t rank_post_pad = tokens_cnts[tid * num_experts + expert_id] + cumsum[expert_id]; @@ -338,14 +343,20 @@ __device__ void _count_and_sort_expert_tokens( int32_t* __restrict__ sorted_token_ids, int32_t* __restrict__ cumsum_buffer, int32_t* __restrict__ expert_map, size_t numel, int32_t num_experts, int32_t max_num_tokens_padded, int32_t* __restrict__ token_mask, - int32_t model_offset, int32_t topk_num, bool has_expert_map) { + int32_t model_offset, int32_t topk_num, bool has_expert_map, + int32_t* scatter_idx = nullptr) { const size_t tid = blockIdx.y * blockDim.x + threadIdx.x; const size_t stride = blockDim.x * gridDim.y; for (size_t i = tid; i < numel; i += stride) { - if (int expert_id = get_local_expert_id(i, topk_ids, expert_map, - num_experts, has_expert_map); - expert_id != -1) { + int expert_id = get_local_expert_id(i, topk_ids, expert_map, num_experts, + has_expert_map); + bool valid = + expert_id != -1 && (token_mask == nullptr || token_mask[i / topk_num]); + if (scatter_idx != nullptr) { + scatter_idx[i] = valid ? static_cast(i) : -1; + } + if (expert_id != -1) { if (token_mask == nullptr || token_mask[i / topk_num]) { int32_t rank_post_pad = atomicAdd( &cumsum_buffer[(model_offset * (num_experts + 1)) + expert_id], 1); @@ -377,10 +388,11 @@ __global__ void count_and_sort_expert_tokens_kernel( const scalar_t* __restrict__ topk_ids, int32_t* __restrict__ sorted_token_ids, int32_t* __restrict__ cumsum_buffer, int32_t* __restrict__ expert_map, size_t numel, int32_t num_experts, - int32_t max_num_tokens_padded, int32_t topk_num, bool has_expert_map) { + int32_t max_num_tokens_padded, int32_t topk_num, bool has_expert_map, + int32_t* scatter_idx) { _count_and_sort_expert_tokens( topk_ids, sorted_token_ids, cumsum_buffer, expert_map, numel, num_experts, - max_num_tokens_padded, nullptr, 0, topk_num, has_expert_map); + max_num_tokens_padded, nullptr, 0, topk_num, has_expert_map, scatter_idx); } // Reduce the topk expert outputs per token (summed in fp32). The output is @@ -533,12 +545,12 @@ __global__ void moe_align_block_size_small_batch_expert_kernel( int32_t* __restrict__ total_tokens_post_pad, int32_t* __restrict__ expert_map, int32_t num_experts, int32_t block_size, size_t numel, int32_t max_num_tokens_padded, int32_t topk_num, - bool has_expert_map) { + bool has_expert_map, int32_t* scatter_idx) { _moe_align_block_size_small_batch_expert( topk_ids, sorted_token_ids, expert_ids, total_tokens_post_pad, expert_map, num_experts, block_size, numel, max_num_tokens_padded, CEILDIV(max_num_tokens_padded, block_size), -1, 0, topk_num, nullptr, - has_expert_map); + has_expert_map, scatter_idx); } template @@ -658,16 +670,31 @@ __global__ void moe_lora_align_block_size_small_batch_expert_kernel( // taken from // https://github.com/sgl-project/sglang/blob/8b5f83ed3b7d2a49ad5c5cd5aa61c5d502f47dbc -void moe_align_block_size( - torch::stable::Tensor topk_ids, int64_t num_experts, int64_t block_size, - torch::stable::Tensor sorted_token_ids, torch::stable::Tensor experts_ids, - torch::stable::Tensor num_tokens_post_pad, - std::optional maybe_expert_map) { +void moe_align_block_size(torch::stable::Tensor topk_ids, int64_t num_experts, + int64_t block_size, + torch::stable::Tensor sorted_token_ids, + torch::stable::Tensor experts_ids, + torch::stable::Tensor num_tokens_post_pad, + std::optional maybe_expert_map, + std::optional scatter_idx) { const torch::stable::accelerator::DeviceGuard device_guard( topk_ids.get_device_index()); const cudaStream_t stream = get_current_cuda_stream(topk_ids.get_device_index()); + int32_t* scatter_ptr = nullptr; + if (scatter_idx.has_value()) { + STD_TORCH_CHECK( + scatter_idx->scalar_type() == torch::headeronly::ScalarType::Int, + "scatter_idx must be int32"); + STD_TORCH_CHECK(scatter_idx->device() == topk_ids.device() && + scatter_idx->is_contiguous() && + scatter_idx->numel() == topk_ids.numel(), + "scatter_idx must be contiguous, on the same device, and " + "match topk_ids.numel()"); + scatter_ptr = reinterpret_cast(scatter_idx->mutable_data_ptr()); + } + int64_t padded_num_experts = ((num_experts + WARP_SIZE - 1) / WARP_SIZE) * WARP_SIZE; int experts_per_warp = WARP_SIZE; @@ -713,7 +740,8 @@ void moe_align_block_size( num_tokens_post_pad.mutable_data_ptr()), reinterpret_cast(expert_map.mutable_data_ptr()), num_experts, block_size, topk_ids.numel(), - sorted_token_ids.size(0), topk_ids.size(1), has_expert_map); + sorted_token_ids.size(0), topk_ids.size(1), has_expert_map, + scatter_ptr); } else { torch::stable::Tensor cumsum_buffer = torch::stable::new_empty( topk_ids, {num_experts + 1}, torch::headeronly::ScalarType::Int); @@ -753,7 +781,7 @@ void moe_align_block_size( reinterpret_cast(cumsum_buffer.mutable_data_ptr()), reinterpret_cast(expert_map.mutable_data_ptr()), topk_ids.numel(), num_experts, sorted_token_ids.size(0), - topk_ids.size(1), has_expert_map); + topk_ids.size(1), has_expert_map, scatter_ptr); } }); } diff --git a/csrc/libtorch_stable/moe/moe_ops.h b/csrc/libtorch_stable/moe/moe_ops.h index 03057c3b14..549a19e981 100644 --- a/csrc/libtorch_stable/moe/moe_ops.h +++ b/csrc/libtorch_stable/moe/moe_ops.h @@ -36,11 +36,13 @@ void moe_sum(torch::stable::Tensor& input, torch::stable::Tensor& output, std::optional topk_ids, std::optional expert_map); -void moe_align_block_size( - torch::stable::Tensor topk_ids, int64_t num_experts, int64_t block_size, - torch::stable::Tensor sorted_token_ids, torch::stable::Tensor experts_ids, - torch::stable::Tensor num_tokens_post_pad, - std::optional maybe_expert_map); +void moe_align_block_size(torch::stable::Tensor topk_ids, int64_t num_experts, + int64_t block_size, + torch::stable::Tensor sorted_token_ids, + torch::stable::Tensor experts_ids, + torch::stable::Tensor num_tokens_post_pad, + std::optional maybe_expert_map, + std::optional scatter_idx); void batched_moe_align_block_size( int64_t max_tokens_per_batch, int64_t block_size, diff --git a/csrc/libtorch_stable/moe/moe_permute_unpermute_op.cu b/csrc/libtorch_stable/moe/moe_permute_unpermute_op.cu index 251097456f..957cd60e49 100644 --- a/csrc/libtorch_stable/moe/moe_permute_unpermute_op.cu +++ b/csrc/libtorch_stable/moe/moe_permute_unpermute_op.cu @@ -56,20 +56,18 @@ int64_t moe_permute_sort_workspace_size(int64_t num_expanded_rows, CubKeyValueSorter::getWorkspaceSize(num_expanded_rows, n_expert)); } -void moe_permute_impl( - const torch::stable::Tensor& input, // [n_token, hidden] - const torch::stable::Tensor& topk_ids, // [n_token, topk] - const torch::stable::Tensor& token_expert_indices, // [n_token, topk] +torch::stable::Tensor moe_sort_routing( + const torch::stable::Tensor& topk_ids, // [n_token, topk] + const torch::stable::Tensor& token_expert_indices, // [n_token, topk] const std::optional& expert_map, // [n_expert] int64_t n_expert, int64_t n_local_expert, int64_t topk, - torch::stable::Tensor& permuted_input, // [permuted_size, hidden] torch::stable::Tensor& expert_first_token_offset, // [n_local_expert + 1] torch::stable::Tensor& inv_permuted_idx, // [n_token, topk] - torch::stable::Tensor& permuted_idx, // [permute_size] const std::optional& maybe_sort_workspace, const std::optional& maybe_permuted_experts_id, const std::optional& maybe_sorted_row_idx, - const std::optional& maybe_topk_ids_for_sort) { + const std::optional& maybe_topk_ids_for_sort, + bool build_inverse) { STD_TORCH_CHECK(expert_first_token_offset.scalar_type() == torch::headeronly::ScalarType::Long, "expert_first_token_offset must be int64"); @@ -88,12 +86,11 @@ void moe_permute_impl( "token_expert_indices shape must be same as inv_permuted_idx"); const torch::stable::accelerator::DeviceGuard device_guard( - input.get_device_index()); - auto device = input.device(); - auto n_token = input.sizes()[0]; - auto n_hidden = input.sizes()[1]; + topk_ids.get_device_index()); + auto device = topk_ids.device(); + auto n_token = topk_ids.sizes()[0]; auto expanded_rows = n_token * topk; - auto stream = get_current_cuda_stream(input.get_device_index()); + auto stream = get_current_cuda_stream(topk_ids.get_device_index()); auto sorter_size = moe_permute_sort_workspace_size(expanded_rows, n_expert); auto sort_workspace = maybe_allocate_tensor( @@ -107,13 +104,11 @@ void moe_permute_impl( torch::headeronly::ScalarType::Int, device, "sorted_row_idx"); CubKeyValueSorter sorter{}; - int64_t* valid_num_ptr = nullptr; torch::stable::Tensor topk_ids_for_sort = topk_ids; - if (expert_map.has_value()) { - const int* expert_map_ptr = get_ptr(expert_map.value()); - valid_num_ptr = - get_ptr(expert_first_token_offset) + n_local_expert; + if (expert_map.has_value() || build_inverse) { + const int* expert_map_ptr = + expert_map.has_value() ? get_ptr(expert_map.value()) : nullptr; topk_ids_for_sort = maybe_allocate_tensor( maybe_topk_ids_for_sort, topk_ids.sizes(), torch::headeronly::ScalarType::Int, device, "topk_ids_for_sort"); @@ -126,7 +121,39 @@ void moe_permute_impl( get_ptr(topk_ids_for_sort), get_ptr(token_expert_indices), get_ptr(permuted_experts_id), get_ptr(sorted_row_idx), get_ptr(expert_first_token_offset), n_token, n_expert, - n_local_expert, topk, sorter, get_ptr(sort_workspace), stream); + n_local_expert, topk, sorter, get_ptr(sort_workspace), stream, + build_inverse ? get_ptr(inv_permuted_idx) : nullptr); + return sorted_row_idx; +} + +void moe_permute_impl( + const torch::stable::Tensor& input, // [n_token, hidden] + const torch::stable::Tensor& topk_ids, // [n_token, topk] + const torch::stable::Tensor& token_expert_indices, // [n_token, topk] + const std::optional& expert_map, // [n_expert] + int64_t n_expert, int64_t n_local_expert, int64_t topk, + torch::stable::Tensor& permuted_input, // [permuted_size, hidden] + torch::stable::Tensor& expert_first_token_offset, // [n_local_expert + 1] + torch::stable::Tensor& inv_permuted_idx, // [n_token, topk] + torch::stable::Tensor& permuted_idx, // [permute_size] + const std::optional& maybe_sort_workspace, + const std::optional& maybe_permuted_experts_id, + const std::optional& maybe_sorted_row_idx, + const std::optional& maybe_topk_ids_for_sort) { + const torch::stable::accelerator::DeviceGuard device_guard( + input.get_device_index()); + auto sorted_row_idx = moe_sort_routing( + topk_ids, token_expert_indices, expert_map, n_expert, n_local_expert, + topk, expert_first_token_offset, inv_permuted_idx, maybe_sort_workspace, + maybe_permuted_experts_id, maybe_sorted_row_idx, maybe_topk_ids_for_sort, + false); + auto stream = get_current_cuda_stream(input.get_device_index()); + auto n_token = input.size(0); + auto n_hidden = input.size(1); + auto valid_num_ptr = + expert_map.has_value() + ? get_ptr(expert_first_token_offset) + n_local_expert + : nullptr; MOE_DISPATCH(input.scalar_type(), [&] { expandInputRowsKernelLauncher( @@ -171,6 +198,22 @@ void moe_permute_with_scratch( topk_ids_for_sort); } +void moe_prepare_scatter(const torch::stable::Tensor& topk_ids, + const torch::stable::Tensor& token_expert_indices, + const std::optional& expert_map, + int64_t n_expert, int64_t n_local_expert, + torch::stable::Tensor& expert_first_token_offset, + torch::stable::Tensor& scatter_idx, + torch::stable::Tensor& sort_workspace, + torch::stable::Tensor& sorted_experts, + torch::stable::Tensor& sorted_rows, + torch::stable::Tensor& topk_ids_for_sort) { + moe_sort_routing(topk_ids, token_expert_indices, expert_map, n_expert, + n_local_expert, topk_ids.size(1), expert_first_token_offset, + scatter_idx, sort_workspace, sorted_experts, sorted_rows, + topk_ids_for_sort, true); +} + void moe_unpermute( const torch::stable::Tensor& permuted_hidden_states, // [n_token * topk, hidden] @@ -315,6 +358,19 @@ void moe_unpermute( STD_TORCH_CHECK(false, "moe_unpermute is not supported on CUDA < 12.0"); } +void moe_prepare_scatter(const torch::stable::Tensor& topk_ids, + const torch::stable::Tensor& token_expert_indices, + const std::optional& expert_map, + int64_t n_expert, int64_t n_local_expert, + torch::stable::Tensor& expert_first_token_offset, + torch::stable::Tensor& scatter_idx, + torch::stable::Tensor& sort_workspace, + torch::stable::Tensor& sorted_experts, + torch::stable::Tensor& sorted_rows, + torch::stable::Tensor& topk_ids_for_sort) { + STD_TORCH_CHECK(false, "moe_prepare_scatter is not supported on CUDA < 12.0"); +} + #endif bool moe_permute_unpermute_supported() { @@ -328,5 +384,6 @@ bool moe_permute_unpermute_supported() { STABLE_TORCH_LIBRARY_IMPL(_moe_C, CUDA, m) { m.impl("moe_permute", TORCH_BOX(&moe_permute)); m.impl("moe_permute_with_scratch", TORCH_BOX(&moe_permute_with_scratch)); + m.impl("moe_prepare_scatter", TORCH_BOX(&moe_prepare_scatter)); m.impl("moe_unpermute", TORCH_BOX(&moe_unpermute)); } diff --git a/csrc/libtorch_stable/moe/moe_wna16_utils.h b/csrc/libtorch_stable/moe/moe_wna16_utils.h index 8ef03f0e60..6ce52ee5b1 100644 --- a/csrc/libtorch_stable/moe/moe_wna16_utils.h +++ b/csrc/libtorch_stable/moe/moe_wna16_utils.h @@ -34,10 +34,6 @@ class ScalarType { static __host__ __device__ float2 inline num22float2(const half2 x) { return __half22float2(x); } - - static __host__ __device__ half2 inline float22num2(const float2 x) { - return __float22half2_rn(x); - } }; template <> @@ -71,10 +67,6 @@ class ScalarType { static __host__ __device__ float2 inline num22float2(const nv_bfloat162 x) { return __bfloat1622float2(x); } - - static __host__ __device__ nv_bfloat162 inline float22num2(const float2 x) { - return __float22bfloat162_rn(x); - } #endif }; diff --git a/csrc/libtorch_stable/moe/permute_unpermute_kernels/moe_permute_unpermute_kernel.cu b/csrc/libtorch_stable/moe/permute_unpermute_kernels/moe_permute_unpermute_kernel.cu index ec7851da3d..8d8bdf8baa 100644 --- a/csrc/libtorch_stable/moe/permute_unpermute_kernels/moe_permute_unpermute_kernel.cu +++ b/csrc/libtorch_stable/moe/permute_unpermute_kernels/moe_permute_unpermute_kernel.cu @@ -82,34 +82,40 @@ __device__ inline int64_t findTotalEltsLessThanTarget(T const* sorted_indices, return target_location + 1; } -// Calculates the start offset of the tokens for a given expert. The last -// element is the total number of valid tokens -__global__ void computeExpertFirstTokenOffsetKernel( +// Computes expert offsets (ending with the valid row count) and optional +// inverse indices. +__global__ void computeExpertOffsetsAndInverseKernel( int const* sorted_experts, int64_t const sorted_experts_len, - int const num_experts, int64_t* expert_first_token_offset) { - // First, compute the global tid. We only need 1 thread per expert. - int const expert = blockIdx.x * blockDim.x + threadIdx.x; + int const num_experts, int64_t* expert_first_token_offset, + int const* sorted_rows, int* inverse) { + int const index = blockIdx.x * blockDim.x + threadIdx.x; // Note that expert goes [0, num_experts] (inclusive) because we want a count // for the total number of active tokens at the end of the scan. - if (expert >= num_experts + 1) { - return; + if (index < num_experts + 1) { + expert_first_token_offset[index] = + findTotalEltsLessThanTarget(sorted_experts, sorted_experts_len, index); + } + if (inverse != nullptr && index < sorted_experts_len) { + inverse[sorted_rows[index]] = index; } - expert_first_token_offset[expert] = - findTotalEltsLessThanTarget(sorted_experts, sorted_experts_len, expert); } -void computeExpertFirstTokenOffset(int const* sorted_indices, - int const total_indices, - int const num_experts, - int64_t* expert_first_token_offset, - cudaStream_t stream) { - int const num_entries = num_experts + 1; - int const threads = std::min(1024, num_entries); +void computeExpertOffsetsAndInverse(int const* sorted_indices, + int const total_indices, + int const num_experts, + int64_t* expert_first_token_offset, + int const* sorted_rows, int* inverse, + cudaStream_t stream) { + int const num_entries = inverse != nullptr + ? std::max(num_experts + 1, total_indices) + : num_experts + 1; + int const threads = inverse != nullptr ? 256 : std::min(1024, num_entries); int const blocks = (num_entries + threads - 1) / threads; - computeExpertFirstTokenOffsetKernel<<>>( - sorted_indices, total_indices, num_experts, expert_first_token_offset); + computeExpertOffsetsAndInverseKernel<<>>( + sorted_indices, total_indices, num_experts, expert_first_token_offset, + sorted_rows, inverse); } void sortAndScanExpert(const int* expert_for_source_row, const int* source_rows, @@ -117,7 +123,7 @@ void sortAndScanExpert(const int* expert_for_source_row, const int* source_rows, int64_t* expert_first_token_offset, int num_rows, int num_experts, int num_experts_per_node, int k, CubKeyValueSorter& sorter, void* sorter_ws, - cudaStream_t stream) { + cudaStream_t stream, int* inverse) { int64_t const expanded_num_rows = static_cast(k) * num_rows; // We need to use the full num_experts because that is the sentinel value used // by topk for disabled experts @@ -127,9 +133,9 @@ void sortAndScanExpert(const int* expert_for_source_row, const int* source_rows, sorter.run((void*)sorter_ws, sorter_ws_size_bytes, expert_for_source_row, permuted_experts, source_rows, permuted_rows, expanded_num_rows, stream); - computeExpertFirstTokenOffset(permuted_experts, expanded_num_rows, - num_experts_per_node, expert_first_token_offset, - stream); + computeExpertOffsetsAndInverse( + permuted_experts, expanded_num_rows, num_experts_per_node, + expert_first_token_offset, permuted_rows, inverse, stream); } __global__ void preprocessTopkIdKernel(int* topk_id_ptr, int size, @@ -142,7 +148,7 @@ __global__ void preprocessTopkIdKernel(int* topk_id_ptr, int size, extern __shared__ int smem_expert_map[]; // store expert_map in smem for (int i = tidx; i < num_experts; i += blockDim.x) { - smem_expert_map[i] = expert_map_ptr[i]; + smem_expert_map[i] = expert_map_ptr ? expert_map_ptr[i] : i; } __syncthreads(); diff --git a/csrc/libtorch_stable/moe/permute_unpermute_kernels/moe_permute_unpermute_kernel.h b/csrc/libtorch_stable/moe/permute_unpermute_kernels/moe_permute_unpermute_kernel.h index 89c278a4ed..53afad2f5d 100644 --- a/csrc/libtorch_stable/moe/permute_unpermute_kernels/moe_permute_unpermute_kernel.h +++ b/csrc/libtorch_stable/moe/permute_unpermute_kernels/moe_permute_unpermute_kernel.h @@ -43,18 +43,19 @@ class CubKeyValueSorter { int num_bits_; }; -void computeExpertFirstTokenOffset(int const* sorted_indices, - int const total_indices, - int const num_experts, - int64_t* expert_first_token_offset, - cudaStream_t stream); +void computeExpertOffsetsAndInverse(int const* sorted_indices, + int const total_indices, + int const num_experts, + int64_t* expert_first_token_offset, + int const* sorted_rows, int* inverse, + cudaStream_t stream); void sortAndScanExpert(const int* expert_for_source_row, const int* source_rows, int* permuted_experts, int* permuted_rows, int64_t* expert_first_token_offset, int num_rows, int num_experts, int num_experts_per_node, int k, CubKeyValueSorter& sorter, void* sorter_ws, - cudaStream_t stream); + cudaStream_t stream, int* inverse = nullptr); template void expandInputRowsKernelLauncher( diff --git a/csrc/libtorch_stable/moe/torch_bindings.cpp b/csrc/libtorch_stable/moe/torch_bindings.cpp index c0e966b2be..2feec38389 100644 --- a/csrc/libtorch_stable/moe/torch_bindings.cpp +++ b/csrc/libtorch_stable/moe/torch_bindings.cpp @@ -39,7 +39,8 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_moe_C, m) { " int block_size, Tensor! sorted_token_ids," " Tensor! experts_ids," " Tensor! num_tokens_post_pad," - " Tensor? maybe_expert_map) -> ()"); + " Tensor? maybe_expert_map," + " Tensor(a!)? scatter_idx=None) -> ()"); // Aligning the number of tokens to be processed by each expert such // that it is divisible by the block size, but for the batched case. @@ -105,6 +106,13 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_moe_C, m) { "permuted_idx, Tensor! sort_workspace, Tensor! permuted_experts_id, " "Tensor! sorted_row_idx, Tensor! topk_ids_for_sort)->()"); + m.def( + "moe_prepare_scatter(Tensor topk_ids, Tensor token_expert_indices, " + "Tensor? expert_map, int n_expert, int n_local_expert, " + "Tensor! expert_first_token_offset, Tensor! scatter_idx, " + "Tensor! sort_workspace, Tensor! sorted_experts, Tensor! sorted_rows, " + "Tensor! topk_ids_for_sort)->()"); + m.def( "moe_unpermute(Tensor permuted_hidden_states, Tensor topk_weights," "Tensor inv_permuted_idx, Tensor? expert_first_token_offset, " diff --git a/csrc/libtorch_stable/quantization/gptq/matrix_view.cuh b/csrc/libtorch_stable/quantization/gptq/matrix_view.cuh index 3d82e43931..abe3368cb1 100644 --- a/csrc/libtorch_stable/quantization/gptq/matrix_view.cuh +++ b/csrc/libtorch_stable/quantization/gptq/matrix_view.cuh @@ -74,18 +74,9 @@ class MatrixView_half_rw { const int width) : data(data), height(height), width(width) {} - __device__ __forceinline__ half item(int row, int column) const { - return data[row * width + column]; - } __device__ __forceinline__ half* item_ptr(int row, int column) { return &data[row * width + column]; } - __device__ __forceinline__ void set(int row, int column, half value) { - data[row * width + column] = value; - } - __device__ __forceinline__ void set_half2(int row, int column, half2 value) { - ((half2*)data)[(row * width + column) / 2] = value; - } __device__ __forceinline__ void set4(int row, int column, half v0, half v1, half v2, half v3) { @@ -113,14 +104,6 @@ class MatrixView_q4_row { return (data[row * width / 8 + column / 8] >> shift) & 0x0f; } - __device__ __forceinline__ void item2(int (&items)[2], int row, - int column) const { - int shift = (column & 0x07) * 4; - uint32_t d = data[row * width / 8 + column / 8] >> shift; - items[0] = d & 0x0f; - items[1] = (d >> 4) & 0x0f; - } - __device__ __forceinline__ void item4(int (&items)[4], int row, int column) const { int shift = (column & 0x07) * 4; @@ -148,14 +131,6 @@ class MatrixView_q2_row { return (data[row * width / 16 + column / 16] >> shift) & 0x03; } - __device__ __forceinline__ void item2(int (&items)[2], int row, - int column) const { - int shift = (column & 0x0f) * 2; - uint32_t d = data[row * width / 16 + column / 16] >> shift; - items[0] = d & 0x03; - items[1] = (d >> 2) & 0x03; - } - __device__ __forceinline__ void item4(int (&items)[4], int row, int column) const { int shift = (column & 0x0f) * 2; @@ -237,14 +212,6 @@ class MatrixView_q8_row { return (data[row * width / 4 + column / 4] >> shift) & 0xff; } - __device__ __forceinline__ void item2(int (&items)[2], int row, - int column) const { - int shift = (column & 0x03) * 8; - uint32_t d = data[row * width / 4 + column / 4] >> shift; - items[0] = d & 0xff; - items[1] = (d >> 8) & 0xff; - } - __device__ __forceinline__ void item4(int (&items)[4], int row, int column) const { int shift = (column & 0x03) * 2; diff --git a/csrc/libtorch_stable/quantization/marlin/marlin.cuh b/csrc/libtorch_stable/quantization/marlin/marlin.cuh index bfb65e874b..af4d4730ba 100644 --- a/csrc/libtorch_stable/quantization/marlin/marlin.cuh +++ b/csrc/libtorch_stable/quantization/marlin/marlin.cuh @@ -67,14 +67,6 @@ __device__ inline void cp_async2_ca_pred(void* smem_ptr, const void* glob_ptr, } } -__device__ inline void cp_async4_ca_pred(void* smem_ptr, const void* glob_ptr, - bool pred = true) { - if (pred) { - reinterpret_cast(smem_ptr)[0] = - reinterpret_cast(glob_ptr)[0]; - } -} - __device__ inline void cp_async4_pred(void* smem_ptr, const void* glob_ptr, bool pred = true) { if (pred) { @@ -121,19 +113,6 @@ __device__ inline void cp_async2_ca_pred(void* smem_ptr, const void* glob_ptr, "r"(smem), "l"(glob_ptr), "n"(BYTES)); } -__device__ inline void cp_async4_ca_pred(void* smem_ptr, const void* glob_ptr, - bool pred = true) { - const int BYTES = 16; - uint32_t smem = static_cast(__cvta_generic_to_shared(smem_ptr)); - asm volatile( - "{\n" - " .reg .pred p;\n" - " setp.ne.b32 p, %0, 0;\n" - " @p cp.async.ca.shared.global [%1], [%2], %3;\n" - "}\n" ::"r"((int)pred), - "r"(smem), "l"(glob_ptr), "n"(BYTES)); -} - __device__ inline void cp_async4_pred(void* smem_ptr, const void* glob_ptr, bool pred = true) { const int BYTES = 16; @@ -170,4 +149,4 @@ __device__ inline void cp_async_wait() { } // namespace MARLIN_NAMESPACE_NAME -#endif \ No newline at end of file +#endif diff --git a/csrc/libtorch_stable/sampler.cu b/csrc/libtorch_stable/sampler.cu index b8d64f7af1..703a748ce6 100644 --- a/csrc/libtorch_stable/sampler.cu +++ b/csrc/libtorch_stable/sampler.cu @@ -45,11 +45,6 @@ __global__ void apply_repetition_penalties_kernel( } } -__device__ __forceinline__ auto convert_to_uint32(float x) -> uint32_t { - uint32_t bits = __float_as_uint(x); - return (bits & 0x80000000) ? bits : ~bits & 0x7fffffff; -} - template static inline __device__ uint32_t extractBinIdx(float x) { if constexpr (step == 0) { diff --git a/csrc/quantization/w8a8/fp8/amd/quant_utils.cuh b/csrc/quantization/w8a8/fp8/amd/quant_utils.cuh index 58d19de734..32925c0c3c 100644 --- a/csrc/quantization/w8a8/fp8/amd/quant_utils.cuh +++ b/csrc/quantization/w8a8/fp8/amd/quant_utils.cuh @@ -67,104 +67,6 @@ using fp8_type = __hip_fp8_e4m3_fnuz; using fp8x2_type = __hip_fp8x2_e4m3_fnuz; #endif -// fp8 -> half -template <> -__inline__ __device__ uint16_t -vec_conversion(const uint8_t& a) { - return __hip_cvt_fp8_to_halfraw(a, fp8_type::__default_interpret).x; -} - -// fp8x2 -> half2 -template <> -__inline__ __device__ uint32_t -vec_conversion(const uint16_t& a) { - union { - __half2_raw h2r; - uint32_t ui32; - } tmp; - tmp.h2r = __hip_cvt_fp8x2_to_halfraw2(a, fp8_type::__default_interpret); - return tmp.ui32; -} - -// fp8x4 -> half2x2 -template <> -__inline__ __device__ uint2 vec_conversion(const uint32_t& a) { - union { - uint2 u32x2; - uint32_t u32[2]; - } tmp; - tmp.u32[0] = vec_conversion((uint16_t)a); - tmp.u32[1] = vec_conversion((uint16_t)(a >> 16U)); - return tmp.u32x2; -} - -// fp8x8 -> half2x4 -template <> -__inline__ __device__ uint4 vec_conversion(const uint2& a) { - union { - uint4 u64x2; - uint2 u64[2]; - } tmp; - tmp.u64[0] = vec_conversion(a.x); - tmp.u64[1] = vec_conversion(a.y); - return tmp.u64x2; -} - -using __nv_bfloat16 = __hip_bfloat16; - -// fp8 -> __nv_bfloat16 -template <> -__inline__ __device__ __nv_bfloat16 -vec_conversion<__nv_bfloat16, uint8_t>(const uint8_t& a) { - fp8_type f8; - f8.__x = a; - return __float2bfloat16(static_cast(f8)); -} - -using __nv_bfloat162 = __hip_bfloat162; - -// fp8x2 -> __nv_bfloat162 -template <> -__inline__ __device__ __nv_bfloat162 -vec_conversion<__nv_bfloat162, uint16_t>(const uint16_t& a) { - __nv_bfloat162 res; - res.x = vec_conversion<__nv_bfloat16, uint8_t>((uint8_t)a); - res.y = vec_conversion<__nv_bfloat16, uint8_t>((uint8_t)(a >> 8U)); - return res; -} - -// fp8x4 -> bf16_4_t -template <> -__inline__ __device__ bf16_4_t -vec_conversion(const uint32_t& a) { - bf16_4_t res; - res.x = vec_conversion<__nv_bfloat162, uint16_t>((uint16_t)a); - res.y = vec_conversion<__nv_bfloat162, uint16_t>((uint16_t)(a >> 16U)); - return res; -} - -// fp8x8 -> bf16_8_t -template <> -__inline__ __device__ bf16_8_t vec_conversion(const uint2& a) { - bf16_4_t tmp1, tmp2; - tmp1 = vec_conversion(a.x); - tmp2 = vec_conversion(a.y); - bf16_8_t res; - res.x = tmp1.x; - res.y = tmp1.y; - res.z = tmp2.x; - res.w = tmp2.y; - return res; -} - -// fp8 -> float -template <> -__inline__ __device__ float vec_conversion(const uint8_t& a) { - fp8_type f8; - f8.__x = a; - return static_cast(f8); -} - // fp8x2 -> float2 template <> __inline__ __device__ float2 @@ -193,138 +95,6 @@ vec_conversion(const uint32_t& a) { return res; } -// fp8x8 -> float8 -template <> -__inline__ __device__ Float8_ vec_conversion(const uint2& a) { - Float4_ tmp1, tmp2; - tmp1 = vec_conversion(a.x); - tmp2 = vec_conversion(a.y); - Float8_ res; - res.x = tmp1.x; - res.y = tmp1.y; - res.z = tmp2.x; - res.w = tmp2.y; - return res; -} - -// half -> fp8 -template <> -__inline__ __device__ uint8_t -vec_conversion(const uint16_t& a) { - __half_raw tmp; - tmp.x = a; - return __hip_cvt_halfraw_to_fp8(tmp, fp8_type::__default_saturation, - fp8_type::__default_interpret); -} - -template <> -__inline__ __device__ uint16_t -vec_conversion(const uint32_t& a) { - union { - uint32_t ui32; - __half2_raw h2r; - } tmp; - tmp.ui32 = a; - return __hip_cvt_halfraw2_to_fp8x2(tmp.h2r, fp8_type::__default_saturation, - fp8_type::__default_interpret); -} - -// bf16 -> fp8 -template <> -__inline__ __device__ uint8_t -vec_conversion(const __nv_bfloat16& a) { - return __hip_cvt_float_to_fp8(__bfloat162float(a), - fp8_type::__default_saturation, - fp8_type::__default_interpret); -} - -// float -> fp8 -template <> -__inline__ __device__ uint8_t vec_conversion(const float& a) { - return __hip_cvt_float_to_fp8(a, fp8_type::__default_saturation, - fp8_type::__default_interpret); -} - -// float2 -> half2 -template <> -__inline__ __device__ uint32_t -vec_conversion(const float2& a) { - union { - half2 float16; - uint32_t uint32; - }; - - float16 = __float22half2_rn(a); - return uint32; -} - -// Float4 -> half2x2 -template <> -__inline__ __device__ uint2 vec_conversion(const Float4_& a) { - uint2 b; - float2 val; - val.x = a.x.x; - val.y = a.x.y; - b.x = vec_conversion(val); - - val.x = a.y.x; - val.y = a.y.y; - b.y = vec_conversion(val); - return b; -} - -// Float4 -> float4 -template <> -__inline__ __device__ float4 vec_conversion(const Float4_& a) { - float4 b; - b.x = a.x.x; - b.y = a.x.y; - b.z = a.y.x; - b.w = a.y.y; - return b; -} - -// Float8 -> half2x4 -template <> -__inline__ __device__ uint4 vec_conversion(const Float8_& a) { - uint4 b; - b.x = vec_conversion(a.x); - b.y = vec_conversion(a.y); - b.z = vec_conversion(a.z); - b.w = vec_conversion(a.w); - return b; -} - -// float2 -> bfloat162 -template <> -__inline__ __device__ __nv_bfloat162 -vec_conversion<__nv_bfloat162, float2>(const float2& a) { - __nv_bfloat162 b = __float22bfloat162_rn(a); - return b; -} - -// Float4 -> bfloat162x2 -template <> -__inline__ __device__ bf16_4_t -vec_conversion(const Float4_& a) { - bf16_4_t b; - b.x = __float22bfloat162_rn(a.x); - b.y = __float22bfloat162_rn(a.y); - return b; -} - -// Float8 -> bfloat162x4 -template <> -__inline__ __device__ bf16_8_t -vec_conversion(const Float8_& a) { - bf16_8_t b; - b.x = __float22bfloat162_rn(a.x); - b.y = __float22bfloat162_rn(a.y); - b.z = __float22bfloat162_rn(a.z); - b.w = __float22bfloat162_rn(a.w); - return b; -} - /* Scaled and vectorized conversions, for data exchange between high and low precision domains @@ -335,6 +105,7 @@ vec_conversion(const Float8_& a) { */ using __nv_bfloat16 = __hip_bfloat16; +using __nv_bfloat162 = __hip_bfloat162; // fp8 -> __nv_bfloat16 template <> @@ -613,17 +384,6 @@ scaled_vec_conversion(const float4& a, float scale) { } #endif // ENABLE_FP8 -template -__inline__ __device__ Tout convert(const Tin& x) { - #ifdef ENABLE_FP8 - if constexpr (kv_dt == Fp8KVCacheDataType::kFp8E4M3) { - return vec_conversion(x); - } - #endif - assert(false); - return {}; // Squash missing return statement warning -} - template __inline__ __device__ Tout scaled_convert(const Tin& x, const float scale) { #ifdef ENABLE_FP8 diff --git a/csrc/quantization/w8a8/fp8/nvidia/quant_utils.cuh b/csrc/quantization/w8a8/fp8/nvidia/quant_utils.cuh index 11a9b0818d..7fcf541fd6 100644 --- a/csrc/quantization/w8a8/fp8/nvidia/quant_utils.cuh +++ b/csrc/quantization/w8a8/fp8/nvidia/quant_utils.cuh @@ -282,19 +282,6 @@ __inline__ __device__ float4 scaled_vec_conversion( } #endif // ENABLE_FP8 -template -__inline__ __device__ Tout convert(const Tin& x) { - #if 0 // Disable the following code to reduce the binary size. - if constexpr (kv_dt == Fp8KVCacheDataType::kFp8E4M3) { - return vec_conversion(x, __NV_E4M3); - } else if constexpr (kv_dt == Fp8KVCacheDataType::kFp8E5M2) { - return vec_conversion(x, __NV_E5M2); - } - #endif - assert(false); - __builtin_unreachable(); // Suppress missing return statement warning -} - template __inline__ __device__ Tout scaled_convert(const Tin& x, const float scale) { #ifdef ENABLE_FP8 diff --git a/csrc/rocm/attention.cu b/csrc/rocm/attention.cu index b747b870e2..2a11653a58 100644 --- a/csrc/rocm/attention.cu +++ b/csrc/rocm/attention.cu @@ -342,22 +342,6 @@ typedef union u64_cvt { int64_t i64; } _T8x8; -__device__ __forceinline__ _B8x8 convert_b16x8(const _B16x8& input, - _T8x8& Mtemp) { - _T8x8 Qtmp8x8; - - for (int i = 0; i < 2; i++) { - floatx4 q_out = {0, 0, 0, 0}; - q_out = gcn_mfma16x16x16_instr<_Float16, 0, 0, 0>(Mtemp.b64, input.xy[i], - q_out); - Qtmp8x8.b16x4[i * 2] = - __builtin_amdgcn_cvt_pk_fp8_f32(q_out[0], q_out[1], 0, false); - Qtmp8x8.b16x4[i * 2 + 1] = - __builtin_amdgcn_cvt_pk_fp8_f32(q_out[2], q_out[3], 0, false); - } - return Qtmp8x8.b8x8; -} - __device__ float warpReduceMax(float val) { for (int offset = warpSize / 2; offset > 0; offset /= 2) { val = max( @@ -2583,23 +2567,6 @@ __device__ __forceinline__ float to_float(const T& inp) { } } -template -__device__ __forceinline__ float to_float_b16(const bit16_t& inp) { - union tmpcvt { - bit16_t u; - _Float16 f; - __hip_bfloat16 b; - } t16; - t16.u = inp; - if constexpr (std::is_same::value) { - return (float)t16.f; - } else if constexpr (std::is_same::value) { - return __bfloat162float(t16.b); - } else { - static_assert(false, "unsupported 16b dtype"); - } -} - template __device__ __forceinline__ T from_float(const float& inp) { if constexpr (std::is_same::value) { diff --git a/csrc/rocm/moe_q_gemm_rdna3.cu b/csrc/rocm/moe_q_gemm_rdna3.cu index 6c25ed7e4b..61ffbb9987 100644 --- a/csrc/rocm/moe_q_gemm_rdna3.cu +++ b/csrc/rocm/moe_q_gemm_rdna3.cu @@ -67,21 +67,6 @@ __forceinline__ __device__ float dot22_8_f(half2 (&dq)[4], const half* a_ptr) { return result; } -__forceinline__ __device__ float dot22_8_f(float (&dq)[8], - const bf16_t* a_ptr) { - float result = 0.0f; - #pragma unroll - for (int i = 0; i < 4; i++) { - uint32_t aw; - __builtin_memcpy(&aw, a_ptr + 2 * i, sizeof(uint32_t)); - float a_x = __uint_as_float((aw & 0xFFFFu) << 16); - float a_y = __uint_as_float(aw & 0xFFFF0000u); - result = __fmaf_rn(dq[2 * i + 0], a_x, result); - result = __fmaf_rn(dq[2 * i + 1], a_y, result); - } - return result; -} - __forceinline__ __device__ void atomic_add_pk4_f16(half* addr, half2 v01, half2 v23) { unsigned long long* addr_u = reinterpret_cast(addr); diff --git a/csrc/rocm/q_gemm_rdna3.cu b/csrc/rocm/q_gemm_rdna3.cu index 104af66d1c..bf23682699 100644 --- a/csrc/rocm/q_gemm_rdna3.cu +++ b/csrc/rocm/q_gemm_rdna3.cu @@ -115,61 +115,6 @@ __forceinline__ __device__ float dot22_8_f(half2 (&dq)[4], const half* a_ptr) { return result; } -__forceinline__ __device__ float dot22_8_f(bf162_t (&dq)[4], - const bf16_t* a_ptr) { - // RDNA3 (gfx1100) lacks a packed bf16 FMA: there is no v_pk_fma_bf16 in - // the gfx11 ISA (it only landed on CDNA3+ / gfx94x and later). hipcc - // therefore lowers __hfma2(bf162_t, bf162_t, bf162_t) to a serialised - // fallback (single-element FMAs or fp32 round-trips), which empirically - // runs ~2× the cycle count of v_pk_fma_f16 on the same VALU. The bf16 - // decode path was paying that tax in full, scaling linearly with M (the - // fp16 path scales sub-linearly because its v_pk_fma_f16 is full rate - // and the kernel becomes memory-bound). - // - // Fix: widen bf16 → fp32 explicitly (a left-shift by 16, free in VGPRs) - // and accumulate with v_fma_f32, which IS full rate on RDNA3. Same FMA - // count, but each FMA is fast. Bonus: the accumulator is now fp32 - // throughout instead of bf16, which is also numerically more accurate - // (no compounding bf16-rounding inside the dot loop). - float result = 0.0f; - #pragma unroll - for (int i = 0; i < 4; i++) { - uint32_t aw, dw; - __builtin_memcpy(&aw, a_ptr + 2 * i, sizeof(uint32_t)); - __builtin_memcpy(&dw, &dq[i], sizeof(uint32_t)); - // bf16 in low 16 bits → fp32 by left-shifting into the upper half. - // bf16 in high 16 bits → already aligned with fp32's upper half. - float a_x = __uint_as_float((aw & 0xFFFFu) << 16); - float a_y = __uint_as_float(aw & 0xFFFF0000u); - float d_x = __uint_as_float((dw & 0xFFFFu) << 16); - float d_y = __uint_as_float(dw & 0xFFFF0000u); - result = __fmaf_rn(d_x, a_x, result); - result = __fmaf_rn(d_y, a_y, result); - } - return result; -} - -// fp32-input dot product: paired with dequant_4bit_8_bf16_f32 which already -// produces fp32 dq[8]. Saves the bf16→fp32 widening that the bf162_t -// overload above does for dq (still need to widen A from bf16). Wins more -// at high N: the bf162_t version's per-call widening cost scales with the -// number of dequants × M_COUNT × 4 dot calls; the fp32 version pays only -// for A widening (M_COUNT × 4 × 4 widens, half as many). -__forceinline__ __device__ float dot22_8_f(float (&dq)[8], - const bf16_t* a_ptr) { - float result = 0.0f; - #pragma unroll - for (int i = 0; i < 4; i++) { - uint32_t aw; - __builtin_memcpy(&aw, a_ptr + 2 * i, sizeof(uint32_t)); - float a_x = __uint_as_float((aw & 0xFFFFu) << 16); - float a_y = __uint_as_float(aw & 0xFFFF0000u); - result = __fmaf_rn(dq[2 * i + 0], a_x, result); - result = __fmaf_rn(dq[2 * i + 1], a_y, result); - } - return result; -} - // --------------------------------------------------------------------------- // Packed atomic-add via CAS-loop on a 64-bit word (4 fp16/bf16 lanes per CAS). // RDNA3 (gfx11) does NOT have native v_global_atomic_pk_add_f16 / _bf16 (those @@ -327,9 +272,8 @@ __global__ void gemm_q4_kernel_rdna3(const T* __restrict__ a, // Per-column dequant constants. We hold one set of (z, y) pairs per column. // fp16 uses the exllama (z1z16, y1y16) double-pair to enable the upper- - // nibble-*16 trick. bf16 uses fp32 scalars (z, y) because the dequant - // produces fp32 directly — see prep_zero_scale_bf16_f32 / the FMA - // bypass for the missing v_pk_fma_bf16 on gfx11. + // nibble-*16 trick. bf16 uses fp32 scalars (z, y) for the inline v_dot2 + // bias correction. half2 z1z16_h[4][2], y1y16_h[4][2]; float z_b_f[4], y_b_f[4]; diff --git a/csrc/rocm/qdq_4_rdna3.cuh b/csrc/rocm/qdq_4_rdna3.cuh index f4be668c1f..43a7b33ba1 100644 --- a/csrc/rocm/qdq_4_rdna3.cuh +++ b/csrc/rocm/qdq_4_rdna3.cuh @@ -113,63 +113,6 @@ __forceinline__ __device__ void dequant_4bit_8_fp16(uint32_t qa, half2 (&dq)[4], // bf16 path // --------------------------------------------------------------------------- -// Bit-trick magic for bf16: -// bf16(128) == 0x4300 (sign 0, exp 134, mantissa 0). -// For nibble n in [0..15], bits [3:0] of mantissa hold n exactly because -// bf16's ULP at 128 is 1 (mantissa step = 2^(7-7) = 1). So -// ((qa & 0x000F000F) | 0x43004300) bitcasts to bfloat162(128+n_lo, 128+n_hi). -// -// Because bf16's mantissa is only 7 bits, we cannot use the fp16 "upper nibble -// * 16" trick. Instead each pair of nibbles is shifted down to [3:0]/[19:16] -// via a single 4/8/12-bit right-shift before the OR. That costs one extra -// shift per pair vs fp16, but keeps the FMA structure identical. -__forceinline__ __device__ void prep_zero_scale_bf16(uint32_t zero, - bf16_t scale, - bf162_t& z_prep, - bf162_t& y_prep) { - // z = scale * -(128 + zero); y = scale. - float scale_f = __bfloat162float(scale); - float zf = -(128.0f + (float)zero) * scale_f; - bf16_t zb = __float2bfloat16(zf); - z_prep = __bfloat162bfloat162(zb); - y_prep = __bfloat162bfloat162(scale); -} - -__forceinline__ __device__ void dequant_4bit_8_bf16(uint32_t qa, - bf162_t (&dq)[4], - bf162_t z_prep, - bf162_t y_prep) { - const uint32_t c0 = 0x43004300; - - union { - uint32_t u; - bf162_t b2; - } q0, q1, q2, q3; - q0.u = ((qa >> 0) & 0x000F000F) | c0; // bf162(128+q[0], 128+q[1]) - q1.u = ((qa >> 4) & 0x000F000F) | c0; // bf162(128+q[2], 128+q[3]) - q2.u = ((qa >> 8) & 0x000F000F) | c0; // bf162(128+q[4], 128+q[5]) - q3.u = ((qa >> 12) & 0x000F000F) | c0; // bf162(128+q[6], 128+q[7]) - - // dq = q_b * scale + (-(128+zero)*scale) = (q - zero) * scale - dq[0] = __hfma2(q0.b2, y_prep, z_prep); - dq[1] = __hfma2(q1.b2, y_prep, z_prep); - dq[2] = __hfma2(q2.b2, y_prep, z_prep); - dq[3] = __hfma2(q3.b2, y_prep, z_prep); -} - -// --------------------------------------------------------------------------- -// bf16-input → fp32-output dequant (RDNA3 scalar path). -// -// RDNA3 (gfx1100) has no v_pk_fma_bf16; packed bf16 FMA lowers to a slow -// fallback. Rather than computing dq in bf16 and widening at FMA time in -// the dot product, we widen to fp32 here once (a free left-shift by 16) and -// emit the (q - zero) * scale FMA directly in fp32. This: -// * Replaces 4× slow bf16 packed FMA with 8× fast fp32 FMA per int32. -// * Eliminates 4× bf16→fp32 widens that the dot product would do. -// * Keeps the dot product accumulator in fp32 without a roundtrip. -// -// Output: fp32 dq[8], one element per K position (consumed by the -// fp32-overload of dot22_8_f in q_gemm_rdna3.cu). __forceinline__ __device__ void prep_zero_scale_bf16_f32(uint32_t zero, bf16_t scale, float& z_prep, @@ -179,60 +122,6 @@ __forceinline__ __device__ void prep_zero_scale_bf16_f32(uint32_t zero, y_prep = scale_f; } -// Pure-q dequant for the M_COUNT=1 factored path: outputs the unscaled fp32 -// values 128+nibble, without folding scale/zero. The caller folds scale/zb -// into the accumulator outside the inner loop using a precomputed sum_a, -// which saves ~27% of the FMA count vs the per-col-dequant approach above -// (only beneficial at M_COUNT=1; break-even at M_COUNT=2). -// -// Cost: 0 FMAs (pure bit-trick + as_float reinterprets). -__forceinline__ __device__ void dequant_4bit_8_bf16_q_only(uint32_t qa, - float (&q_f32)[8]) { - const uint32_t c0 = 0x43004300; - const uint32_t q0 = ((qa >> 0) & 0x000F000F) | c0; - const uint32_t q1 = ((qa >> 4) & 0x000F000F) | c0; - const uint32_t q2 = ((qa >> 8) & 0x000F000F) | c0; - const uint32_t q3 = ((qa >> 12) & 0x000F000F) | c0; - q_f32[0] = __uint_as_float((q0 & 0xFFFFu) << 16); - q_f32[1] = __uint_as_float(q0 & 0xFFFF0000u); - q_f32[2] = __uint_as_float((q1 & 0xFFFFu) << 16); - q_f32[3] = __uint_as_float(q1 & 0xFFFF0000u); - q_f32[4] = __uint_as_float((q2 & 0xFFFFu) << 16); - q_f32[5] = __uint_as_float(q2 & 0xFFFF0000u); - q_f32[6] = __uint_as_float((q3 & 0xFFFFu) << 16); - q_f32[7] = __uint_as_float(q3 & 0xFFFF0000u); -} - -__forceinline__ __device__ void dequant_4bit_8_bf16_f32(uint32_t qa, - float (&dq)[8], - float z_prep, - float y_prep) { - const uint32_t c0 = 0x43004300; - const uint32_t q0 = ((qa >> 0) & 0x000F000F) | c0; - const uint32_t q1 = ((qa >> 4) & 0x000F000F) | c0; - const uint32_t q2 = ((qa >> 8) & 0x000F000F) | c0; - const uint32_t q3 = ((qa >> 12) & 0x000F000F) | c0; - // bf16(128+nibble) bits → fp32(128+nibble) bits via left-shift by 16 - // (just zero-extends the mantissa from 7 to 23 bits; exponent preserved). - const float q0x = __uint_as_float((q0 & 0xFFFFu) << 16); - const float q0y = __uint_as_float(q0 & 0xFFFF0000u); - const float q1x = __uint_as_float((q1 & 0xFFFFu) << 16); - const float q1y = __uint_as_float(q1 & 0xFFFF0000u); - const float q2x = __uint_as_float((q2 & 0xFFFFu) << 16); - const float q2y = __uint_as_float(q2 & 0xFFFF0000u); - const float q3x = __uint_as_float((q3 & 0xFFFFu) << 16); - const float q3y = __uint_as_float(q3 & 0xFFFF0000u); - // dq[i] = q_f32 * scale + (-(128+zero)*scale) = (nibble - zero) * scale - dq[0] = __fmaf_rn(q0x, y_prep, z_prep); - dq[1] = __fmaf_rn(q0y, y_prep, z_prep); - dq[2] = __fmaf_rn(q1x, y_prep, z_prep); - dq[3] = __fmaf_rn(q1y, y_prep, z_prep); - dq[4] = __fmaf_rn(q2x, y_prep, z_prep); - dq[5] = __fmaf_rn(q2y, y_prep, z_prep); - dq[6] = __fmaf_rn(q3x, y_prep, z_prep); - dq[7] = __fmaf_rn(q3y, y_prep, z_prep); -} - } // namespace gptq_rdna3 } // namespace vllm diff --git a/docker/Dockerfile b/docker/Dockerfile index ad010285b1..895968601f 100644 --- a/docker/Dockerfile +++ b/docker/Dockerfile @@ -477,6 +477,24 @@ RUN --mount=type=cache,target=/root/.cache/ccache \ #################### EXTENSIONS BUILD IMAGE #################### # Build DeepEP - runs in PARALLEL with csrc-build # This stage is independent and doesn't affect csrc cache +# Kernel symbol map for CI test selection: which source file produced each GPU +# kernel symbol, read from the objects this stage just compiled. Only when +# VLLM_KERNEL_SYMBOL_MAP=1 (nightly), and placed after the compile layers so +# their cache is untouched. Never fails the build; an empty map is a signal. +ARG VLLM_KERNEL_SYMBOL_MAP="" +ARG buildkite_commit="" +COPY tools/ci/kernel_symbol_map.py tools/ci/kernel_symbol_map.py +RUN if [ "${VLLM_KERNEL_SYMBOL_MAP}" = "1" ]; then \ + mkdir -p dist && python3 tools/ci/kernel_symbol_map.py \ + --build-root build --source-root . --commit "${buildkite_commit}" \ + --out dist/kernel_symbol_map.json.gz; \ + else echo "Skipping kernel symbol map (VLLM_KERNEL_SYMBOL_MAP != 1)"; fi + +# Exported on its own by the CI bake target of the same name (ci-infra +# docker/ci.hcl) with output=type=local, so the map becomes a build artifact. +FROM scratch AS kernel-symbol-map +COPY --from=csrc-build /workspace/dist/kernel_symbol_map.json.gz /kernel_symbol_map.json.gz + FROM base AS extensions-build ARG TARGETARCH ARG CUDA_VERSION @@ -735,6 +753,8 @@ RUN --mount=type=cache,target=/opt/uv/cache \ FROM ${FINAL_BASE_IMAGE} AS vllm-runtime-base ARG INSTALL_RUBIN_PRERELEASE +ARG TARGETPLATFORM +ARG max_jobs=2 ARG CUDA_VERSION ARG PYTHON_VERSION ARG NCCL_VERSION @@ -854,6 +874,20 @@ RUN if [ "${INSTALL_RUBIN_PRERELEASE}" = "true" ]; then \ && ldconfig; \ fi +# Install the initialized-snapshot runtime on the architecture supported by +# `vllm snapshot`. CRIU and its CUDA plugin are built from the same pinned +# source. The `cuda-checkpoint`-compatible helper is built from NVIDIA Dynamo's +# Apache-2.0 source against the public CUDA Driver checkpoint APIs instead of +# redistributing NVIDIA's proprietary cuda-checkpoint binary. +COPY tools/install_snapshot_runtime.sh /usr/local/bin/install_snapshot_runtime.sh +RUN if [ "${TARGETPLATFORM}" = "linux/amd64" ]; then \ + MAX_JOBS="${max_jobs}" /usr/local/bin/install_snapshot_runtime.sh; \ + else \ + echo "Initialized snapshots are not installed for ${TARGETPLATFORM}"; \ + fi + +ENV CRIU_CUDA_PLUGIN_DIR=/usr/local/lib/criu + # Install uv for faster pip installs RUN python3 -m pip install uv @@ -1047,6 +1081,10 @@ RUN --mount=type=bind,from=build,src=/workspace/dist,target=/vllm-workspace/dist --extra-index-url ${PYTORCH_CUDA_INDEX_BASE_URL}/cu$(echo $CUDA_VERSION | cut -d. -f1,2 | tr -d '.'); \ fi +# Expose the Rust frontend bundled in the wheel on the container's PATH. +RUN ln -s "$(python3 -c 'import sysconfig; print(sysconfig.get_path("platlib"))')/vllm/vllm-rs" /usr/local/bin/vllm-rs \ + && vllm-rs --version + RUN --mount=type=cache,target=/opt/uv/cache \ . /etc/environment && \ uv pip list @@ -1226,7 +1264,8 @@ ENV UV_HTTP_TIMEOUT=500 ARG torch_cuda_arch_list='7.5 8.0 8.6 8.9 9.0 10.0 11.0 12.0' ENV TORCH_CUDA_ARCH_LIST=${torch_cuda_arch_list} RUN --mount=type=cache,target=/opt/uv/cache \ - --mount=type=bind,source=requirements/kv_connectors.txt,target=/tmp/kv_connectors.txt,ro \ + --mount=type=bind,source=requirements,target=/tmp/vllm-requirements,ro \ + set -eu; \ CUDA_MAJOR="${CUDA_VERSION%%.*}"; \ CUDA_VERSION_DASH=$(echo $CUDA_VERSION | cut -d. -f1,2 | tr '.' '-'); \ CUDA_HOME=/usr/local/cuda; \ @@ -1235,11 +1274,19 @@ RUN --mount=type=cache,target=/opt/uv/cache \ libcublas-dev-${CUDA_VERSION_DASH} \ libcusolver-dev-${CUDA_VERSION_DASH}"; \ if [ "$INSTALL_KV_CONNECTORS" = "true" ]; then \ - uv pip install --system -r /tmp/kv_connectors.txt --no-build || ( \ + KV_REQUIREMENTS=/tmp/vllm-requirements/kv_connectors.txt; \ + if [ "${CUDA_MAJOR}" = "12" ]; then \ + KV_REQUIREMENTS=/tmp/vllm-requirements/kv_connectors_cu12.txt; \ + fi; \ + # Preserve the stack used to compile vLLM when resolving optional dependencies. \ + uv pip freeze --system > /tmp/kv-installed.txt; \ + grep -E '^(torch|torchaudio|torchvision|triton|cuda-bindings|cuda-python|cuda-toolkit)==' \ + /tmp/kv-installed.txt > /tmp/kv-constraints.txt; \ + uv pip install --system -c /tmp/kv-constraints.txt -r "${KV_REQUIREMENTS}" --no-build || ( \ # if the above fails, install from source apt-get update -y && \ apt-get install -y --no-install-recommends --allow-change-held-packages ${BUILD_PKGS} && \ - uv pip install --system -r /tmp/kv_connectors.txt --no-build-isolation && \ + uv pip install --system -c /tmp/kv-constraints.txt -r "${KV_REQUIREMENTS}" --no-build-isolation && \ apt-get purge -y ${BUILD_PKGS} && \ # clean up -dev packages, keep runtime libraries rm -rf /var/lib/apt/lists/* \ @@ -1250,8 +1297,9 @@ RUN --mount=type=cache,target=/opt/uv/cache \ MOONCAKE_VERSION=$(python3 -c "import importlib.metadata as m; print(m.version('mooncake-transfer-engine'))" 2>/dev/null || true) && \ if [ "${CUDA_MAJOR}" = "13" ] && [ -n "${MOONCAKE_VERSION}" ]; then \ uv pip uninstall --system mooncake-transfer-engine 2>/dev/null || true; \ - uv pip install --system "mooncake-transfer-engine-cuda13==${MOONCAKE_VERSION}"; \ + uv pip install --system -c /tmp/kv-constraints.txt "mooncake-transfer-engine-cuda13==${MOONCAKE_VERSION}"; \ fi; \ + rm -f /tmp/kv-installed.txt /tmp/kv-constraints.txt; \ fi # Mooncake registers GPU memory for RDMA via nvidia-peermem by default. Hosts diff --git a/docker/Dockerfile.cpu b/docker/Dockerfile.cpu index f81ae88096..d61e186dac 100644 --- a/docker/Dockerfile.cpu +++ b/docker/Dockerfile.cpu @@ -140,6 +140,12 @@ RUN --mount=type=cache,target=/root/.cache/uv \ git clone --recurse-submodules --shallow-submodules --filter=blob:none "https://github.com/triton-lang/triton-cpu.git"; \ cd triton-cpu; \ git checkout "270e696d"; \ + if [ ! -f third_party/sleef/CMakeLists.txt ]; then \ + echo "sleef submodule empty after shallow fetch; retrying unshallow"; \ + git submodule foreach --recursive 'git fetch --unshallow origin || true'; \ + git submodule update --init --recursive --force; \ + fi; \ + test -f third_party/sleef/CMakeLists.txt || { echo "FATAL: sleef submodule unavailable"; exit 1; }; \ if [ "$USE_SCCACHE" = "1" ]; then \ if [ "$SCCACHE_LOCAL_ONLY" = "1" ]; then \ unset SCCACHE_BUCKET SCCACHE_REGION; \ @@ -465,6 +471,10 @@ ADD ./benchmarks/ ./benchmarks/ ADD ./vllm/collect_env.py . ADD ./docker/ ./docker/ ADD ./.buildkite/ ./.buildkite/ +# tools/ carries repo-tooling code that a few tests exercise directly +# (e.g. tests/tools/test_check_test_tethering.py imports +# tools.pre_commit.check_test_tethering). +ADD ./tools/ ./tools/ # install development dependencies (for testing) RUN --mount=type=cache,target=/root/.cache/uv \ diff --git a/docker/Dockerfile.rock b/docker/Dockerfile.rock index 6e52199fa3..16e01ad723 100644 --- a/docker/Dockerfile.rock +++ b/docker/Dockerfile.rock @@ -914,7 +914,7 @@ RUN case "$(which python3)" in \ *) ;; esac RUN --mount=type=cache,target=/root/.cache/uv \ - uv pip install --system --upgrade huggingface-hub[cli] + uv pip install --system --upgrade huggingface-hub # Install vLLM using uv (inherited from base stage) # Note: No -U flag to avoid upgrading PyTorch ROCm to CUDA version diff --git a/docker/Dockerfile.rocm b/docker/Dockerfile.rocm index 483aa9bb97..0bcc8a6ca6 100644 --- a/docker/Dockerfile.rocm +++ b/docker/Dockerfile.rocm @@ -938,7 +938,7 @@ RUN case "$(which python3)" in \ *) ;; esac RUN --mount=type=cache,target=/root/.cache/uv \ - uv pip install --system --upgrade huggingface-hub[cli] + uv pip install --system --upgrade huggingface-hub # Install vLLM using uv (inherited from base stage) # Note: No -U flag to avoid upgrading PyTorch ROCm to CUDA version diff --git a/docker/Dockerfile.rocm_gfx1250 b/docker/Dockerfile.rocm_gfx1250 index c68fdfebfe..ecb618d990 100644 --- a/docker/Dockerfile.rocm_gfx1250 +++ b/docker/Dockerfile.rocm_gfx1250 @@ -671,7 +671,7 @@ RUN case "$(which python3)" in \ *) ;; esac RUN --mount=type=cache,target=/root/.cache/uv \ - uv pip install --system --upgrade huggingface-hub[cli] + uv pip install --system --upgrade huggingface-hub # Install vLLM using uv (inherited from base stage) # Note: No -U flag to avoid upgrading PyTorch ROCm to CUDA version diff --git a/docs/assets/contributing/dockerfile-stages-dependency.png b/docs/assets/contributing/dockerfile-stages-dependency.png index 9eca3ae913..8ab9103801 100644 Binary files a/docs/assets/contributing/dockerfile-stages-dependency.png and b/docs/assets/contributing/dockerfile-stages-dependency.png differ diff --git a/docs/contributing/profiling.md b/docs/contributing/profiling.md index 880d065a2c..86bd68acef 100644 --- a/docs/contributing/profiling.md +++ b/docs/contributing/profiling.md @@ -22,6 +22,10 @@ To use the `torch.profiler` module, set the `profiler` entry to `'torch'` and `t - `torch_profiler_with_flops` to enable recording FLOPs, off by default - `torch_profiler_use_gzip` to control gzip-compressing profiling files, on by default - `torch_profiler_dump_cuda_time_total` to control dumping and printing the aggregated CUDA self time table, on by default +- `torch_profiler_activities` to select worker activities. Defaults are + platform-specific: `["CPU"]` on CPU, `["CPU", "CUDA"]` on NVIDIA GPUs, and + `["CPU", "XPU"]` on Intel GPUs. Selecting `["CUDA"]` omits CPU annotations + and the AsyncLLM CPU trace to reduce profiling overhead and trace size. When using `vllm bench serve`, you can enable profiling by passing the `--profile` flag. diff --git a/docs/deployment/integrations/kthena.md b/docs/deployment/integrations/kthena.md index 0285b733c6..73dcb3dafd 100644 --- a/docs/deployment/integrations/kthena.md +++ b/docs/deployment/integrations/kthena.md @@ -2,7 +2,9 @@ [**Kthena**](https://github.com/volcano-sh/kthena) is a Kubernetes-native LLM inference platform that transforms how organizations deploy and manage Large Language Models in production. Built with declarative model lifecycle management and intelligent request routing, it provides high performance and enterprise-grade scalability for LLM inference workloads. -This guide shows how to deploy a production-grade, **multi-node vLLM** service on Kubernetes. +Kthena manages inference workloads through `ModelServing` and provides model-aware request routing through `ModelServer` and `ModelRoute`. + +This guide shows how to deploy a **multi-node vLLM** service on Kubernetes using `ModelServing` and expose it through a Kubernetes Service. We’ll: @@ -18,7 +20,7 @@ You need: - A Kubernetes cluster with **GPU nodes**. - `kubectl` access with cluster-admin or equivalent permissions. -- **Volcano** installed for gang scheduling. +- **Volcano 1.14 or later** for the role-level gang scheduling used in this example. - **Kthena** installed with the `ModelServing` CRD available. - A valid **Hugging Face token** if loading models from Hugging Face Hub. @@ -30,21 +32,22 @@ helm repo update helm install volcano volcano-sh/volcano -n volcano-system --create-namespace ``` -This provides the gang-scheduling and network topology features used by Kthena. +Volcano provides the gang scheduling used in this example. It is optional for Kthena deployments that do not use Volcano scheduling. ### 1.2 Install Kthena ```bash -helm install kthena oci://ghcr.io/volcano-sh/charts/kthena --version v0.1.0 --namespace kthena-system --create-namespace +helm install kthena oci://ghcr.io/volcano-sh/charts/kthena --version v1.0.0 --namespace kthena-system --create-namespace ``` - The `kthena-system` namespace is created. -- Kthena controllers and CRDs, including `ModelServing`, are installed and healthy. +- Kthena controllers and CRDs, including `ModelServing`, are installed. Validate: ```bash -kubectl get crd | grep modelserving +kubectl get pods -n kthena-system +kubectl get crd modelservings.workload.serving.volcano.sh ``` You should see: @@ -57,7 +60,7 @@ modelservings.workload.serving.volcano.sh ... ## 2. The Multi-Node vLLM `ModelServing` Example -Kthena provides an example manifest to deploy a **multi-node vLLM cluster running Llama**. Conceptually this is equivalent to the vLLM production stack Helm deployment, but expressed with `ModelServing`. +Kthena provides an example manifest to deploy a **multi-node vLLM cluster running Llama**. Each role replica contains an entry pod and its worker pods. A simplified version of the example (`llama-multinode`) looks like: @@ -116,10 +119,9 @@ Commands: - -c - > bash /vllm-workspace/examples/ray_serving/multi-node-serving.sh - leader --ray_cluster_size=2; python3 -m - vllm.entrypoints.launchers.api_server.entry --port 8080 --model - meta-llama/Llama-3.1-405B-Instruct --tensor-parallel-size 8 - --pipeline-parallel-size 2 + leader --ray_cluster_size=2; + vllm serve meta-llama/Llama-3.1-405B-Instruct --port 8080 --tensor-parallel-size 8 + --pipeline-parallel-size 2 --distributed-executor-backend ray ``` Worker: @@ -166,9 +168,9 @@ Save one of the following manifests to `modelserving.yaml`: restartGracePeriodSeconds: 60 gangPolicy: minRoleReplicas: - 405b: 1 + llama-405b: 1 roles: - - name: 405b + - name: llama-405b replicas: 2 entryTemplate: spec: @@ -256,9 +258,9 @@ Save one of the following manifests to `modelserving.yaml`: restartGracePeriodSeconds: 60 gangPolicy: minRoleReplicas: - 405b: 1 + llama-405b: 1 roles: - - name: 405b + - name: llama-405b replicas: 2 entryTemplate: spec: @@ -275,7 +277,7 @@ Save one of the following manifests to `modelserving.yaml`: - sh - -c - "bash /vllm-workspace/examples/ray_serving/multi-node-serving.sh leader --ray_cluster_size=2; - vllm serve meta-llama/Llama-3.1-405B-Instruct --port 8080 --tensor-parallel-size 8 --pipeline-parallel-size 2" + vllm serve meta-llama/Llama-3.1-405B-Instruct --port 8080 --tensor-parallel-size 8 --pipeline-parallel-size 2 --distributed-executor-backend ray" resources: limits: nvidia.com/gpu: "8" @@ -383,24 +385,24 @@ kubectl get pod -owide -l modelserving.volcano.sh/name=llama-multinode Example output (from docs): ```text -NAMESPACE NAME READY STATUS RESTARTS AGE IP NODE ... -default llama-multinode-0-405b-0-0 1/1 Running 0 15m 10.244.0.56 192.168.5.12 ... -default llama-multinode-0-405b-0-1 1/1 Running 0 15m 10.244.0.58 192.168.5.43 ... -default llama-multinode-0-405b-1-0 1/1 Running 0 15m 10.244.0.57 192.168.5.58 ... -default llama-multinode-0-405b-1-1 1/1 Running 0 15m 10.244.0.53 192.168.5.36 ... +NAMESPACE NAME READY STATUS RESTARTS AGE IP NODE ... +default llama-multinode-0-llama-405b-0-0 1/1 Running 0 15m 10.244.0.56 192.168.5.12 ... +default llama-multinode-0-llama-405b-0-1 1/1 Running 0 15m 10.244.0.58 192.168.5.43 ... +default llama-multinode-0-llama-405b-1-0 1/1 Running 0 15m 10.244.0.57 192.168.5.58 ... +default llama-multinode-0-llama-405b-1-1 1/1 Running 0 15m 10.244.0.53 192.168.5.36 ... ``` Pod name pattern: - `llama-multinode----`. -The first number indicates `ServingGroup`. The second (`405b`) is the `Role`. The remaining indices identify the pod within the role. +The first index identifies the `ServingGroup`, followed by the role name (`llama-405b`), role replica index, and pod index. --- -## 6. Accessing the vLLM OpenAI-Compatible API +## 5. Accessing the vLLM OpenAI-Compatible API -Expose the entry via a Service: +Save the following Service as `service.yaml` to expose the entry pods: ```yaml apiVersion: v1 @@ -412,7 +414,6 @@ spec: selector: modelserving.volcano.sh/name: llama-multinode modelserving.volcano.sh/entry: "true" - # optionally further narrow to leader role if you label it ports: - name: http port: 80 @@ -420,9 +421,10 @@ spec: type: ClusterIP ``` -Port-forward from your local machine: +Apply the Service, then port-forward from your local machine: ```bash +kubectl apply -f service.yaml kubectl port-forward svc/llama-multinode-openai 30080:80 -n default ``` @@ -450,11 +452,12 @@ You should see an OpenAI-style response from vLLM. --- -## 7. Clean Up +## 6. Clean Up To remove the deployment and its resources: ```bash +kubectl delete -f service.yaml kubectl delete modelserving llama-multinode -n default ``` @@ -464,3 +467,5 @@ If you’re done with the entire stack: helm uninstall kthena -n kthena-system # or your Kthena release name helm uninstall volcano -n volcano-system ``` + +For model-aware routing and prefill-decode disaggregation, see the [Kthena documentation](https://kthena.volcano.sh/docs/user-guide/router-routing). diff --git a/docs/features/automatic_prefix_caching.md b/docs/features/automatic_prefix_caching.md index 03cca3b38d..cfc8c19acb 100644 --- a/docs/features/automatic_prefix_caching.md +++ b/docs/features/automatic_prefix_caching.md @@ -15,7 +15,7 @@ Set `enable_prefix_caching=True` in vLLM engine to enable APC. Here is an exampl ## Hybrid Mamba models -Under `--mamba-cache-mode align`, Mamba state is stored only on the Mamba block grid, so a prefix-cache hit can resume only at a block boundary. `--enable-mamba-fine-grained-prefix-cache` also stores a checkpoint at the shared-prefix junction, the point where an earlier request with the same prefix stopped. Requests whose shared prefix ends inside a block can then reuse it. +Under `--mamba-cache-mode align`, Mamba state is stored only on the Mamba block grid, so a prefix-cache hit can resume only at a block boundary. `--enable-mamba-shared-prefix-checkpoint` also stores a checkpoint at the shared-prefix junction, the point where an earlier request with the same prefix stopped. Requests whose shared prefix ends inside a block can then reuse it. This helps when many requests share a long system prompt and then diverge. It is off by default, and takes effect only when all of the following hold: @@ -28,7 +28,7 @@ This helps when many requests share a long system prompt and then diverge. It is vllm serve \ --mamba-cache-mode align \ --prefix-match-unit 64 \ - --enable-mamba-fine-grained-prefix-cache + --enable-mamba-shared-prefix-checkpoint ``` `--prefix-match-unit` is required. It sets the granularity at which prefix-cache keys are computed. When unset it defaults to the greatest common divisor of the prefix-cacheable KV cache group block sizes. Under `align` that is the block size itself, so no sub-block boundary exists and the flag has no effect. diff --git a/docs/features/disagg_encoder.md b/docs/features/disagg_encoder.md index a4b2ece241..383e7d45ff 100644 --- a/docs/features/disagg_encoder.md +++ b/docs/features/disagg_encoder.md @@ -32,8 +32,9 @@ Design doc: