Compare commits

...
Author SHA1 Message Date
avesedandkhluu 0fc695fc6d [Bugfix][Frontend] Cap fastapi < 0.137 to avoid prometheus-fastapi-instrumentator crash on serve startup
A fresh pip install pulls fastapi 0.137, whose new _IncludedRouter route type
(a BaseRoute with no .path) breaks prometheus-fastapi-instrumentator's route
iteration -> 'vllm serve' dies at startup with:
  AttributeError: '_IncludedRouter' object has no attribute 'path'

requirements/common.txt pinned fastapi[standard]>=0.115.0 with no upper bound.
Cap to <0.137; lift once the instrumentator handles _IncludedRouter
(https://github.com/trallnag/prometheus-fastapi-instrumentator/issues/370).

Signed-off-by: avesed <trevor.lijunlin@gmail.com>
2026-06-14 20:35:17 -07:00
Kevin H. Luu 91df0fad4d [Bugfix][CPU] Don't build triton-cpu on arm64 release image (#45401)
Signed-off-by: khluu <khluu000@gmail.com>
(cherry picked from commit e3e31e54b0)
2026-06-12 14:52:17 -07:00
Mohammad Miadh Angkadandkhluu 78743ab5bf [Docker] Fix CUTLASS DSL cu13 install order in Dockerfile (#45204)
Signed-off-by: Mohammad Miadh Angkad <176301910+mmangkad@users.noreply.github.com>
(cherry picked from commit 40e065e86a)
2026-06-11 16:24:31 -07:00
Rohan Potdarandkhluu b2d7294b0f [ROCm][Bugfix] Make intermediate_pad TP-aware in rocm_aiter_fused_experts (#44679)
Signed-off-by: Rohan138 <rohanpotdar138@gmail.com>
Co-authored-by: Andreas Karatzas <akaratza@amd.com>
(cherry picked from commit 5b6b536fdc)
2026-06-11 16:24:26 -07:00
TJianandkhluu 741ba421d8 [Bugfix] [DSV4] [ROCm] Pin apache-tvm-ffi version to 0.1.10 (#45169)
Signed-off-by: tjtanaa <tunjian.tan@embeddedllm.com>
(cherry picked from commit bfe1001ab6)
2026-06-11 16:24:21 -07:00
xaguilar-amdandkhluu ac94893da3 [ROCm][MLA][Bugfix] Reserve FP8 prefill workspace before lock for Kimi-K2.5 (#42978)
Signed-off-by: Xavier Aguilar <xavier.aguilarfruto@amd.com>
Co-authored-by: Douglas Lehr <91553416+dllehr-amd@users.noreply.github.com>
Co-authored-by: Andreas Karatzas <akaratza@amd.com>
(cherry picked from commit baacbfcebf)
2026-06-11 16:24:16 -07:00
7 changed files with 96 additions and 9 deletions
+35
View File
@@ -218,6 +218,10 @@ COPY requirements/common.txt requirements/common.txt
COPY requirements/cuda.txt requirements/cuda.txt
COPY use_existing_torch.py use_existing_torch.py
COPY pyproject.toml pyproject.toml
# nvidia-cutlass-dsl[cu13] installs -libs-base and -libs-cu13 wheels that
# share paths with different content. uv can extract them in either order,
# leaving base files that break CUDA 13 CuTe DSL JIT.
# TODO(mmangkad): Remove this after NVIDIA/cutlass#3259 is fixed.
RUN --mount=type=cache,target=/opt/uv/cache \
if [ "$(echo $CUDA_VERSION | cut -d. -f1)" = "12" ]; then \
sed -i 's/^nvidia-cutlass-dsl\[cu13\]/nvidia-cutlass-dsl/' requirements/cuda.txt; \
@@ -234,6 +238,13 @@ RUN --mount=type=cache,target=/opt/uv/cache \
else \
uv pip install --python /opt/venv/bin/python3 -r requirements/cuda.txt \
--extra-index-url ${PYTORCH_CUDA_INDEX_BASE_URL}/cu$(echo $CUDA_VERSION | cut -d. -f1,2 | tr -d '.'); \
fi \
&& if [ "$(echo $CUDA_VERSION | cut -d. -f1)" = "13" ]; then \
CUTLASS_DSL_VERSION=$(uv pip show --python /opt/venv/bin/python3 nvidia-cutlass-dsl 2>/dev/null | awk '/^Version:/{print $2}') && \
if [ -n "$CUTLASS_DSL_VERSION" ]; then \
uv pip install --python /opt/venv/bin/python3 --force-reinstall --no-deps \
"nvidia-cutlass-dsl-libs-cu13==${CUTLASS_DSL_VERSION}"; \
fi; \
fi
# Track PyTorch lib versions used during build and match in downstream instances.
@@ -745,6 +756,10 @@ ENV VLLM_ENABLE_CUDA_COMPATIBILITY=0
ARG PYTORCH_CUDA_INDEX_BASE_URL
COPY requirements/common.txt /tmp/common.txt
COPY requirements/cuda.txt /tmp/requirements-cuda.txt
# nvidia-cutlass-dsl[cu13] installs -libs-base and -libs-cu13 wheels that
# share paths with different content. uv can extract them in either order,
# leaving base files that break CUDA 13 CuTe DSL JIT.
# TODO(mmangkad): Remove this after NVIDIA/cutlass#3259 is fixed.
RUN --mount=type=cache,target=/opt/uv/cache \
if [ "$(echo $CUDA_VERSION | cut -d. -f1)" = "12" ]; then \
sed -i 's/^nvidia-cutlass-dsl\[cu13\]/nvidia-cutlass-dsl/' /tmp/requirements-cuda.txt; \
@@ -752,6 +767,13 @@ RUN --mount=type=cache,target=/opt/uv/cache \
fi && \
uv pip install --system -r /tmp/requirements-cuda.txt \
--extra-index-url ${PYTORCH_CUDA_INDEX_BASE_URL}/cu$(echo $CUDA_VERSION | cut -d. -f1,2 | tr -d '.') && \
if [ "$(echo $CUDA_VERSION | cut -d. -f1)" = "13" ]; then \
CUTLASS_DSL_VERSION=$(uv pip show --system nvidia-cutlass-dsl 2>/dev/null | awk '/^Version:/{print $2}') && \
if [ -n "$CUTLASS_DSL_VERSION" ]; then \
uv pip install --system --force-reinstall --no-deps \
"nvidia-cutlass-dsl-libs-cu13==${CUTLASS_DSL_VERSION}"; \
fi; \
fi && \
rm /tmp/requirements-cuda.txt /tmp/common.txt
# Install FlashInfer JIT cache (requires CUDA-version-specific index URL)
@@ -842,6 +864,19 @@ RUN --mount=type=bind,from=build,src=/tmp/ep_kernels_workspace/dist,target=/vllm
uv pip install --system ep_kernels/dist/*.whl --verbose \
--extra-index-url ${PYTORCH_CUDA_INDEX_BASE_URL}/cu$(echo $CUDA_VERSION | cut -d. -f1,2 | tr -d '.')
# nvidia-cutlass-dsl[cu13] installs -libs-base and -libs-cu13 wheels that
# share paths with different content. Force -libs-cu13 last after runtime
# dependency installs so uv cannot leave base files behind.
# TODO(mmangkad): Remove this after NVIDIA/cutlass#3259 is fixed.
RUN --mount=type=cache,target=/opt/uv/cache \
if [ "$(echo $CUDA_VERSION | cut -d. -f1)" = "13" ]; then \
CUTLASS_DSL_VERSION=$(uv pip show --system nvidia-cutlass-dsl 2>/dev/null | awk '/^Version:/{print $2}') && \
if [ -n "$CUTLASS_DSL_VERSION" ]; then \
uv pip install --system --force-reinstall --no-deps \
"nvidia-cutlass-dsl-libs-cu13==${CUTLASS_DSL_VERSION}"; \
fi; \
fi
# Download FlashInfer precompiled cubins AFTER all pip installs are done.
# This must run after the vLLM wheel and EP kernels installs above, because
# those can reinstall/touch flashinfer packages. Downloading cubins earlier
+11
View File
@@ -168,6 +168,12 @@ RUN --mount=type=cache,target=/root/.cache/uv \
######################### TRITON-CPU BUILD IMAGE #########################
FROM base AS vllm-triton-cpu-build
# Support for cross-compilation with x86 ISA including AVX2 and AVX512: docker build --build-arg VLLM_CPU_X86="true" ...
# Re-declared here because this stage is `FROM base` (not `vllm-build`), so it
# does not inherit the ARG/ENV defined there. Without it, the guard below would
# see an empty value and build triton-cpu on non-x86 targets (e.g. arm64).
ARG VLLM_CPU_X86=0
WORKDIR /vllm-workspace
RUN mkdir dist
@@ -269,6 +275,11 @@ ENV HF_HUB_DOWNLOAD_TIMEOUT 60
######################### RELEASE IMAGE #########################
FROM base AS vllm-openai
# Re-declared here because this stage is `FROM base` (not `vllm-build`), so the
# RUN below that gates the triton-cpu wheel install on $VLLM_CPU_X86 would
# otherwise see an empty value and try to install it on non-x86 targets.
ARG VLLM_CPU_X86=0
WORKDIR /vllm-workspace
RUN --mount=type=cache,target=/root/.cache/uv \
+1 -1
View File
@@ -11,7 +11,7 @@ transformers >= 4.56.0, != 5.0.*, != 5.1.*, != 5.2.*, != 5.3.*, != 5.4.*, != 5.5
tokenizers >= 0.21.1 # Required for fast incremental detokenization.
safetensors >= 0.6.2 # MXFP4/MXFP6 dtype support (F8_E8M0, F4) added in 0.6.0: https://github.com/huggingface/safetensors/pull/611
protobuf >= 5.29.6, !=6.30.*, !=6.31.*, !=6.32.*, !=6.33.0.*, !=6.33.1.*, !=6.33.2.*, !=6.33.3.*, !=6.33.4.* # Required by LlamaTokenizer, gRPC. CVE-2026-0994
fastapi[standard] >= 0.115.0 # Required by FastAPI's form models in the OpenAI API server's audio transcriptions endpoint.
fastapi[standard] >= 0.115.0, < 0.137 # Required by FastAPI's form models in the OpenAI API server's audio transcriptions endpoint; <0.137 due to prometheus-fastapi-instrumentator#370.
aiohttp >= 3.13.3
openai >= 2.0.0 # For Responses API with reasoning content
pydantic >= 2.12.0
+2 -1
View File
@@ -23,6 +23,7 @@ timm>=1.0.17
# To be consistent with test_quark.py
amd-quark>=0.8.99
tilelang==0.1.10
# Required apache-tvm-ffi matching tilelang version
apache-tvm-ffi==0.1.10
# Required for faster safetensors model loading
fastsafetensors >= 0.3.2
+1
View File
@@ -44,6 +44,7 @@ anyio==4.13.0
# watchfiles
apache-tvm-ffi==0.1.10
# via
# -c requirements/rocm.txt
# tilelang
# xgrammar
arctic-inference==0.1.1
@@ -340,6 +340,16 @@ def rocm_aiter_fused_experts(
moe_config.intermediate_size_per_partition
- moe_config.intermediate_size_per_partition_unpadded
)
# Round hidden_pad/intermediate_pad to match AITER's CK/FlyDSL MoE
# dispatch (currently pinned to v0.1.13.post1):
# https://github.com/ROCm/aiter/blob/v0.1.13.post1/aiter/fused_moe.py#L1073
# https://github.com/ROCm/aiter/blob/v0.1.13.post1/aiter/fused_moe.py#L1099
# TODO: Revisit this once we bump AITER to 0.1.15 with padding fixes
# for CK/FlyDSL MoE GEMM e.g. https://github.com/ROCm/aiter/pull/3401
hidden_pad = hidden_pad // 128 * 128
intermediate_pad = (
intermediate_pad // 64 * 64 * (2 if moe_config.tp_size == 1 else 1)
)
return rocm_aiter_ops.fused_moe(
hidden_states,
@@ -357,8 +367,8 @@ def rocm_aiter_fused_experts(
doweight_stage1=apply_router_weight_on_input,
num_local_tokens=num_local_tokens,
output_dtype=output_dtype,
hidden_pad=hidden_pad // 128 * 128,
intermediate_pad=intermediate_pad // 64 * 64 * 2,
hidden_pad=hidden_pad,
intermediate_pad=intermediate_pad,
bias1=quant_config.w1_bias if quant_config.use_mxfp4_w4a16 else None,
bias2=quant_config.w2_bias if quant_config.use_mxfp4_w4a16 else None,
moe_sorting_dispatch_policy=moe_sorting_dispatch_policy,
@@ -242,7 +242,12 @@ class AiterMLAMetadataBuilder(MLACommonMetadataBuilder[AiterMLAMetadata]):
vllm_config.model_config.max_model_len,
vllm_config.scheduler_config.max_num_batched_tokens,
)
self._init_fp8_prefill_ps_buffers(max_num_reqs, max_prefill_qlen, device)
self._init_fp8_prefill_ps_buffers(
max_num_reqs,
max_prefill_qlen,
vllm_config.scheduler_config.max_num_batched_tokens,
device,
)
if self.compilation_config.cudagraph_mode.has_full_cudagraphs():
self.paged_kv_indptr = torch.zeros(
@@ -257,21 +262,29 @@ class AiterMLAMetadataBuilder(MLACommonMetadataBuilder[AiterMLAMetadata]):
self,
max_num_reqs: int,
max_prefill_qlen: int,
max_num_batched_tokens: int,
device: torch.device,
) -> None:
"""Pre-allocate persistent buffers for FP8 MLA prefill PS metadata.
Uses ``get_ps_metadata_info_v1`` with max values so the buffers are
large enough for any batch. ``get_ps_metadata_v1`` fills them
per-batch in ``build()``.
per-batch in ``build()``. The FP8 prefill forward path also uses the
global workspace manager for per-call scratch, so reserve its maximum
shape here before the workspace manager is locked after warmup.
Args:
max_num_reqs: Maximum number of concurrent requests.
max_prefill_qlen: Maximum Q-length for a single request in one
prefill batch. Should be ``min(max_model_len,
max_num_batched_tokens)`` — the chunked-prefill scheduler
never emits more than ``max_num_batched_tokens`` new tokens
per batch.
max_num_batched_tokens)`` — a single request never exceeds
``max_model_len`` tokens, nor the per-batch token budget.
max_num_batched_tokens: Maximum number of tokens scheduled in one
batch. The ``final_lse`` scratch is sized by ``total_q`` (the
summed Q-length over all prefill requests in the batch), which
is bounded by this budget rather than by a single request's
``max_prefill_qlen`` — concurrent requests can sum to more than
``max_model_len`` when ``max_model_len < max_num_batched_tokens``.
device: Target device for the buffers.
"""
from aiter import get_ps_metadata_info_v1
@@ -279,6 +292,7 @@ class AiterMLAMetadataBuilder(MLACommonMetadataBuilder[AiterMLAMetadata]):
# After kv_b_proj decompression, K has num_heads heads (same as Q).
# So gqa_ratio=1 and num_head_k=num_heads for the PS kernel.
num_head_k = self.num_heads
v_head_dim = self.mla_dims.v_head_dim
# gqa_ratio = 1
# qlen_granularity = _FP8_PREFILL_TILE_Q // max(gqa_ratio, 1)
qlen_granularity = _FP8_PREFILL_TILE_Q
@@ -318,6 +332,21 @@ class AiterMLAMetadataBuilder(MLACommonMetadataBuilder[AiterMLAMetadata]):
device=device,
)
from vllm.v1.worker.workspace import current_workspace_manager
max_num_partial_tiles = reduce_partial_map_size
current_workspace_manager().get_simultaneous(
(
(max_num_partial_tiles * _FP8_PREFILL_TILE_Q, num_head_k, v_head_dim),
torch.float32,
),
(
(max_num_partial_tiles * _FP8_PREFILL_TILE_Q, num_head_k),
torch.float32,
),
((max_num_batched_tokens, num_head_k), torch.float32),
)
logger.info(
"FP8 MLA prefill PS buffers allocated "
"(max_batch=%d, max_qlen=%d, num_head_k=%d)",