forked from Karylab-cklius/vllm
Compare commits
25
Commits
| Author | SHA1 | Date | |
|---|---|---|---|
|
|
3a981fbbdb | ||
|
|
7fc97042c3 | ||
|
|
18c4067a54 | ||
|
|
550218b136 | ||
|
|
5c342876a6 | ||
|
|
9427c45386 | ||
|
|
43c8cbf79b | ||
|
|
62286308c9 | ||
|
|
26587f9519 | ||
|
|
93e3bc8f30 | ||
|
|
c2c9f7c5e2 | ||
|
|
1be6e937b2 | ||
|
|
b3cfca996c | ||
|
|
487dfb3418 | ||
|
|
107a03ba63 | ||
|
|
56a357ed33 | ||
|
|
bea70c7cfc | ||
|
|
75fe92a316 | ||
|
|
b7b58d1eba | ||
|
|
36484e464a | ||
|
|
9e57de7197 | ||
|
|
8c5dafcd09 | ||
|
|
05fa8183a6 | ||
|
|
d973cce3ca | ||
|
|
775c1589ea |
@@ -18,7 +18,7 @@ steps:
|
||||
- label: "XPU example Test"
|
||||
depends_on:
|
||||
- image-build-xpu
|
||||
timeout_in_minutes: 30
|
||||
timeout_in_minutes: 50
|
||||
optional: true
|
||||
device: intel_gpu
|
||||
agent_tags:
|
||||
@@ -39,7 +39,7 @@ steps:
|
||||
- label: "XPU V1 test"
|
||||
depends_on:
|
||||
- image-build-xpu
|
||||
timeout_in_minutes: 30
|
||||
timeout_in_minutes: 70
|
||||
optional: true
|
||||
device: intel_gpu
|
||||
agent_tags:
|
||||
@@ -60,7 +60,7 @@ steps:
|
||||
- label: "XPU server test"
|
||||
depends_on:
|
||||
- image-build-xpu
|
||||
timeout_in_minutes: 30
|
||||
timeout_in_minutes: 45
|
||||
optional: true
|
||||
device: intel_gpu
|
||||
agent_tags:
|
||||
|
||||
@@ -3,7 +3,7 @@ depends_on:
|
||||
- image-build-xpu
|
||||
steps:
|
||||
- label: XPU Sleep Mode
|
||||
timeout_in_minutes: 30
|
||||
timeout_in_minutes: 45
|
||||
device: intel_gpu
|
||||
agent_tags:
|
||||
label: production
|
||||
|
||||
@@ -86,7 +86,7 @@ steps:
|
||||
pytest -v -s lora/test_punica_ops.py::test_add_lora_fused_moe_early_exit'
|
||||
|
||||
- label: LoRA Punica FP8/XPU Ops
|
||||
timeout_in_minutes: 45
|
||||
timeout_in_minutes: 60
|
||||
device: intel_gpu
|
||||
agent_tags:
|
||||
label: production
|
||||
|
||||
@@ -3,7 +3,7 @@ depends_on:
|
||||
- image-build-xpu
|
||||
steps:
|
||||
- label: V1 Core + KV + Metrics
|
||||
timeout_in_minutes: 30
|
||||
timeout_in_minutes: 45
|
||||
device: intel_gpu
|
||||
agent_tags:
|
||||
label: production
|
||||
@@ -33,7 +33,7 @@ steps:
|
||||
pytest -v -s v1/executor'
|
||||
|
||||
- label: V1 Sample + Logits
|
||||
timeout_in_minutes: 30
|
||||
timeout_in_minutes: 90
|
||||
device: intel_gpu
|
||||
agent_tags:
|
||||
label: production
|
||||
@@ -152,7 +152,7 @@ steps:
|
||||
|
||||
- label: Regression
|
||||
key: regression
|
||||
timeout_in_minutes: 30
|
||||
timeout_in_minutes: 50
|
||||
device: intel_gpu
|
||||
agent_tags:
|
||||
label: production
|
||||
@@ -186,7 +186,7 @@ steps:
|
||||
|
||||
- label: Metrics, Tracing (2 GPUs)
|
||||
key: metrics-tracing-2-gpus
|
||||
timeout_in_minutes: 30
|
||||
timeout_in_minutes: 45
|
||||
num_devices: 2
|
||||
device: intel_gpu
|
||||
agent_tags:
|
||||
@@ -222,7 +222,7 @@ steps:
|
||||
|
||||
- label: Async Engine, Inputs, Utils, Worker
|
||||
key: async-engine-inputs-utils-worker
|
||||
timeout_in_minutes: 30
|
||||
timeout_in_minutes: 55
|
||||
device: intel_gpu
|
||||
agent_tags:
|
||||
label: production
|
||||
|
||||
@@ -4,7 +4,7 @@ depends_on:
|
||||
steps:
|
||||
- label: Distributed Model Tests (2 GPUs)
|
||||
key: distributed-model-tests-2-gpus
|
||||
timeout_in_minutes: 50
|
||||
timeout_in_minutes: 65
|
||||
device: intel_gpu
|
||||
agent_tags:
|
||||
label: production
|
||||
|
||||
@@ -4,7 +4,7 @@ depends_on:
|
||||
steps:
|
||||
- label: "Multi-Modal Models (Standard) 1: qwen2"
|
||||
key: multi-modal-models-standard-1-qwen2
|
||||
timeout_in_minutes: 45
|
||||
timeout_in_minutes: 70
|
||||
device: intel_gpu
|
||||
agent_tags:
|
||||
label: production
|
||||
@@ -29,7 +29,7 @@ steps:
|
||||
|
||||
- label: "Multi-Modal Models (Standard) 2: qwen3 + gemma"
|
||||
key: multi-modal-models-standard-2-qwen3-gemma
|
||||
timeout_in_minutes: 45
|
||||
timeout_in_minutes: 70
|
||||
device: intel_gpu
|
||||
agent_tags:
|
||||
label: production
|
||||
@@ -52,7 +52,7 @@ steps:
|
||||
|
||||
- label: "Multi-Modal Models (Standard) 3: llava + qwen2_vl"
|
||||
key: multi-modal-models-standard-3-llava-qwen2-vl
|
||||
timeout_in_minutes: 45
|
||||
timeout_in_minutes: 65
|
||||
device: intel_gpu
|
||||
agent_tags:
|
||||
label: production
|
||||
@@ -100,7 +100,7 @@ steps:
|
||||
|
||||
- label: Multi-Modal Processor # 44min
|
||||
key: multi-modal-processor
|
||||
timeout_in_minutes: 45
|
||||
timeout_in_minutes: 60
|
||||
device: intel_gpu
|
||||
agent_tags:
|
||||
label: production
|
||||
|
||||
@@ -17,7 +17,7 @@ steps:
|
||||
- label: "XPU example Test"
|
||||
depends_on:
|
||||
- image-build-xpu
|
||||
timeout_in_minutes: 30
|
||||
timeout_in_minutes: 50
|
||||
device: intel_gpu
|
||||
agent_tags:
|
||||
label: production
|
||||
@@ -76,7 +76,7 @@ steps:
|
||||
- label: "XPU V1 test"
|
||||
depends_on:
|
||||
- image-build-xpu
|
||||
timeout_in_minutes: 30
|
||||
timeout_in_minutes: 70
|
||||
device: intel_gpu
|
||||
agent_tags:
|
||||
label: production
|
||||
@@ -104,7 +104,7 @@ steps:
|
||||
- label: "XPU server test"
|
||||
depends_on:
|
||||
- image-build-xpu
|
||||
timeout_in_minutes: 30
|
||||
timeout_in_minutes: 45
|
||||
device: intel_gpu
|
||||
agent_tags:
|
||||
label: production
|
||||
|
||||
@@ -15,9 +15,9 @@ set -euo pipefail
|
||||
|
||||
DEFAULT_REPO_SLUG="vllm-project/vllm"
|
||||
DEFAULT_CI_HCL_SOURCE="docker/ci-rocm.hcl"
|
||||
DEFAULT_CI_BASE_CONTENT_FILES="requirements/common.txt requirements/rocm.txt requirements/test/rocm.txt docker/Dockerfile.rocm_base docker/ci-rocm.hcl docker/docker-bake-rocm.hcl tools/install_torchcodec_rocm.sh tests/vllm_test_utils .buildkite/scripts/ci-bake-rocm.sh .buildkite/scripts/rocm/build-ci-base.sh"
|
||||
DEFAULT_CI_BASE_CONTENT_FILES="requirements/common.txt requirements/rocm.txt requirements/test/rocm.txt docker/Dockerfile.rocm_base docker/ci-rocm.hcl docker/docker-bake-rocm.hcl tools/install_torchcodec_rocm.sh tools/install_protoc.sh rust-toolchain.toml tests/vllm_test_utils .buildkite/scripts/ci-bake-rocm.sh .buildkite/scripts/rocm/build-ci-base.sh"
|
||||
DEFAULT_CI_BASE_DOCKERFILE="docker/Dockerfile.rocm"
|
||||
DEFAULT_CI_BASE_DOCKERFILE_STAGES="base build_rixl build_rocshmem build_deepep mori_base ci_base"
|
||||
DEFAULT_CI_BASE_DOCKERFILE_STAGES="base rust_toolchain_input_0 rust_toolchain_input_1 rust-toolchain-input rust-toolchain build_rixl build_rocshmem build_deepep mori_base ci_base"
|
||||
DEFAULT_CI_BASE_METADATA_VERSION="1"
|
||||
IMAGE_EXISTED_BEFORE_BUILD=0
|
||||
|
||||
@@ -764,7 +764,7 @@ configure_ci_base_image_refs() {
|
||||
fi
|
||||
set_buildkite_metadata "rocm-ci-base-image" "${CI_BASE_IMAGE_TAG}"
|
||||
set_buildkite_metadata "rocm-ci-base-image-content" "${content_tag}"
|
||||
set_buildkite_metadata "rocm-ci-base-image-commit" "${CI_BASE_IMAGE_TAG_COMMIT:-}"
|
||||
set_buildkite_metadata "rocm-ci-base-image-commit" "${CI_BASE_IMAGE_TAG_COMMIT_REF:-}"
|
||||
set_buildkite_metadata "rocm-ci-base-image-stable" "${CI_BASE_IMAGE_TAG_STABLE:-}"
|
||||
return 0
|
||||
fi
|
||||
@@ -1211,12 +1211,24 @@ uses_rocm_csrc_cache() {
|
||||
esac
|
||||
}
|
||||
|
||||
uses_rocm_rust_cache() {
|
||||
case "${TARGET}" in
|
||||
rust-rocm-ci|test-rocm-ci|test-rocm-ci-with-wheel|test-rocm-ci-with-artifacts|export-wheel-rocm)
|
||||
return 0
|
||||
;;
|
||||
*)
|
||||
return 1
|
||||
;;
|
||||
esac
|
||||
}
|
||||
|
||||
compute_rocm_csrc_content_hash() {
|
||||
local bake_dir=""
|
||||
local dockerfile_rocm=""
|
||||
local -a content_paths=(
|
||||
"requirements/common.txt"
|
||||
"requirements/rocm.txt"
|
||||
"pyproject.toml"
|
||||
"setup.py"
|
||||
"CMakeLists.txt"
|
||||
"cmake"
|
||||
@@ -1260,6 +1272,56 @@ compute_rocm_csrc_content_hash_if_needed() {
|
||||
echo "ROCm csrc content cache ref: ${ROCM_CSRC_CONTENT_CACHE_REF}"
|
||||
}
|
||||
|
||||
compute_rocm_rust_content_hash() {
|
||||
local bake_dir=""
|
||||
local dockerfile_rocm=""
|
||||
local -a content_paths=(
|
||||
"requirements/build/rust.txt"
|
||||
"rust/Cargo.lock"
|
||||
"rust/Cargo.toml"
|
||||
"rust/proto"
|
||||
"rust/src"
|
||||
"rust-toolchain.toml"
|
||||
"tools/build_rust.py"
|
||||
"tools/install_protoc.sh"
|
||||
"build_rust.sh"
|
||||
)
|
||||
local -a content_args=()
|
||||
|
||||
bake_dir=$(dirname "${VLLM_BAKE_FILE}")
|
||||
dockerfile_rocm="${bake_dir}/Dockerfile.rocm"
|
||||
mapfile -t content_args < <(
|
||||
get_content_arg_names "${dockerfile_rocm}" "base rust_toolchain_input_0 rust_toolchain_input_1 rust-toolchain-input rust_input_0 rust_input_1 rust-input rust-toolchain rust-build" "${ROCM_RUST_CONTENT_ARGS:-}"
|
||||
)
|
||||
|
||||
{
|
||||
printf 'rust-input-files-hash:%s\n' "$(compute_content_hash "${content_paths[@]}")"
|
||||
printf 'dockerfile:%s\n' "${dockerfile_rocm}"
|
||||
printf 'resolved-build-args:\n'
|
||||
hash_dockerfile_arg_values "${dockerfile_rocm}" "${content_args[@]}"
|
||||
printf 'dockerfile-stages:base rust_toolchain_input_0 rust_toolchain_input_1 rust-toolchain-input rust_input_0 rust_input_1 rust-input rust-toolchain rust-build\n'
|
||||
if [[ -f "${dockerfile_rocm}" ]]; then
|
||||
hash_dockerfile_stages "${dockerfile_rocm}" "base rust_toolchain_input_0 rust_toolchain_input_1 rust-toolchain-input rust_input_0 rust_input_1 rust-input rust-toolchain rust-build"
|
||||
else
|
||||
printf 'missing:%s\n' "${dockerfile_rocm}"
|
||||
fi
|
||||
} | sha256sum | cut -d' ' -f1
|
||||
}
|
||||
|
||||
compute_rocm_rust_content_hash_if_needed() {
|
||||
local cache_repo="${DOCKERHUB_CACHE_REPO:-rocm/vllm-ci-cache}"
|
||||
|
||||
if [[ "${ROCM_RUST_CONTENT_CACHE:-1}" == "0" ]] || ! uses_rocm_rust_cache; then
|
||||
return 0
|
||||
fi
|
||||
|
||||
ROCM_RUST_CONTENT_HASH=$(compute_rocm_rust_content_hash)
|
||||
ROCM_RUST_CONTENT_CACHE_REF="${cache_repo}:rust-rocm-input-${ROCM_RUST_CONTENT_HASH}"
|
||||
export ROCM_RUST_CONTENT_HASH
|
||||
export ROCM_RUST_CONTENT_CACHE_REF
|
||||
echo "ROCm Rust content cache ref: ${ROCM_RUST_CONTENT_CACHE_REF}"
|
||||
}
|
||||
|
||||
write_hcl_string_list_entries() {
|
||||
local indent="$1"
|
||||
local value=""
|
||||
@@ -1317,6 +1379,7 @@ write_rocm_build_arg_override() {
|
||||
"${CI_BASE_DOCKERFILE_STAGES:-${DEFAULT_CI_BASE_DOCKERFILE_STAGES}}" \
|
||||
"${CI_BASE_CONTENT_ARGS:-}"
|
||||
get_content_arg_names "${dockerfile_rocm}" "base csrc-build" "${ROCM_CSRC_CONTENT_ARGS:-}"
|
||||
get_content_arg_names "${dockerfile_rocm}" "base rust_toolchain_input_0 rust_toolchain_input_1 rust-toolchain-input rust_input_0 rust_input_1 rust-input rust-toolchain rust-build" "${ROCM_RUST_CONTENT_ARGS:-}"
|
||||
} | awk 'NF && !seen[$0]++'
|
||||
)
|
||||
|
||||
@@ -1365,46 +1428,133 @@ validate_cache_export_mode() {
|
||||
esac
|
||||
}
|
||||
|
||||
validate_content_cache_export_mode() {
|
||||
local mode="$1"
|
||||
local env_name="$2"
|
||||
|
||||
case "${mode}" in
|
||||
missing|always|never)
|
||||
;;
|
||||
*)
|
||||
echo "Error: ${env_name} must be one of: missing, always, never"
|
||||
exit 1
|
||||
;;
|
||||
esac
|
||||
}
|
||||
|
||||
should_export_content_cache_ref() {
|
||||
local cache_ref="$1"
|
||||
local cache_name="$2"
|
||||
local mode="${ROCM_CONTENT_CACHE_EXPORT_MODE:-missing}"
|
||||
|
||||
case "${mode}" in
|
||||
always)
|
||||
echo "${cache_name} content cache export mode is always; exporting ${cache_ref}"
|
||||
return 0
|
||||
;;
|
||||
never)
|
||||
echo "${cache_name} content cache export mode is never; not exporting ${cache_ref}"
|
||||
return 1
|
||||
;;
|
||||
missing|"")
|
||||
if docker buildx imagetools inspect "${cache_ref}" >/dev/null 2>&1; then
|
||||
echo "${cache_name} content cache exists; not re-exporting ${cache_ref}"
|
||||
return 1
|
||||
fi
|
||||
echo "${cache_name} content cache missing; will export ${cache_ref}"
|
||||
return 0
|
||||
;;
|
||||
*)
|
||||
echo "Error: ROCM_CONTENT_CACHE_EXPORT_MODE must be one of: missing, always, never"
|
||||
exit 1
|
||||
;;
|
||||
esac
|
||||
}
|
||||
|
||||
write_rocm_cache_override() {
|
||||
local cache_repo="${DOCKERHUB_CACHE_REPO:-rocm/vllm-ci-cache}"
|
||||
local content_cache_export_mode="${ROCM_CONTENT_CACHE_EXPORT_MODE:-missing}"
|
||||
local csrc_cache_to_mode="${ROCM_CSRC_CACHE_TO_MODE:-max}"
|
||||
local rust_cache_to_mode="${ROCM_RUST_CACHE_TO_MODE:-max}"
|
||||
local rocm_cache_to_mode="${ROCM_FINAL_CACHE_TO_MODE:-min}"
|
||||
local -a content_cache_from=()
|
||||
local -a csrc_content_cache_from=()
|
||||
local -a rust_content_cache_from=()
|
||||
local -a combined_content_cache_from=()
|
||||
local -a csrc_cache_to=()
|
||||
local -a rust_cache_to=()
|
||||
local -a rocm_cache_to=()
|
||||
local -a export_wheel_cache_to=()
|
||||
local export_csrc_cache=1
|
||||
local export_rust_cache=1
|
||||
|
||||
if ! uses_rocm_csrc_cache; then
|
||||
if ! uses_rocm_csrc_cache && ! uses_rocm_rust_cache; then
|
||||
return 0
|
||||
fi
|
||||
|
||||
validate_content_cache_export_mode \
|
||||
"${content_cache_export_mode}" \
|
||||
"ROCM_CONTENT_CACHE_EXPORT_MODE"
|
||||
validate_cache_export_mode "${csrc_cache_to_mode}" "ROCM_CSRC_CACHE_TO_MODE"
|
||||
validate_cache_export_mode "${rust_cache_to_mode}" "ROCM_RUST_CACHE_TO_MODE"
|
||||
validate_cache_export_mode "${rocm_cache_to_mode}" "ROCM_FINAL_CACHE_TO_MODE"
|
||||
echo "ROCm content cache export mode: ${content_cache_export_mode}"
|
||||
echo "ROCm csrc cache export mode: ${csrc_cache_to_mode}"
|
||||
echo "ROCm Rust cache export mode: ${rust_cache_to_mode}"
|
||||
echo "ROCm final image cache export mode: ${rocm_cache_to_mode}"
|
||||
|
||||
if [[ -n "${ROCM_CSRC_CONTENT_CACHE_REF:-}" ]]; then
|
||||
content_cache_from+=("type=registry,ref=${ROCM_CSRC_CONTENT_CACHE_REF}")
|
||||
csrc_cache_to+=(
|
||||
"type=registry,ref=${ROCM_CSRC_CONTENT_CACHE_REF},mode=${csrc_cache_to_mode},ignore-error=true"
|
||||
)
|
||||
csrc_content_cache_from+=("type=registry,ref=${ROCM_CSRC_CONTENT_CACHE_REF}")
|
||||
if should_export_content_cache_ref "${ROCM_CSRC_CONTENT_CACHE_REF}" "ROCm csrc"; then
|
||||
csrc_cache_to+=(
|
||||
"type=registry,ref=${ROCM_CSRC_CONTENT_CACHE_REF},mode=${csrc_cache_to_mode},ignore-error=true"
|
||||
)
|
||||
else
|
||||
export_csrc_cache=0
|
||||
fi
|
||||
fi
|
||||
|
||||
if [[ -n "${ROCM_RUST_CONTENT_CACHE_REF:-}" ]]; then
|
||||
rust_content_cache_from+=("type=registry,ref=${ROCM_RUST_CONTENT_CACHE_REF}")
|
||||
if should_export_content_cache_ref "${ROCM_RUST_CONTENT_CACHE_REF}" "ROCm Rust"; then
|
||||
rust_cache_to+=(
|
||||
"type=registry,ref=${ROCM_RUST_CONTENT_CACHE_REF},mode=${rust_cache_to_mode},ignore-error=true"
|
||||
)
|
||||
else
|
||||
export_rust_cache=0
|
||||
fi
|
||||
fi
|
||||
|
||||
combined_content_cache_from=("${csrc_content_cache_from[@]}" "${rust_content_cache_from[@]}")
|
||||
|
||||
# Docker Hub cache exports are best-effort. A cache-only target failure can
|
||||
# otherwise cancel the sibling image target before its manifest is pushed.
|
||||
if [[ -n "${BUILDKITE_COMMIT:-}" ]]; then
|
||||
csrc_cache_to+=(
|
||||
"type=registry,ref=${cache_repo}:csrc-rocm-${BUILDKITE_COMMIT},mode=${csrc_cache_to_mode},ignore-error=true"
|
||||
)
|
||||
if [[ ${export_csrc_cache} -eq 1 ]]; then
|
||||
csrc_cache_to+=(
|
||||
"type=registry,ref=${cache_repo}:csrc-rocm-${BUILDKITE_COMMIT},mode=${csrc_cache_to_mode},ignore-error=true"
|
||||
)
|
||||
fi
|
||||
if [[ ${export_rust_cache} -eq 1 ]]; then
|
||||
rust_cache_to+=(
|
||||
"type=registry,ref=${cache_repo}:rust-rocm-${BUILDKITE_COMMIT},mode=${rust_cache_to_mode},ignore-error=true"
|
||||
)
|
||||
fi
|
||||
rocm_cache_to+=(
|
||||
"type=registry,ref=${cache_repo}:rocm-${BUILDKITE_COMMIT},mode=${rocm_cache_to_mode},ignore-error=true"
|
||||
)
|
||||
fi
|
||||
|
||||
if [[ -n "${ROCM_CACHE_BRANCH_TAG:-}" ]]; then
|
||||
csrc_cache_to+=(
|
||||
"type=registry,ref=${cache_repo}:csrc-rocm-branch-${ROCM_CACHE_BRANCH_TAG},mode=${csrc_cache_to_mode},ignore-error=true"
|
||||
)
|
||||
if [[ ${export_csrc_cache} -eq 1 ]]; then
|
||||
csrc_cache_to+=(
|
||||
"type=registry,ref=${cache_repo}:csrc-rocm-branch-${ROCM_CACHE_BRANCH_TAG},mode=${csrc_cache_to_mode},ignore-error=true"
|
||||
)
|
||||
fi
|
||||
if [[ ${export_rust_cache} -eq 1 ]]; then
|
||||
rust_cache_to+=(
|
||||
"type=registry,ref=${cache_repo}:rust-rocm-branch-${ROCM_CACHE_BRANCH_TAG},mode=${rust_cache_to_mode},ignore-error=true"
|
||||
)
|
||||
fi
|
||||
rocm_cache_to+=(
|
||||
"type=registry,ref=${cache_repo}:rocm-branch-${ROCM_CACHE_BRANCH_TAG},mode=${rocm_cache_to_mode},ignore-error=true"
|
||||
)
|
||||
@@ -1422,7 +1572,7 @@ target "csrc-rocm-ci" {
|
||||
cache-from = concat(
|
||||
get_cache_from_rocm_csrc(),
|
||||
EOF
|
||||
write_hcl_string_list " " "${content_cache_from[@]}"
|
||||
write_hcl_string_list " " "${csrc_content_cache_from[@]}"
|
||||
cat <<EOF
|
||||
)
|
||||
EOF
|
||||
@@ -1430,11 +1580,23 @@ EOF
|
||||
cat <<EOF
|
||||
}
|
||||
|
||||
target "rust-rocm-ci" {
|
||||
cache-from = concat(
|
||||
get_cache_from_rocm_rust(),
|
||||
EOF
|
||||
write_hcl_string_list " " "${rust_content_cache_from[@]}"
|
||||
cat <<EOF
|
||||
)
|
||||
EOF
|
||||
write_hcl_string_list_attr " " "cache-to" "${rust_cache_to[@]}"
|
||||
cat <<EOF
|
||||
}
|
||||
|
||||
target "test-rocm-ci" {
|
||||
cache-from = concat(
|
||||
get_cache_from_rocm(),
|
||||
EOF
|
||||
write_hcl_string_list " " "${content_cache_from[@]}"
|
||||
write_hcl_string_list " " "${combined_content_cache_from[@]}"
|
||||
cat <<EOF
|
||||
)
|
||||
EOF
|
||||
@@ -1446,7 +1608,7 @@ target "export-wheel-rocm" {
|
||||
cache-from = concat(
|
||||
get_cache_from_rocm(),
|
||||
EOF
|
||||
write_hcl_string_list " " "${content_cache_from[@]}"
|
||||
write_hcl_string_list " " "${combined_content_cache_from[@]}"
|
||||
cat <<EOF
|
||||
)
|
||||
EOF
|
||||
@@ -1920,6 +2082,7 @@ main() {
|
||||
compute_dependency_cache_keys
|
||||
write_ci_base_label_override
|
||||
compute_rocm_csrc_content_hash_if_needed
|
||||
compute_rocm_rust_content_hash_if_needed
|
||||
write_rocm_cache_override
|
||||
resolve_ci_base_dependency_targets
|
||||
print_bake_config
|
||||
|
||||
@@ -210,7 +210,7 @@ repos:
|
||||
name: Check SPDX headers
|
||||
entry: python tools/pre_commit/check_spdx_header.py
|
||||
language: python
|
||||
types: [python]
|
||||
types_or: [python, rust, proto]
|
||||
- id: check-root-lazy-imports
|
||||
name: Check root lazy imports
|
||||
entry: python tools/pre_commit/check_init_lazy_imports.py
|
||||
|
||||
@@ -1116,6 +1116,164 @@ void fused_sigmoid_gating_delta_rule_update_kernel_impl(
|
||||
});
|
||||
}
|
||||
|
||||
// Speculative-decode variant: processes a varlen batch where each sequence has
|
||||
// ``q_len`` draft tokens, runs the recurrence sequentially over those tokens
|
||||
// (inside the kernel, so one dispatch handles the whole draft block), reads the
|
||||
// initial state from cache slot ``num_accepted-1`` and stores the state *after*
|
||||
// token ``t`` into cache slot ``t`` (multi-slot rollback, matching the GPU
|
||||
// kernel). Parallelized over (sequence, v_head); the per-sequence token loop is
|
||||
// sequential as required by the recurrence.
|
||||
template <typename scalar_t, typename param_t>
|
||||
void fused_sigmoid_gating_delta_rule_update_spec_kernel_impl(
|
||||
const scalar_t* __restrict__ q_ptr, // [T, HK, EK]
|
||||
const scalar_t* __restrict__ k_ptr, // [T, HK, EK]
|
||||
const scalar_t* __restrict__ v_ptr, // [T, HV, EV]
|
||||
const param_t* __restrict__ A_log_ptr,
|
||||
const scalar_t* __restrict__ a_ptr, // [T, HV]
|
||||
const scalar_t* __restrict__ dt_bias_ptr,
|
||||
const scalar_t* __restrict__ b_ptr, // [T, HV]
|
||||
const int32_t* __restrict__ spec_indices_ptr, // [N, S]
|
||||
const int32_t* __restrict__ num_accepted_ptr, // [N]
|
||||
const int32_t* __restrict__ cu_seqlens_ptr, // [N + 1]
|
||||
float* __restrict__ state_ptr,
|
||||
scalar_t* __restrict__ o_ptr, // [T, HV, EV]
|
||||
float* __restrict__ qk_scale_buf, // [2, T, HK]
|
||||
int64_t total_tokens,
|
||||
int64_t batch_size,
|
||||
int64_t spec_stride,
|
||||
int64_t num_heads,
|
||||
int64_t head_dim,
|
||||
int64_t v_num_heads,
|
||||
int64_t v_head_dim,
|
||||
int64_t q_strideT,
|
||||
int64_t q_strideH,
|
||||
int64_t k_strideT,
|
||||
int64_t k_strideH,
|
||||
int64_t v_strideT,
|
||||
int64_t v_strideH,
|
||||
int64_t state_slot_stride,
|
||||
bool use_qk_l2norm_in_kernel,
|
||||
double softplus_threshold) {
|
||||
using bVec = at::vec::Vectorized<scalar_t>;
|
||||
using fVec = at::vec::Vectorized<float>;
|
||||
constexpr int64_t VecSize = bVec::size();
|
||||
constexpr int64_t fVecSize = fVec::size();
|
||||
int64_t group_size = v_num_heads / num_heads;
|
||||
double scale = 1 / std::sqrt((double)head_dim);
|
||||
fVec scale_vec = fVec((float)scale);
|
||||
|
||||
if (use_qk_l2norm_in_kernel) {
|
||||
float eps = 1e-5f;
|
||||
at::parallel_for(0, total_tokens * num_heads, 0, [&](int64_t begin, int64_t end) {
|
||||
for (int64_t i = begin; i < end; ++i) {
|
||||
int64_t ti = i / num_heads;
|
||||
int64_t ni = i % num_heads;
|
||||
const scalar_t* qp = q_ptr + ti * q_strideT + ni * q_strideH;
|
||||
const scalar_t* kp = k_ptr + ti * k_strideT + ni * k_strideH;
|
||||
float sq = 0.f, sk = 0.f;
|
||||
for (int64_t d = 0; d < head_dim; ++d) {
|
||||
float qv = (float)qp[d];
|
||||
sq += qv * qv;
|
||||
float kv = (float)kp[d];
|
||||
sk += kv * kv;
|
||||
}
|
||||
qk_scale_buf[ti * num_heads + ni] = 1.f / std::sqrt(sq + eps);
|
||||
qk_scale_buf[total_tokens * num_heads + ti * num_heads + ni] = 1.f / std::sqrt(sk + eps);
|
||||
}
|
||||
});
|
||||
}
|
||||
|
||||
at::parallel_for(0, batch_size * v_num_heads, 0, [&](int64_t begin, int64_t end) {
|
||||
for (int64_t idx = begin; idx < end; ++idx) {
|
||||
int64_t bi = idx / v_num_heads;
|
||||
int64_t ni = idx % v_num_heads;
|
||||
int64_t kh = ni / group_size;
|
||||
int64_t q_start = cu_seqlens_ptr[bi];
|
||||
int64_t q_len = cu_seqlens_ptr[bi + 1] - q_start;
|
||||
if (q_len <= 0) {
|
||||
continue;
|
||||
}
|
||||
int64_t acc = (int64_t)num_accepted_ptr[bi];
|
||||
// Clamp acc-1 to >=0: when num_accepted is 0 the unclamped index reads
|
||||
// out of bounds and yields an arbitrary prev_slot used to index the SSM
|
||||
// state. Mirrors the GPU guard tl.maximum(num_accepted - 1, 0).
|
||||
int64_t prev_slot =
|
||||
(int64_t)spec_indices_ptr[bi * spec_stride + (acc > 0 ? acc - 1 : 0)];
|
||||
for (int64_t t = 0; t < q_len; ++t) {
|
||||
int64_t cur_slot = (int64_t)spec_indices_ptr[bi * spec_stride + t];
|
||||
int64_t token = q_start + t;
|
||||
const float* src = state_ptr + prev_slot * state_slot_stride + ni * head_dim * v_head_dim;
|
||||
float* dst = state_ptr + cur_slot * state_slot_stride + ni * head_dim * v_head_dim;
|
||||
float g_val = -std::exp((float)A_log_ptr[ni]) *
|
||||
softplus((float)a_ptr[token * v_num_heads + ni] + (float)dt_bias_ptr[ni], softplus_threshold);
|
||||
float g_val_exp = std::exp(g_val);
|
||||
fVec g_val_exp_vec = fVec(g_val_exp);
|
||||
float beta_val = 1.f / (1.f + std::exp(-(float)b_ptr[token * v_num_heads + ni]));
|
||||
fVec beta_vec = fVec(beta_val);
|
||||
int64_t q_offset = token * q_strideT + kh * q_strideH;
|
||||
int64_t k_offset = token * k_strideT + kh * k_strideH;
|
||||
float q_scale = use_qk_l2norm_in_kernel ? qk_scale_buf[token * num_heads + kh] : 1.f;
|
||||
float k_scale =
|
||||
use_qk_l2norm_in_kernel ? qk_scale_buf[total_tokens * num_heads + token * num_heads + kh] : 1.f;
|
||||
int64_t v_offset = token * v_strideT + ni * v_strideH;
|
||||
int64_t o_offset = (token * v_num_heads + ni) * v_head_dim;
|
||||
int64_t dvi = 0;
|
||||
for (; dvi <= v_head_dim - VecSize; dvi += VecSize) {
|
||||
fVec kv_mem_vec0 = fVec(0.f);
|
||||
fVec kv_mem_vec1 = fVec(0.f);
|
||||
for (int di = 0; di < head_dim; ++di) {
|
||||
fVec k_val_vec = fVec((float)k_ptr[k_offset + di] * k_scale);
|
||||
fVec sv0 = fVec::loadu(src + di * v_head_dim + dvi);
|
||||
fVec sv1 = fVec::loadu(src + di * v_head_dim + dvi + fVecSize);
|
||||
kv_mem_vec0 = kv_mem_vec0 + sv0 * g_val_exp_vec * k_val_vec;
|
||||
kv_mem_vec1 = kv_mem_vec1 + sv1 * g_val_exp_vec * k_val_vec;
|
||||
}
|
||||
bVec v_bvec = bVec::loadu(v_ptr + v_offset + dvi);
|
||||
fVec v_vec0, v_vec1;
|
||||
std::tie(v_vec0, v_vec1) = at::vec::convert_to_float(v_bvec);
|
||||
fVec dt_vec0 = (v_vec0 - kv_mem_vec0) * beta_vec;
|
||||
fVec dt_vec1 = (v_vec1 - kv_mem_vec1) * beta_vec;
|
||||
fVec o_vec0 = fVec(0.f);
|
||||
fVec o_vec1 = fVec(0.f);
|
||||
for (int di = 0; di < head_dim; ++di) {
|
||||
fVec q_vec = fVec((float)q_ptr[q_offset + di] * q_scale);
|
||||
fVec k_vec = fVec((float)k_ptr[k_offset + di] * k_scale);
|
||||
fVec sv0 = fVec::loadu(src + di * v_head_dim + dvi);
|
||||
fVec sv1 = fVec::loadu(src + di * v_head_dim + dvi + fVecSize);
|
||||
sv0 = sv0 * g_val_exp_vec + k_vec * dt_vec0;
|
||||
sv1 = sv1 * g_val_exp_vec + k_vec * dt_vec1;
|
||||
o_vec0 = o_vec0 + sv0 * q_vec * scale_vec;
|
||||
o_vec1 = o_vec1 + sv1 * q_vec * scale_vec;
|
||||
sv0.store(dst + di * v_head_dim + dvi);
|
||||
sv1.store(dst + di * v_head_dim + dvi + fVecSize);
|
||||
}
|
||||
bVec o_vec = at::vec::convert_from_float<scalar_t>(o_vec0, o_vec1);
|
||||
o_vec.store(o_ptr + o_offset + dvi);
|
||||
}
|
||||
for (; dvi < v_head_dim; ++dvi) {
|
||||
float kv_mem_val = 0.f;
|
||||
for (int di = 0; di < head_dim; ++di) {
|
||||
float k_val = (float)k_ptr[k_offset + di] * k_scale;
|
||||
kv_mem_val += src[di * v_head_dim + dvi] * g_val_exp * k_val;
|
||||
}
|
||||
float v_val = (float)v_ptr[v_offset + dvi];
|
||||
float dt_val = (v_val - kv_mem_val) * beta_val;
|
||||
float o_val = 0.f;
|
||||
for (int di = 0; di < head_dim; ++di) {
|
||||
float q_val = (float)q_ptr[q_offset + di] * q_scale;
|
||||
float k_val = (float)k_ptr[k_offset + di] * k_scale;
|
||||
float ns = src[di * v_head_dim + dvi] * g_val_exp + k_val * dt_val;
|
||||
dst[di * v_head_dim + dvi] = ns;
|
||||
o_val += ns * q_val * scale;
|
||||
}
|
||||
o_ptr[o_offset + dvi] = (scalar_t)o_val;
|
||||
}
|
||||
prev_slot = cur_slot;
|
||||
}
|
||||
}
|
||||
});
|
||||
}
|
||||
|
||||
template <typename scalar_t>
|
||||
void fused_gdn_gating_kernel_impl(
|
||||
float* __restrict__ A_log,
|
||||
@@ -1500,6 +1658,103 @@ at::Tensor fused_sigmoid_gating_delta_rule_update_cpu(
|
||||
return core_attn_out;
|
||||
}
|
||||
|
||||
// Speculative-decode update (multi-token, multi-slot rollback).
|
||||
// q: [T, HK, EK] k: [T, HK, EK] v: [T, HV, EV]
|
||||
// a: [T, HV] b: [T, HV]
|
||||
// initial_state_source: [N_slots, HV, EK, EV] FP32 (updated in place)
|
||||
// spec_state_indices: [batch, S] INT32 (S = num_spec + 1)
|
||||
// num_accepted_tokens: [batch] INT32
|
||||
// cu_seqlens: [batch + 1] INT32
|
||||
// Returns output: [T, HV, EV]
|
||||
at::Tensor fused_sigmoid_gating_delta_rule_update_spec_cpu(
|
||||
const at::Tensor& A_log,
|
||||
const at::Tensor& dt_bias,
|
||||
const at::Tensor& q,
|
||||
const at::Tensor& k,
|
||||
const at::Tensor& v,
|
||||
const at::Tensor& a,
|
||||
const at::Tensor& b,
|
||||
at::Tensor& initial_state_source,
|
||||
const at::Tensor& spec_state_indices,
|
||||
const at::Tensor& num_accepted_tokens,
|
||||
const at::Tensor& cu_seqlens,
|
||||
bool use_qk_l2norm_in_kernel,
|
||||
double softplus_beta = 1.0,
|
||||
double softplus_threshold = 20.0) {
|
||||
CHECK_DIM(3, q);
|
||||
CHECK_DIM(3, v);
|
||||
CHECK_LAST_DIM_CONTIGUOUS_INPUT(q);
|
||||
int64_t total_tokens = q.size(0);
|
||||
int64_t num_heads = q.size(1);
|
||||
int64_t head_dim = q.size(2);
|
||||
int64_t v_num_heads = v.size(1);
|
||||
int64_t v_head_dim = v.size(2);
|
||||
int64_t batch_size = cu_seqlens.size(0) - 1;
|
||||
int64_t spec_stride = spec_state_indices.stride(0);
|
||||
CHECK_INPUT_SHAPE_DTYPE<true>(k, {total_tokens, num_heads, head_dim}, q.scalar_type());
|
||||
CHECK_INPUT_SHAPE_DTYPE<true>(v, {total_tokens, v_num_heads, v_head_dim}, q.scalar_type());
|
||||
CHECK_INPUT_SHAPE_DTYPE<true>(a, {total_tokens, v_num_heads}, q.scalar_type());
|
||||
CHECK_INPUT_SHAPE_DTYPE<true>(b, {total_tokens, v_num_heads}, q.scalar_type());
|
||||
CHECK_INPUT_SHAPE_DTYPE<true>(dt_bias, {v_num_heads}, q.scalar_type());
|
||||
CHECK_INPUT_SHAPE_DTYPE<true>(num_accepted_tokens, {batch_size}, at::kInt);
|
||||
CHECK_INPUT_SHAPE_DTYPE<true>(cu_seqlens, {batch_size + 1}, at::kInt);
|
||||
CHECK_EQ(v_num_heads % num_heads, 0);
|
||||
TORCH_CHECK(A_log.sizes() == at::IntArrayRef({v_num_heads}));
|
||||
CHECK_INPUT_SHAPE_DTYPE<true>(
|
||||
initial_state_source,
|
||||
{initial_state_source.size(0), v_num_heads, head_dim, v_head_dim},
|
||||
at::kFloat);
|
||||
TORCH_CHECK(initial_state_source.size(0) >= batch_size,
|
||||
"initial_state_source capacity too small: size(0)=",
|
||||
initial_state_source.size(0), ", batch_size=", batch_size);
|
||||
|
||||
int64_t q_strideT = q.stride(0);
|
||||
int64_t q_strideH = q.stride(1);
|
||||
int64_t k_strideT = k.stride(0);
|
||||
int64_t k_strideH = k.stride(1);
|
||||
int64_t v_strideT = v.stride(0);
|
||||
int64_t v_strideH = v.stride(1);
|
||||
int64_t state_slot_stride = initial_state_source.stride(0);
|
||||
|
||||
at::Tensor o = at::empty({total_tokens, v_num_heads, v_head_dim}, q.options());
|
||||
at::Tensor qk_scale_buf = at::empty({2, total_tokens, num_heads}, at::kFloat);
|
||||
|
||||
CPU_DISPATCH_REDUCED_FLOATING_TYPES_EXT(
|
||||
q.scalar_type(), A_log.scalar_type(), "fused_sigmoid_gating_delta_rule_update_spec_kernel_impl", [&] {
|
||||
fused_sigmoid_gating_delta_rule_update_spec_kernel_impl<scalar_t, param_t>(
|
||||
q.data_ptr<scalar_t>(),
|
||||
k.data_ptr<scalar_t>(),
|
||||
v.data_ptr<scalar_t>(),
|
||||
A_log.data_ptr<param_t>(),
|
||||
a.data_ptr<scalar_t>(),
|
||||
dt_bias.data_ptr<scalar_t>(),
|
||||
b.data_ptr<scalar_t>(),
|
||||
spec_state_indices.data_ptr<int32_t>(),
|
||||
num_accepted_tokens.data_ptr<int32_t>(),
|
||||
cu_seqlens.data_ptr<int32_t>(),
|
||||
initial_state_source.data_ptr<float>(),
|
||||
o.data_ptr<scalar_t>(),
|
||||
qk_scale_buf.data_ptr<float>(),
|
||||
total_tokens,
|
||||
batch_size,
|
||||
spec_stride,
|
||||
num_heads,
|
||||
head_dim,
|
||||
v_num_heads,
|
||||
v_head_dim,
|
||||
q_strideT,
|
||||
q_strideH,
|
||||
k_strideT,
|
||||
k_strideH,
|
||||
v_strideT,
|
||||
v_strideH,
|
||||
state_slot_stride,
|
||||
use_qk_l2norm_in_kernel,
|
||||
softplus_threshold);
|
||||
});
|
||||
return o;
|
||||
}
|
||||
|
||||
// A_log: [num_v_heads]
|
||||
// a: [batch, num_v_heads]
|
||||
// b: [batch, num_v_heads]
|
||||
|
||||
@@ -120,6 +120,14 @@ at::Tensor fused_sigmoid_gating_delta_rule_update_cpu(
|
||||
bool use_qk_l2norm_in_kernel, double softplus_beta = 1.0,
|
||||
double softplus_threshold = 20.0);
|
||||
|
||||
at::Tensor fused_sigmoid_gating_delta_rule_update_spec_cpu(
|
||||
const at::Tensor& A_log, const at::Tensor& dt_bias, const at::Tensor& q,
|
||||
const at::Tensor& k, const at::Tensor& v, const at::Tensor& a,
|
||||
const at::Tensor& b, at::Tensor& initial_state_source,
|
||||
const at::Tensor& spec_state_indices, const at::Tensor& num_accepted_tokens,
|
||||
const at::Tensor& cu_seqlens, bool use_qk_l2norm_in_kernel,
|
||||
double softplus_beta = 1.0, double softplus_threshold = 20.0);
|
||||
|
||||
std::tuple<at::Tensor, at::Tensor> fused_gdn_gating_cpu(
|
||||
const at::Tensor& A_log, const at::Tensor& a, const at::Tensor& b,
|
||||
const at::Tensor& dt_bias);
|
||||
@@ -508,6 +516,15 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
|
||||
"softplus_threshold=20.0) -> Tensor");
|
||||
ops.impl("fused_sigmoid_gating_delta_rule_update_cpu", torch::kCPU,
|
||||
&fused_sigmoid_gating_delta_rule_update_cpu);
|
||||
ops.def(
|
||||
"fused_sigmoid_gating_delta_rule_update_spec_cpu(Tensor A_log, Tensor "
|
||||
"dt_bias, Tensor q, Tensor k, Tensor v, Tensor a, Tensor b, "
|
||||
"Tensor(a!) initial_state_source, Tensor spec_state_indices, "
|
||||
"Tensor num_accepted_tokens, Tensor cu_seqlens, bool "
|
||||
"use_qk_l2norm_in_kernel, float softplus_beta=1.0, float "
|
||||
"softplus_threshold=20.0) -> Tensor");
|
||||
ops.impl("fused_sigmoid_gating_delta_rule_update_spec_cpu", torch::kCPU,
|
||||
&fused_sigmoid_gating_delta_rule_update_spec_cpu);
|
||||
ops.def(
|
||||
"fused_gdn_gating_cpu(Tensor A_log, Tensor a, Tensor b, Tensor dt_bias) "
|
||||
"-> (Tensor, Tensor)");
|
||||
|
||||
@@ -691,19 +691,20 @@ void launchFusedDeepseekV4QNormRopeKVRopeQuantInsert(
|
||||
//
|
||||
// Sibling to the FlashMLA kernel above, used by the FlashInfer V4 sparse-MLA
|
||||
// backend. Differences from the legacy path:
|
||||
// * No Q head padding — output Q layout matches the input num_heads_q.
|
||||
// * Optional FP8 Q head padding, with zero-filled padding heads written
|
||||
// directly from this fused kernel.
|
||||
// * KV is written as a *contiguous* 512-wide row per token (token-strided),
|
||||
// not the legacy UE8M0 paged layout with a separate scale tail.
|
||||
// * Q/KV are stored either as bf16 or as per-tensor E4M3 FP8 (one global
|
||||
// scale), selected by the STORE_Q_FP8 / STORE_KV_FP8 template flags.
|
||||
//
|
||||
// Grid: 1D, gridDim.x = ceil(num_tokens_full * (num_heads_q + 1) / warps).
|
||||
// Each warp handles one (token, slot): slot < num_heads_q → Q, slot ==
|
||||
// num_heads_q → KV.
|
||||
// Grid: 1D, gridDim.x = ceil(num_tokens_full * (num_heads_q_padded + 1) /
|
||||
// warps). Each warp handles one (token, slot): slot < num_heads_q → live Q,
|
||||
// num_heads_q <= slot < num_heads_q_padded → padded Q, final slot → KV.
|
||||
template <typename scalar_t_in, bool STORE_Q_FP8, bool STORE_KV_FP8>
|
||||
__global__ void fusedDeepseekV4FullCacheKernel(
|
||||
scalar_t_in* __restrict__ q_inout, // [N, H, 512], in place (bf16)
|
||||
uint8_t* __restrict__ q_fp8_out, // [N, H, 512] fp8, optional
|
||||
uint8_t* __restrict__ q_fp8_out, // [N, H_out, 512] fp8, optional
|
||||
int64_t const q_fp8_stride0, // elements (fp8 == bytes)
|
||||
int64_t const q_fp8_stride1, // elements (fp8 == bytes)
|
||||
scalar_t_in const* __restrict__ kv_in, // [N, 512] bf16
|
||||
@@ -716,7 +717,8 @@ __global__ void fusedDeepseekV4FullCacheKernel(
|
||||
float const eps,
|
||||
int const num_tokens_full, // = q.size(0) = kv.size(0)
|
||||
int const num_tokens_insert, // = slot_mapping.size(0)
|
||||
int const num_heads_q, // H (no padding)
|
||||
int const num_heads_q, // real H
|
||||
int const num_heads_q_padded, // output/kernel H
|
||||
int const cache_block_size, // tokens per cache block
|
||||
int64_t const kv_block_stride, // bytes per cache block
|
||||
int64_t const kv_token_stride) { // bytes per cache token
|
||||
@@ -731,11 +733,12 @@ __global__ void fusedDeepseekV4FullCacheKernel(
|
||||
int const laneId = threadIdx.x % 32;
|
||||
int const globalWarpIdx = blockIdx.x * warpsPerBlock + warpId;
|
||||
|
||||
int const slotsPerToken = num_heads_q + 1;
|
||||
int const slotsPerToken = num_heads_q_padded + 1;
|
||||
int const tokenIdx = globalWarpIdx / slotsPerToken;
|
||||
int const slotIdx = globalWarpIdx % slotsPerToken;
|
||||
if (tokenIdx >= num_tokens_full) return;
|
||||
bool const isKV = (slotIdx == num_heads_q);
|
||||
bool const isKV = (slotIdx == num_heads_q_padded);
|
||||
bool const isQPad = (slotIdx >= num_heads_q) && !isKV;
|
||||
// KV branch: skip DP-padded tokens (no slot reserved for them).
|
||||
if (isKV && tokenIdx >= num_tokens_insert) return;
|
||||
|
||||
@@ -744,6 +747,19 @@ __global__ void fusedDeepseekV4FullCacheKernel(
|
||||
#endif
|
||||
|
||||
int const dim_base = laneId * kElemsPerLane; // in [0, 512) step 16
|
||||
|
||||
if (isQPad) {
|
||||
// padded position zero_
|
||||
uint4 const zero{};
|
||||
if constexpr (STORE_Q_FP8) {
|
||||
uint8_t* dst = q_fp8_out +
|
||||
static_cast<int64_t>(tokenIdx) * q_fp8_stride0 +
|
||||
static_cast<int64_t>(slotIdx) * q_fp8_stride1 + dim_base;
|
||||
*reinterpret_cast<uint4*>(dst) = zero;
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
||||
scalar_t_in const* src_ptr;
|
||||
if (isKV) {
|
||||
src_ptr = kv_in + static_cast<int64_t>(tokenIdx) * kHeadDim + dim_base;
|
||||
@@ -842,7 +858,8 @@ __global__ void fusedDeepseekV4FullCacheKernel(
|
||||
}
|
||||
scalar_t_in* dst =
|
||||
q_inout +
|
||||
(static_cast<int64_t>(tokenIdx) * num_heads_q + slotIdx) * kHeadDim +
|
||||
(static_cast<int64_t>(tokenIdx) * num_heads_q + slotIdx) *
|
||||
kHeadDim +
|
||||
dim_base;
|
||||
*reinterpret_cast<uint4*>(dst) = out0;
|
||||
*reinterpret_cast<uint4*>(dst + 8) = out1;
|
||||
@@ -898,12 +915,13 @@ static void launchFullCacheKernel(
|
||||
float const* cos_sin_cache, float const* fp8_scale,
|
||||
float const* q_fp8_scale_inv, float const eps, int const num_tokens_full,
|
||||
int const num_tokens_insert, int const num_heads_q,
|
||||
int const cache_block_size, int64_t const kv_block_stride,
|
||||
int64_t const kv_token_stride, char const* op_name, cudaStream_t stream) {
|
||||
int const num_heads_q_padded, int const cache_block_size,
|
||||
int64_t const kv_block_stride, int64_t const kv_token_stride,
|
||||
char const* op_name, cudaStream_t stream) {
|
||||
constexpr int kBlockSize = 256;
|
||||
constexpr int kWarpsPerBlock = kBlockSize / 32;
|
||||
int64_t const total_warps =
|
||||
static_cast<int64_t>(num_tokens_full) * (num_heads_q + 1);
|
||||
static_cast<int64_t>(num_tokens_full) * (num_heads_q_padded + 1);
|
||||
int const grid =
|
||||
static_cast<int>((total_warps + kWarpsPerBlock - 1) / kWarpsPerBlock);
|
||||
auto* kernel =
|
||||
@@ -926,13 +944,14 @@ static void launchFullCacheKernel(
|
||||
q_fp8_stride1, kv_in, k_cache, slot_mapping, position_ids,
|
||||
cos_sin_cache, fp8_scale, q_fp8_scale_inv, eps,
|
||||
num_tokens_full, num_tokens_insert, num_heads_q,
|
||||
cache_block_size, kv_block_stride, kv_token_stride);
|
||||
num_heads_q_padded, cache_block_size, kv_block_stride,
|
||||
kv_token_stride);
|
||||
#else
|
||||
kernel<<<grid, kBlockSize, 0, stream>>>(
|
||||
q_inout, q_fp8_out, q_fp8_stride0, q_fp8_stride1, kv_in, k_cache,
|
||||
slot_mapping, position_ids, cos_sin_cache, fp8_scale, q_fp8_scale_inv,
|
||||
eps, num_tokens_full, num_tokens_insert, num_heads_q, cache_block_size,
|
||||
kv_block_stride, kv_token_stride);
|
||||
slot_mapping, position_ids, cos_sin_cache, fp8_scale,
|
||||
q_fp8_scale_inv, eps, num_tokens_full, num_tokens_insert, num_heads_q,
|
||||
num_heads_q_padded, cache_block_size, kv_block_stride, kv_token_stride);
|
||||
#endif
|
||||
}
|
||||
|
||||
@@ -1090,8 +1109,8 @@ void fused_deepseek_v4_qnorm_rope_kv_rope_full_cache_bf16_insert(
|
||||
position_ids.const_data_ptr<int64_t>(),
|
||||
cos_sin_cache.const_data_ptr<float>(), nullptr, nullptr,
|
||||
static_cast<float>(eps), num_tokens_full, num_tokens_insert,
|
||||
num_heads_q, static_cast<int>(cache_block_size), kv_block_stride,
|
||||
kv_token_stride,
|
||||
num_heads_q, num_heads_q,
|
||||
static_cast<int>(cache_block_size), kv_block_stride, kv_token_stride,
|
||||
"fused_deepseek_v4_qnorm_rope_kv_rope_full_cache_bf16_insert",
|
||||
stream);
|
||||
});
|
||||
@@ -1116,8 +1135,9 @@ void fused_deepseek_v4_qnorm_rope_kv_rope_full_cache_fp8_insert(
|
||||
STD_TORCH_CHECK(q_fp8.device().is_cuda() && q_fp8.is_contiguous() &&
|
||||
q_fp8.scalar_type() == ScalarType::Float8_e4m3fn &&
|
||||
q_fp8.dim() == 3 && q_fp8.size(0) == q.size(0) &&
|
||||
q_fp8.size(1) == q.size(1) && q_fp8.size(2) == q.size(2),
|
||||
"q_fp8 must be a contiguous float8_e4m3fn tensor matching q");
|
||||
q_fp8.size(1) >= q.size(1) && q_fp8.size(2) == q.size(2),
|
||||
"q_fp8 must be a contiguous float8_e4m3fn tensor with shape "
|
||||
"[N, H_out, 512] and H_out >= q.size(1)");
|
||||
STD_TORCH_CHECK(k_cache.device().is_cuda(), "k_cache must be CUDA");
|
||||
STD_TORCH_CHECK(slot_mapping.device().is_cuda() &&
|
||||
slot_mapping.scalar_type() == ScalarType::Long,
|
||||
@@ -1155,6 +1175,7 @@ void fused_deepseek_v4_qnorm_rope_kv_rope_full_cache_fp8_insert(
|
||||
STD_TORCH_CHECK(num_tokens_insert <= num_tokens_full,
|
||||
"slot_mapping must not exceed q row count");
|
||||
int const num_heads_q = static_cast<int>(q.size(1));
|
||||
int const num_heads_q_padded = static_cast<int>(q_fp8.size(1));
|
||||
|
||||
const torch::stable::accelerator::DeviceGuard device_guard(
|
||||
q.get_device_index());
|
||||
@@ -1167,8 +1188,7 @@ void fused_deepseek_v4_qnorm_rope_kv_rope_full_cache_fp8_insert(
|
||||
true>(
|
||||
// q is read-only in the fp8 path (the kernel writes q_fp8); the
|
||||
// launcher signature is non-const, so cast away const on the ptr.
|
||||
reinterpret_cast<scalar_t*>(
|
||||
const_cast<void*>(q.const_data_ptr())),
|
||||
reinterpret_cast<scalar_t*>(const_cast<void*>(q.const_data_ptr())),
|
||||
reinterpret_cast<uint8_t*>(q_fp8.mutable_data_ptr()),
|
||||
q_fp8.stride(0), q_fp8.stride(1),
|
||||
reinterpret_cast<scalar_t const*>(kv.const_data_ptr()),
|
||||
@@ -1179,7 +1199,7 @@ void fused_deepseek_v4_qnorm_rope_kv_rope_full_cache_fp8_insert(
|
||||
fp8_scale.const_data_ptr<float>(),
|
||||
q_fp8_scale_inv.const_data_ptr<float>(), static_cast<float>(eps),
|
||||
num_tokens_full, num_tokens_insert, num_heads_q,
|
||||
static_cast<int>(cache_block_size),
|
||||
num_heads_q_padded, static_cast<int>(cache_block_size),
|
||||
// fp8 cache: 1 byte/element -> stride already in bytes.
|
||||
k_cache.stride(0), k_cache.stride(1),
|
||||
"fused_deepseek_v4_qnorm_rope_kv_rope_full_cache_fp8_insert",
|
||||
|
||||
+111
-17
@@ -1,5 +1,7 @@
|
||||
# default base image
|
||||
ARG REMOTE_VLLM="0"
|
||||
ARG VLLM_REPO="https://github.com/vllm-project/vllm.git"
|
||||
ARG VLLM_BRANCH="main"
|
||||
ARG COMMON_WORKDIR=/app
|
||||
ARG BASE_IMAGE=rocm/vllm-dev:base
|
||||
ARG CI_BASE_IMAGE=rocm/vllm-dev:ci_base
|
||||
@@ -110,10 +112,8 @@ WORKDIR ${COMMON_WORKDIR}
|
||||
FROM base AS fetch_vllm_0
|
||||
ONBUILD COPY ./ vllm/
|
||||
FROM base AS fetch_vllm_1
|
||||
ARG VLLM_REPO="https://github.com/vllm-project/vllm.git"
|
||||
ARG VLLM_BRANCH="main"
|
||||
ENV VLLM_REPO=${VLLM_REPO}
|
||||
ENV VLLM_BRANCH=${VLLM_BRANCH}
|
||||
ARG VLLM_REPO
|
||||
ARG VLLM_BRANCH
|
||||
ONBUILD RUN git clone ${VLLM_REPO} \
|
||||
&& cd vllm \
|
||||
&& git fetch -v --prune -- origin ${VLLM_BRANCH} \
|
||||
@@ -124,34 +124,118 @@ ONBUILD RUN git clone ${VLLM_REPO} \
|
||||
FROM fetch_vllm_${REMOTE_VLLM} AS fetch_vllm
|
||||
|
||||
# -----------------------
|
||||
# Rust build stage
|
||||
# Builds the `vllm-rs` frontend in a dedicated stage so the wheel build stages
|
||||
# don't need the rust toolchain or protoc.
|
||||
FROM fetch_vllm AS rust-build
|
||||
ARG COMMON_WORKDIR
|
||||
ARG USE_SCCACHE
|
||||
# Rust/protoc toolchain inputs
|
||||
#
|
||||
# Keep these separate from Rust source inputs so Rust source-only changes do not
|
||||
# invalidate the long-lived ci_base image that carries the toolchain.
|
||||
FROM scratch AS rust_toolchain_input_0
|
||||
COPY rust-toolchain.toml /rust-toolchain-input/vllm/rust-toolchain.toml
|
||||
COPY tools/install_protoc.sh /rust-toolchain-input/vllm/tools/install_protoc.sh
|
||||
|
||||
FROM base AS rust_toolchain_input_1
|
||||
ARG VLLM_REPO
|
||||
ARG VLLM_BRANCH
|
||||
RUN git clone --no-checkout --filter=blob:none ${VLLM_REPO} /rust-toolchain-input/vllm \
|
||||
&& cd /rust-toolchain-input/vllm \
|
||||
&& git fetch -v --prune -- origin ${VLLM_BRANCH} \
|
||||
&& git sparse-checkout init --no-cone \
|
||||
&& git sparse-checkout set \
|
||||
rust-toolchain.toml \
|
||||
tools/install_protoc.sh \
|
||||
&& git checkout FETCH_HEAD
|
||||
|
||||
FROM rust_toolchain_input_${REMOTE_VLLM} AS rust-toolchain-input
|
||||
|
||||
# -----------------------
|
||||
# Rust build inputs
|
||||
#
|
||||
# REMOTE_VLLM=0: use only the local files that can affect Rust artifacts.
|
||||
# REMOTE_VLLM=1: fetch only those same files from the requested remote ref.
|
||||
FROM scratch AS rust_input_0
|
||||
COPY requirements/build/rust.txt /rust-input/vllm/requirements/build/rust.txt
|
||||
COPY rust/Cargo.lock rust/Cargo.toml /rust-input/vllm/rust/
|
||||
COPY rust/proto /rust-input/vllm/rust/proto
|
||||
COPY rust/src /rust-input/vllm/rust/src
|
||||
COPY rust-toolchain.toml /rust-input/vllm/rust-toolchain.toml
|
||||
COPY tools/build_rust.py /rust-input/vllm/tools/build_rust.py
|
||||
COPY build_rust.sh /rust-input/vllm/build_rust.sh
|
||||
|
||||
FROM base AS rust_input_1
|
||||
ARG VLLM_REPO
|
||||
ARG VLLM_BRANCH
|
||||
RUN git clone --no-checkout --filter=blob:none ${VLLM_REPO} /rust-input/vllm \
|
||||
&& cd /rust-input/vllm \
|
||||
&& git fetch -v --prune -- origin ${VLLM_BRANCH} \
|
||||
&& git sparse-checkout init --no-cone \
|
||||
&& git sparse-checkout set \
|
||||
requirements/build/rust.txt \
|
||||
rust/Cargo.lock \
|
||||
rust/Cargo.toml \
|
||||
rust/proto/** \
|
||||
rust/src/** \
|
||||
rust-toolchain.toml \
|
||||
tools/build_rust.py \
|
||||
build_rust.sh \
|
||||
&& git checkout FETCH_HEAD
|
||||
|
||||
FROM rust_input_${REMOTE_VLLM} AS rust-input
|
||||
|
||||
# -----------------------
|
||||
# Rust/protoc toolchain
|
||||
FROM base AS rust-toolchain
|
||||
|
||||
ENV CARGO_HOME=/root/.cargo
|
||||
ENV RUSTUP_HOME=/root/.rustup
|
||||
ENV PATH=${CARGO_HOME}/bin:${PATH}
|
||||
|
||||
# protoc is used by tonic-build/prost-build.
|
||||
RUN apt-get update -q -y && apt-get install -q -y --no-install-recommends \
|
||||
ca-certificates curl unzip \
|
||||
&& rm -rf /var/lib/apt/lists/*
|
||||
|
||||
COPY tools/install_protoc.sh /tmp/install_protoc.sh
|
||||
COPY --from=rust-toolchain-input /rust-toolchain-input/vllm/tools/install_protoc.sh /tmp/install_protoc.sh
|
||||
RUN /tmp/install_protoc.sh && rm /tmp/install_protoc.sh
|
||||
|
||||
COPY --from=rust-toolchain-input /rust-toolchain-input/vllm/rust-toolchain.toml /tmp/rust-toolchain.toml
|
||||
RUN TOOLCHAIN="$(grep '^channel' /tmp/rust-toolchain.toml | sed 's/.*= *"\(.*\)"/\1/')" \
|
||||
&& if ! command -v rustup >/dev/null 2>&1; then \
|
||||
curl --proto '=https' --tlsv1.2 -sSf https://sh.rustup.rs \
|
||||
| sh -s -- -y --default-toolchain none; \
|
||||
fi \
|
||||
&& rustup toolchain install "${TOOLCHAIN}" \
|
||||
&& rustup default "${TOOLCHAIN}" \
|
||||
&& rustc --version \
|
||||
&& cargo --version
|
||||
|
||||
# -----------------------
|
||||
# Rust build stage
|
||||
# Builds the `vllm-rs` frontend in a dedicated stage so the wheel build stages
|
||||
# don't need the rust toolchain or protoc.
|
||||
FROM rust-toolchain AS rust-build
|
||||
ARG COMMON_WORKDIR
|
||||
ARG USE_SCCACHE
|
||||
|
||||
# Cap cargo parallelism to avoid exhausting the AMD CI host's open-file limit
|
||||
# (rustc spawns enough concurrent processes to hit RLIMIT_NOFILE otherwise).
|
||||
ENV CARGO_BUILD_JOBS=4
|
||||
ENV CARGO_NET_RETRY=10
|
||||
ENV RUSTUP_MAX_RETRIES=10
|
||||
|
||||
WORKDIR ${COMMON_WORKDIR}/vllm
|
||||
|
||||
COPY --from=rust-input /rust-input/vllm/requirements/build/rust.txt requirements/build/rust.txt
|
||||
|
||||
# BuildKit can run this stage in parallel with ROCm native builds. Keep Rust on
|
||||
# a separate local sccache daemon while sharing the same remote cache backend.
|
||||
ENV SCCACHE_SERVER_PORT=4227
|
||||
|
||||
RUN --mount=type=cache,id=vllm-rocm-uv,target=/root/.cache/uv \
|
||||
cd ${COMMON_WORKDIR}/vllm \
|
||||
&& uv pip install --system -r requirements/build/rust.txt
|
||||
uv pip install --system -r requirements/build/rust.txt
|
||||
|
||||
COPY --from=rust-input /rust-input/vllm/rust rust
|
||||
COPY --from=rust-input /rust-input/vllm/rust-toolchain.toml rust-toolchain.toml
|
||||
COPY --from=rust-input /rust-input/vllm/tools/build_rust.py tools/build_rust.py
|
||||
COPY --from=rust-input /rust-input/vllm/build_rust.sh build_rust.sh
|
||||
|
||||
# Build the release binary. Cargo's registry/git caches can be written by
|
||||
# concurrent BuildKit jobs on shared workers, so lock those cache mounts while
|
||||
@@ -159,8 +243,7 @@ RUN --mount=type=cache,id=vllm-rocm-uv,target=/root/.cache/uv \
|
||||
# can outlive source updates across BuildKit cache reuse.
|
||||
RUN --mount=type=cache,id=vllm-rocm-cargo-registry,target=/root/.cargo/registry,sharing=locked \
|
||||
--mount=type=cache,id=vllm-rocm-cargo-git,target=/root/.cargo/git,sharing=locked \
|
||||
cd ${COMMON_WORKDIR}/vllm \
|
||||
&& if [ "$USE_SCCACHE" = "1" ]; then \
|
||||
if [ "$USE_SCCACHE" = "1" ]; then \
|
||||
export RUSTC_WRAPPER=sccache \
|
||||
&& sccache --show-stats; \
|
||||
fi \
|
||||
@@ -185,8 +268,8 @@ COPY requirements/common.txt requirements/common.txt
|
||||
RUN --mount=type=cache,id=vllm-rocm-uv,target=/root/.cache/uv \
|
||||
uv pip install --system -r requirements/rocm.txt
|
||||
|
||||
# pyproject.toml is bind-mounted in the RUN step so metadata-only changes do
|
||||
# not invalidate the expensive native build layer.
|
||||
# pyproject.toml is bind-mounted so it is not persisted in the csrc-build image,
|
||||
# but it still participates in the BuildKit cache key for this RUN step.
|
||||
COPY setup.py CMakeLists.txt ./
|
||||
COPY tools/build_rust.py tools/build_rust.py
|
||||
COPY cmake cmake/
|
||||
@@ -556,6 +639,17 @@ COPY --from=build_vllm_wheel_release ${COMMON_WORKDIR}/vllm/vllm/v1 /vllm_v1
|
||||
FROM mori_base AS ci_base
|
||||
ARG COMMON_WORKDIR
|
||||
|
||||
ENV CARGO_HOME=/root/.cargo
|
||||
ENV RUSTUP_HOME=/root/.rustup
|
||||
ENV PATH=${CARGO_HOME}/bin:${PATH}
|
||||
|
||||
# Bake the Rust/protoc toolchain into ci_base so Rust build/test steps do not
|
||||
# need to fetch it on each per-commit image build.
|
||||
COPY --from=rust-toolchain /root/.cargo /root/.cargo
|
||||
COPY --from=rust-toolchain /root/.rustup /root/.rustup
|
||||
COPY --from=rust-toolchain /usr/local/bin/protoc /usr/local/bin/protoc
|
||||
COPY --from=rust-toolchain /usr/local/include/google /usr/local/include/google
|
||||
|
||||
# Update rdma-core to support latest rocshmem.
|
||||
ARG DEEPEP_NIC
|
||||
RUN if [ "${DEEPEP_NIC}" = "cx7" ] || [ "${DEEPEP_NIC}" = "io" ]; then \
|
||||
|
||||
+4
-97
@@ -37,37 +37,6 @@ RUN microdnf install -y \
|
||||
python${PYTHON_VERSION}-devel python${PYTHON_VERSION}-pip python${PYTHON_VERSION}-wheel && \
|
||||
python${PYTHON_VERSION} -m venv $VIRTUAL_ENV && pip install --no-cache -U pip wheel uv && microdnf clean all
|
||||
|
||||
FROM python-install AS pyarrow
|
||||
|
||||
# Build Apache Arrow
|
||||
WORKDIR /tmp
|
||||
RUN --mount=type=cache,target=/root/.cache/uv \
|
||||
git clone https://github.com/apache/arrow.git -b maint-19.0.1 && \
|
||||
cd arrow/cpp && \
|
||||
mkdir release && cd release && \
|
||||
cmake -DCMAKE_BUILD_TYPE=Release \
|
||||
-DCMAKE_INSTALL_PREFIX=/usr/local \
|
||||
-DARROW_PYTHON=ON \
|
||||
-DARROW_PARQUET=ON \
|
||||
-DARROW_ORC=ON \
|
||||
-DARROW_FILESYSTEM=ON \
|
||||
-DARROW_WITH_LZ4=ON \
|
||||
-DARROW_WITH_ZSTD=ON \
|
||||
-DARROW_WITH_SNAPPY=ON \
|
||||
-DARROW_JSON=ON \
|
||||
-DARROW_CSV=ON \
|
||||
-DARROW_DATASET=ON \
|
||||
-DPROTOBUF_PROTOC_EXECUTABLE=/usr/bin/protoc \
|
||||
-DARROW_DEPENDENCY_SOURCE=BUNDLED \
|
||||
.. && \
|
||||
make -j$(nproc) && \
|
||||
make install && \
|
||||
cd ../../python && \
|
||||
export PYARROW_PARALLEL=4 && \
|
||||
export ARROW_BUILD_TYPE=release && \
|
||||
uv pip install -r requirements-build.txt && \
|
||||
python setup.py build_ext --build-type=$ARROW_BUILD_TYPE --bundle-arrow-cpp bdist_wheel
|
||||
|
||||
FROM python-install AS rust
|
||||
ENV CARGO_HOME=/root/.cargo
|
||||
ENV RUSTUP_HOME=/root/.rustup
|
||||
@@ -121,60 +90,18 @@ RUN --mount=type=cache,target=/root/.cache/uv \
|
||||
FROM python-install AS numba-builder
|
||||
|
||||
ARG MAX_JOBS
|
||||
ARG NUMBA_VERSION=0.61.2
|
||||
ARG NUMBA_VERSION=0.65.0
|
||||
|
||||
WORKDIR /tmp
|
||||
|
||||
# Clone all required dependencies
|
||||
RUN --mount=type=cache,target=/root/.cache/uv \
|
||||
microdnf install ninja-build gcc gcc-c++ -y && \
|
||||
git clone --recursive https://github.com/llvm/llvm-project.git -b llvmorg-15.0.7 && \
|
||||
git clone --recursive https://github.com/numba/llvmlite.git -b v0.44.0 && \
|
||||
git clone --recursive https://github.com/numba/llvmlite.git -b v0.47.0 && \
|
||||
git clone --recursive https://github.com/numba/numba.git -b ${NUMBA_VERSION} && \
|
||||
cd llvm-project && mkdir build && cd build && \
|
||||
cd llvmlite && \
|
||||
uv pip install 'cmake<4' 'setuptools<70' numpy && \
|
||||
export PREFIX=/usr/local && CMAKE_ARGS="${CMAKE_ARGS} -DLLVM_ENABLE_PROJECTS=lld;libunwind;compiler-rt" \
|
||||
CFLAGS="$(echo $CFLAGS | sed 's/-fno-plt //g')" \
|
||||
CXXFLAGS="$(echo $CXXFLAGS | sed 's/-fno-plt //g')" \
|
||||
CMAKE_ARGS="${CMAKE_ARGS} -DFFI_INCLUDE_DIR=$PREFIX/include" \
|
||||
CMAKE_ARGS="${CMAKE_ARGS} -DFFI_LIBRARY_DIR=$PREFIX/lib" \
|
||||
cmake -DCMAKE_INSTALL_PREFIX="${PREFIX}" \
|
||||
-DCMAKE_BUILD_TYPE=Release \
|
||||
-DCMAKE_LIBRARY_PATH="${PREFIX}" \
|
||||
-DLLVM_ENABLE_LIBEDIT=OFF \
|
||||
-DLLVM_ENABLE_LIBXML2=OFF \
|
||||
-DLLVM_ENABLE_RTTI=ON \
|
||||
-DLLVM_ENABLE_TERMINFO=OFF \
|
||||
-DLLVM_INCLUDE_BENCHMARKS=OFF \
|
||||
-DLLVM_INCLUDE_DOCS=OFF \
|
||||
-DLLVM_INCLUDE_EXAMPLES=OFF \
|
||||
-DLLVM_INCLUDE_GO_TESTS=OFF \
|
||||
-DLLVM_INCLUDE_TESTS=OFF \
|
||||
-DLLVM_INCLUDE_UTILS=ON \
|
||||
-DLLVM_INSTALL_UTILS=ON \
|
||||
-DLLVM_UTILS_INSTALL_DIR=libexec/llvm \
|
||||
-DLLVM_BUILD_LLVM_DYLIB=OFF \
|
||||
-DLLVM_LINK_LLVM_DYLIB=OFF \
|
||||
-DLLVM_EXPERIMENTAL_TARGETS_TO_BUILD=WebAssembly \
|
||||
-DLLVM_ENABLE_FFI=ON \
|
||||
-DLLVM_ENABLE_Z3_SOLVER=OFF \
|
||||
-DLLVM_OPTIMIZED_TABLEGEN=ON \
|
||||
-DCMAKE_POLICY_DEFAULT_CMP0111=NEW \
|
||||
-DCOMPILER_RT_BUILD_BUILTINS=ON \
|
||||
-DCOMPILER_RT_BUILTINS_HIDE_SYMBOLS=OFF \
|
||||
-DCOMPILER_RT_BUILD_LIBFUZZER=OFF \
|
||||
-DCOMPILER_RT_BUILD_CRT=OFF \
|
||||
-DCOMPILER_RT_BUILD_MEMPROF=OFF \
|
||||
-DCOMPILER_RT_BUILD_PROFILE=OFF \
|
||||
-DCOMPILER_RT_BUILD_SANITIZERS=OFF \
|
||||
-DCOMPILER_RT_BUILD_XRAY=OFF \
|
||||
-DCOMPILER_RT_BUILD_GWP_ASAN=OFF \
|
||||
-DCOMPILER_RT_BUILD_ORC=OFF \
|
||||
-DCOMPILER_RT_INCLUDE_TESTS=OFF \
|
||||
${CMAKE_ARGS} -GNinja ../llvm \
|
||||
&& ninja install . && \
|
||||
# build llvmlite
|
||||
cd ../../llvmlite && python setup.py bdist_wheel && \
|
||||
python setup.py bdist_wheel && \
|
||||
cd ../numba && \
|
||||
if ! grep '#include "dynamic_annotations.h"' numba/_dispatcher.cpp; then \
|
||||
sed -i '/#include "internal\/pycore_atomic.h"/i\#include "dynamic_annotations.h"' numba/_dispatcher.cpp; \
|
||||
@@ -192,21 +119,6 @@ RUN --mount=type=cache,target=/root/.cache/uv \
|
||||
cd opencv-python && \
|
||||
python -m build --wheel --installer=uv --outdir /tmp/opencv-python/dist
|
||||
|
||||
## Todo(r3hankhan123): Remove guidance-builder stage once vLLM upgrades to new version of llguidance that fixes s390x issues. See https://github.com/guidance-ai/llguidance/issues/330
|
||||
FROM python-install AS guidance-builder
|
||||
WORKDIR /tmp
|
||||
ENV CARGO_HOME=/root/.cargo
|
||||
ENV RUSTUP_HOME=/root/.rustup
|
||||
ENV PATH="$CARGO_HOME/bin:$RUSTUP_HOME/bin:$PATH"
|
||||
RUN --mount=type=cache,target=/root/.cache/uv \
|
||||
--mount=type=bind,from=rust,source=/root/.cargo,target=/root/.cargo,rw \
|
||||
--mount=type=bind,from=rust,source=/root/.rustup,target=/root/.rustup,rw \
|
||||
git clone https://github.com/guidance-ai/llguidance.git && \
|
||||
cd llguidance && \
|
||||
git checkout s390x-fix-v2 && \
|
||||
uv pip install maturin && \
|
||||
python -m maturin build --release --out dist --compatibility linux
|
||||
|
||||
# # Final build stage
|
||||
FROM python-install AS vllm-cpu
|
||||
ARG PYTHON_VERSION
|
||||
@@ -235,20 +147,16 @@ RUN --mount=type=bind,from=numa-build,src=/tmp/numactl-2.0.19,target=/numactl \
|
||||
RUN --mount=type=cache,target=/root/.cache/uv \
|
||||
--mount=type=bind,from=rust,source=/root/.cargo,target=/root/.cargo,rw \
|
||||
--mount=type=bind,from=rust,source=/root/.rustup,target=/root/.rustup,rw \
|
||||
--mount=type=bind,from=pyarrow,source=/tmp/arrow/python/dist,target=/tmp/arrow-wheels \
|
||||
--mount=type=bind,from=torch-vision,source=/tmp/vision/dist,target=/tmp/vision-wheels/ \
|
||||
--mount=type=bind,from=hf-xet-builder,source=/tmp/hf-xet/dist,target=/tmp/hf-xet-wheels/ \
|
||||
--mount=type=bind,from=numba-builder,source=/tmp/llvmlite/dist,target=/tmp/llvmlite-wheels/ \
|
||||
--mount=type=bind,from=numba-builder,source=/tmp/numba/dist,target=/tmp/numba-wheels/ \
|
||||
--mount=type=bind,from=opencv-builder,source=/tmp/opencv-python/dist,target=/tmp/opencv-wheels/ \
|
||||
--mount=type=bind,from=guidance-builder,source=/tmp/llguidance/dist,target=/tmp/guidance-wheels/ \
|
||||
ARROW_WHL_FILE=$(ls /tmp/arrow-wheels/*.whl) && \
|
||||
VISION_WHL_FILE=$(ls /tmp/vision-wheels/*.whl) && \
|
||||
HF_XET_WHL_FILE=$(ls /tmp/hf-xet-wheels/*.whl) && \
|
||||
LLVM_WHL_FILE=$(ls /tmp/llvmlite-wheels/*.whl) && \
|
||||
NUMBA_WHL_FILE=$(ls /tmp/numba-wheels/*.whl) && \
|
||||
OPENCV_WHL_FILE=$(ls /tmp/opencv-wheels/*.whl) && \
|
||||
GUIDANCE_WHL_FILE=$(ls /tmp/guidance-wheels/*.whl) && \
|
||||
uv pip install -v \
|
||||
$ARROW_WHL_FILE \
|
||||
$VISION_WHL_FILE \
|
||||
@@ -256,7 +164,6 @@ RUN --mount=type=cache,target=/root/.cache/uv \
|
||||
$LLVM_WHL_FILE \
|
||||
$NUMBA_WHL_FILE \
|
||||
$OPENCV_WHL_FILE \
|
||||
$GUIDANCE_WHL_FILE \
|
||||
--torch-backend cpu \
|
||||
--index-strategy unsafe-best-match \
|
||||
-r requirements/build/cpu.txt \
|
||||
|
||||
+45
-2
@@ -141,6 +141,10 @@ variable "ROCM_CSRC_CACHE_TO_MODE" {
|
||||
default = "max"
|
||||
}
|
||||
|
||||
variable "ROCM_RUST_CACHE_TO_MODE" {
|
||||
default = "max"
|
||||
}
|
||||
|
||||
variable "ROCM_FINAL_CACHE_TO_MODE" {
|
||||
default = "min"
|
||||
}
|
||||
@@ -164,6 +168,13 @@ function "get_cache_from_rocm" {
|
||||
VLLM_MERGE_BASE_COMMIT != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:csrc-rocm-${VLLM_MERGE_BASE_COMMIT}" : "",
|
||||
ROCM_CACHE_BRANCH_TAG != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:csrc-rocm-branch-${ROCM_CACHE_BRANCH_TAG}" : "",
|
||||
ROCM_CACHE_UPSTREAM_BRANCH_TAG != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:csrc-rocm-branch-${ROCM_CACHE_UPSTREAM_BRANCH_TAG}" : "",
|
||||
# Import the source-scoped Rust frontend cache so non-Rust changes do not
|
||||
# force a fresh cargo release build.
|
||||
BUILDKITE_COMMIT != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rust-rocm-${BUILDKITE_COMMIT}" : "",
|
||||
PARENT_COMMIT != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rust-rocm-${PARENT_COMMIT}" : "",
|
||||
VLLM_MERGE_BASE_COMMIT != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rust-rocm-${VLLM_MERGE_BASE_COMMIT}" : "",
|
||||
ROCM_CACHE_BRANCH_TAG != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rust-rocm-branch-${ROCM_CACHE_BRANCH_TAG}" : "",
|
||||
ROCM_CACHE_UPSTREAM_BRANCH_TAG != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rust-rocm-branch-${ROCM_CACHE_UPSTREAM_BRANCH_TAG}" : "",
|
||||
# Branch-scoped full image cache - fallback when parent-commit cache is evicted
|
||||
ROCM_CACHE_BRANCH_TAG != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rocm-branch-${ROCM_CACHE_BRANCH_TAG}" : "",
|
||||
ROCM_CACHE_UPSTREAM_BRANCH_TAG != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rocm-branch-${ROCM_CACHE_UPSTREAM_BRANCH_TAG}" : "",
|
||||
@@ -204,6 +215,27 @@ function "get_cache_to_rocm_csrc" {
|
||||
])
|
||||
}
|
||||
|
||||
function "get_cache_from_rocm_rust" {
|
||||
params = []
|
||||
result = compact([
|
||||
BUILDKITE_COMMIT != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rust-rocm-${BUILDKITE_COMMIT}" : "",
|
||||
PARENT_COMMIT != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rust-rocm-${PARENT_COMMIT}" : "",
|
||||
VLLM_MERGE_BASE_COMMIT != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rust-rocm-${VLLM_MERGE_BASE_COMMIT}" : "",
|
||||
ROCM_CACHE_BRANCH_TAG != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rust-rocm-branch-${ROCM_CACHE_BRANCH_TAG}" : "",
|
||||
ROCM_CACHE_UPSTREAM_BRANCH_TAG != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rust-rocm-branch-${ROCM_CACHE_UPSTREAM_BRANCH_TAG}" : "",
|
||||
])
|
||||
}
|
||||
|
||||
function "get_cache_to_rocm_rust" {
|
||||
params = []
|
||||
result = compact([
|
||||
# Export exact-commit and branch-scoped Rust caches. A content-addressed
|
||||
# cache ref is appended by ci-bake-rocm.sh when that wrapper is used.
|
||||
BUILDKITE_COMMIT != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rust-rocm-${BUILDKITE_COMMIT},mode=${ROCM_RUST_CACHE_TO_MODE}" : "",
|
||||
ROCM_CACHE_BRANCH_TAG != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rust-rocm-branch-${ROCM_CACHE_BRANCH_TAG},mode=${ROCM_RUST_CACHE_TO_MODE}" : "",
|
||||
])
|
||||
}
|
||||
|
||||
# Cache functions for upstream dependency stages (RIXL/UCX, ROCShmem, DeepEP).
|
||||
# These stages are pinned to specific upstream commit hashes, so cache keys use
|
||||
# those hashes rather than the Buildkite commit. This means the cache persists
|
||||
@@ -276,6 +308,17 @@ target "csrc-rocm-ci" {
|
||||
output = ["type=cacheonly"]
|
||||
}
|
||||
|
||||
# Cache-only target for the Rust frontend build stage. Final-image cache
|
||||
# exports use mode=min and do not reliably persist intermediate cargo layers,
|
||||
# so Rust gets its own source-scoped cache target.
|
||||
target "rust-rocm-ci" {
|
||||
inherits = ["_common-rocm", "_ci-rocm"]
|
||||
target = "rust-build"
|
||||
cache-from = get_cache_from_rocm_rust()
|
||||
cache-to = get_cache_to_rocm_rust()
|
||||
output = ["type=cacheonly"]
|
||||
}
|
||||
|
||||
# Keep wheel export on the same CI graph as the test image build so the
|
||||
# shared build_vllm/export_vllm stages resolve identically within one bake
|
||||
# invocation. Without this, export-wheel-rocm uses the plain local target
|
||||
@@ -292,13 +335,13 @@ target "export-wheel-rocm" {
|
||||
# Artifact-only vLLM build. GPU test jobs consume this artifact on top of
|
||||
# ci_base, avoiding a per-commit multi-GB image push/pull.
|
||||
group "test-rocm-ci-with-artifacts" {
|
||||
targets = ["csrc-rocm-ci", "export-wheel-rocm"]
|
||||
targets = ["rust-rocm-ci", "csrc-rocm-ci", "export-wheel-rocm"]
|
||||
}
|
||||
|
||||
# Full test image + wheel export. Kept for fallback/debugging when a pushed
|
||||
# per-commit image is useful.
|
||||
group "test-rocm-ci-with-wheel" {
|
||||
targets = ["csrc-rocm-ci", "test-rocm-ci", "export-wheel-rocm"]
|
||||
targets = ["rust-rocm-ci", "csrc-rocm-ci", "test-rocm-ci", "export-wheel-rocm"]
|
||||
}
|
||||
|
||||
# Image tags for the ci_base build. ci-bake-rocm.sh rewrites CI_BASE_IMAGE_TAG
|
||||
|
||||
@@ -232,7 +232,7 @@ MLA decode backends are selected using the standard
|
||||
| `ROCM_AITER_MLA` | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3`, `fp8_e5m2` | %1 | Any | ❌ | ❌ | ❌ | ❌ | ❌ | Decoder | N/A |
|
||||
| `ROCM_AITER_MLA_SPARSE` | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3` | 1, 64 | Any | ❌ | ❌ | ✅ | ❌ | ❌ | Decoder | N/A |
|
||||
| `ROCM_AITER_TRITON_MLA` | fp16, bf16 | `auto` | Any | Any | ❌ | ❌ | ❌ | ❌ | ❌ | Decoder | N/A |
|
||||
| `TOKENSPEED_MLA` | fp16, bf16 | `fp8`, `fp8_e4m3` | 32, 64 | Any | ❌ | ❌ | ❌ | ❌ | ❌ | Decoder | 10.x |
|
||||
| `TOKENSPEED_MLA` | fp16, bf16 | `fp8`, `fp8_e4m3` | 32, 64 | Any | ❌ | ❌ | ❌ | ❌ | ✅ | Decoder | 10.x |
|
||||
| `TRITON_MLA` | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3` | %16 | Any | ❌ | ❌ | ❌ | ❌ | ✅ | Decoder | Any |
|
||||
| `XPU_MLA_SPARSE` | fp16, bf16 | `auto`, `float16`, `bfloat16` | Any | 576 | ❌ | ❌ | ✅ | ❌ | ❌ | Decoder | Any |
|
||||
|
||||
|
||||
@@ -358,6 +358,7 @@ th {
|
||||
| `DeepseekForCausalLM` | DeepSeek | `deepseek-ai/deepseek-llm-67b-base`, `deepseek-ai/deepseek-llm-7b-chat`, etc. | ✅︎ | ✅︎ |
|
||||
| `DeepseekV2ForCausalLM` | DeepSeek-V2 | `deepseek-ai/DeepSeek-V2`, `deepseek-ai/DeepSeek-V2-Chat`, etc. | ✅︎ | ✅︎ |
|
||||
| `DeepseekV3ForCausalLM` | DeepSeek-V3 | `deepseek-ai/DeepSeek-V3`, `deepseek-ai/DeepSeek-R1`, `deepseek-ai/DeepSeek-V3.1`, etc. | ✅︎ | ✅︎ |
|
||||
| `DeepseekV32ForCausalLM` | DeepSeek-V3.2 | `deepseek-ai/DeepSeek-V3.2`, etc. | ✅︎ | ✅︎ |
|
||||
| `DeepseekV4ForCausalLM` | DeepSeek-V4 | `deepseek-ai/DeepSeek-V4-Flash`, `deepseek-ai/DeepSeek-V4-Pro`, etc. | | ✅︎ |
|
||||
| `DotsOCRForCausalLM` | dots_ocr | `rednote-hilab/dots.ocr` | ✅︎ | ✅︎ |
|
||||
| `Ernie4_5ForCausalLM` | Ernie4.5 | `baidu/ERNIE-4.5-0.3B-PT`, etc. | ✅︎ | ✅︎ |
|
||||
|
||||
@@ -21,7 +21,7 @@ pillow # Required for image processing
|
||||
prometheus-fastapi-instrumentator >= 8.0.0 # v8 unblocks starlette >= 1.0
|
||||
tiktoken >= 0.6.0 # Required for DBRX tokenizer
|
||||
lm-format-enforcer == 0.11.3
|
||||
llguidance >= 1.7.0, < 1.8.0; platform_machine == "x86_64" or platform_machine == "arm64" or platform_machine == "aarch64" or platform_machine == "ppc64le"
|
||||
llguidance >= 1.7.0, < 1.8.0; platform_machine == "x86_64" or platform_machine == "arm64" or platform_machine == "aarch64" or platform_machine == "ppc64le" or platform_machine == "s390x"
|
||||
outlines_core == 0.2.14
|
||||
# required for outlines backend disk cache
|
||||
diskcache == 5.6.3
|
||||
|
||||
@@ -26,7 +26,7 @@ nvidia-cutlass-dsl[cu13]==4.5.2
|
||||
quack-kernels>=0.3.3
|
||||
|
||||
# Tokenspeed_MLA for faster mla with spec decode
|
||||
tokenspeed-mla==0.1.2; platform_system == "Linux"
|
||||
tokenspeed-mla==0.1.8; platform_system == "Linux"
|
||||
|
||||
# Humming kernels for quantization gemm
|
||||
humming-kernels[cu13]==0.1.10
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
syntax = "proto3";
|
||||
package vllm;
|
||||
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::time::Duration;
|
||||
|
||||
use anyhow::{Context, Result, bail};
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::sync::Arc;
|
||||
|
||||
use tracing::info;
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::collections::HashMap;
|
||||
use std::sync::Arc;
|
||||
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use thiserror::Error;
|
||||
use thiserror_ext::{AsReport as _, Macro};
|
||||
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::ops::Deref;
|
||||
use std::sync::Arc;
|
||||
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
//! Minimal chat facade above [`vllm_text`].
|
||||
//!
|
||||
//! This crate keeps the northbound boundary intentionally small:
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
//! Chat-layer multimodal media preparation.
|
||||
//!
|
||||
//! This module owns the multimodal path for chat requests: it extracts media
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
//! Prompt placeholder expansion shared across modalities.
|
||||
|
||||
use std::collections::{HashMap, VecDeque};
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
//! Image-modality preparation: batch preprocessing and per-item feature
|
||||
//! build.
|
||||
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::collections::HashMap;
|
||||
|
||||
use half::{bf16, f16};
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
//! Video-modality preparation: per-clip preprocessing, config resolution,
|
||||
//! and per-item feature build.
|
||||
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
//! Default output processing pipeline.
|
||||
|
||||
mod structural_tag;
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
//! Applies xgrammar structural-tag constraints for strict tool calling.
|
||||
|
||||
use thiserror_ext::AsReport;
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
//! Adapts decoded text updates into parsed assistant deltas.
|
||||
//!
|
||||
//! This stage sits between low-level token decoding and final block assembly.
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
//! Native Harmony output processing for `gpt_oss`.
|
||||
//!
|
||||
//! Unlike the default text-first pipeline, this processor consumes
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::sync::Arc;
|
||||
|
||||
use futures::executor::block_on;
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::pin::Pin;
|
||||
use std::sync::Arc;
|
||||
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
//! Adapts parsed assistant updates into structured chat events.
|
||||
//!
|
||||
//! This module remains the final assembly stage in `vllm-chat`. Token-to-text
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
pub mod reasoning;
|
||||
pub mod tool;
|
||||
pub mod unified;
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
//! Reasoning parser registration and selection boundary for `vllm-chat`.
|
||||
|
||||
use std::sync::{Arc, LazyLock};
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::sync::Arc;
|
||||
|
||||
use vllm_tokenizer::test_utils::TestTokenizer;
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
//! Tool parser registration and selection boundary for `vllm-chat`.
|
||||
|
||||
use std::sync::{Arc, LazyLock};
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use vllm_parser::tool::{Result, ToolParserOutput};
|
||||
|
||||
use super::{ToolParser, ToolParserFactory, names};
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
//! Unified parser registration and selection boundary for `vllm-chat`.
|
||||
|
||||
use std::sync::LazyLock;
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
//! DeepSeek V3.2 prompt renderer.
|
||||
|
||||
use std::collections::{HashMap, HashSet};
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
mod encoding;
|
||||
|
||||
use vllm_text::Prompt;
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::path::PathBuf;
|
||||
|
||||
use expect_test::{ExpectFile, expect, expect_file};
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
//! DeepSeek V4 prompt renderer.
|
||||
//!
|
||||
//! Original Python implementation:
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
mod encoding;
|
||||
|
||||
use vllm_text::Prompt;
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::path::PathBuf;
|
||||
|
||||
use expect_test::{ExpectFile, expect, expect_file};
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
//! Shared Harmony encoding helper for the GPT-OSS renderer and output parser.
|
||||
|
||||
use std::sync::LazyLock;
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
//! Native Harmony chat renderer for `gpt_oss`.
|
||||
|
||||
pub(crate) mod encoding;
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::path::PathBuf;
|
||||
|
||||
use expect_test::{ExpectFile, expect, expect_file};
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use thiserror::Error as ThisError;
|
||||
|
||||
#[derive(Debug, ThisError)]
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::collections::{HashSet, VecDeque};
|
||||
use std::fmt;
|
||||
use std::str::FromStr;
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::collections::HashMap;
|
||||
|
||||
use serde::Serialize;
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
//! Chat template support for tokenizers using Jinja2 templates.
|
||||
//!
|
||||
//! This module is inlined from SMG's tokenizer crate with local adaptations:
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use minijinja::value::{Kwargs, ViaDeserialize};
|
||||
use minijinja::{Error as MinijinjaError, ErrorKind, Value};
|
||||
use serde::Deserialize;
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::sync::Arc;
|
||||
|
||||
use indexmap::IndexMap;
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::collections::HashMap;
|
||||
use std::sync::Arc;
|
||||
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::fmt;
|
||||
use std::str::FromStr;
|
||||
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::fs;
|
||||
use std::path::Path;
|
||||
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::collections::HashMap;
|
||||
|
||||
use llm_multimodal::ImageDetail;
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::pin::Pin;
|
||||
use std::sync::Arc;
|
||||
use std::task::{Context, Poll};
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::collections::BTreeSet;
|
||||
use std::fmt;
|
||||
use std::sync::Arc;
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
//! Roundtrip tests for the real chat-template and output-processor pairing.
|
||||
//!
|
||||
//! The invariant under test is that a structured assistant message rendered as history can be
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
//! CLI argument definitions for the `vllm-rs` binary.
|
||||
//!
|
||||
//! Python vLLM references:
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use expect_test::expect;
|
||||
use vllm_engine_core_client::TransportMode;
|
||||
use vllm_server::{Config, HttpListenerMode, ParserSelection, RendererSelection};
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
#![allow(clippy::doc_lazy_continuation)]
|
||||
|
||||
use std::fmt::Display;
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::{env, fmt, process};
|
||||
|
||||
use time::UtcOffset;
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
mod cli;
|
||||
mod logging;
|
||||
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::time::Duration;
|
||||
|
||||
use anyhow::{Context, Result, bail};
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::time::Duration;
|
||||
|
||||
use anyhow::{Context, Result, bail};
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::sync::Arc;
|
||||
use std::time::Duration;
|
||||
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::collections::{BTreeMap, BTreeSet};
|
||||
use std::sync::Arc;
|
||||
use std::time::{SystemTime, UNIX_EPOCH};
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::collections::{BTreeMap, BTreeSet, HashMap};
|
||||
use std::sync::atomic::{AtomicU64, Ordering};
|
||||
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::ops::Deref;
|
||||
use std::pin::Pin;
|
||||
use std::task::{Context, Poll};
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::time::Duration;
|
||||
|
||||
use bytes::Bytes;
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::sync::Arc;
|
||||
|
||||
use serde_tuple::{Deserialize_tuple, Serialize_tuple};
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::sync::Arc;
|
||||
|
||||
use parking_lot::Mutex;
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::sync::Arc;
|
||||
|
||||
use serde_tuple::{Deserialize_tuple, Serialize_tuple};
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
mod bootstrap;
|
||||
mod external;
|
||||
mod handle;
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::sync::Arc;
|
||||
use std::time::Duration;
|
||||
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
mod client;
|
||||
mod coordinator;
|
||||
mod error;
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::collections::BTreeMap;
|
||||
use std::collections::BTreeSet;
|
||||
use std::time::{SystemTime, UNIX_EPOCH};
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::path::Path;
|
||||
use std::time::Duration;
|
||||
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use serde::{Deserialize, Serialize};
|
||||
|
||||
/// Effective model dtype reported by the engine after config resolution.
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::collections::BTreeMap;
|
||||
|
||||
use serde::{Deserialize, Serialize};
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
mod array;
|
||||
#[cfg(test)]
|
||||
mod tests;
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::io::Cursor;
|
||||
|
||||
use byteorder::{BigEndian, LittleEndian, NativeEndian, ReadBytesExt};
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::collections::BTreeSet;
|
||||
|
||||
use bytes::Bytes;
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use serde_tuple::{Deserialize_tuple, Serialize_tuple};
|
||||
|
||||
use crate::protocol::tensor::WireNdArray;
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use serde_tuple::{Deserialize_tuple, Serialize_tuple};
|
||||
|
||||
use crate::protocol::OpaqueValue;
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::any::type_name;
|
||||
use std::io::Cursor;
|
||||
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::collections::BTreeMap;
|
||||
|
||||
use serde::{Deserialize, Serialize};
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::collections::BTreeSet;
|
||||
|
||||
use enum_as_inner::EnumAsInner;
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::collections::{BTreeMap, HashMap};
|
||||
|
||||
use bytes::Bytes;
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::collections::{BTreeSet, HashMap};
|
||||
|
||||
use serde::{Deserialize, Serialize};
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use std::collections::BTreeMap;
|
||||
|
||||
use serde::{Deserialize, Serialize};
|
||||
|
||||
@@ -1,3 +1,6 @@
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
|
||||
|
||||
use enum_as_inner::EnumAsInner;
|
||||
use serde::{Deserialize, Deserializer, Serialize, Serializer};
|
||||
use serde_json::Value;
|
||||
|
||||
Some files were not shown because too many files have changed in this diff Show More
Reference in New Issue
Block a user