diff --git a/cmake/external_projects/vllm_flash_attn.cmake b/cmake/external_projects/vllm_flash_attn.cmake index 97a8cfe87b..28a6b336e2 100644 --- a/cmake/external_projects/vllm_flash_attn.cmake +++ b/cmake/external_projects/vllm_flash_attn.cmake @@ -39,7 +39,7 @@ else() FetchContent_Declare( vllm-flash-attn GIT_REPOSITORY https://github.com/vllm-project/flash-attention.git - GIT_TAG bb9a72e7dde0dc614ffc663e052cd6a19ce73a42 + GIT_TAG caaa4eb59845388a20b1f435ecaafb4bd9517ad8 GIT_PROGRESS TRUE # Don't share the vllm-flash-attn build between build types BINARY_DIR ${CMAKE_BINARY_DIR}/vllm-flash-attn diff --git a/docker/Dockerfile b/docker/Dockerfile index 1263bff436..f008fd29e1 100644 --- a/docker/Dockerfile +++ b/docker/Dockerfile @@ -248,10 +248,6 @@ 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; \ @@ -268,13 +264,6 @@ 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. @@ -792,10 +781,6 @@ 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; \ @@ -803,13 +788,6 @@ 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) @@ -908,19 +886,6 @@ 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 - # CUDA image changed from /usr/local/nvidia to /usr/local/cuda in 12.8 but will # return to /usr/local/nvidia in 13.0 to allow container providers to mount drivers # consistently from the host (see https://github.com/vllm-project/vllm/issues/18859). diff --git a/requirements/cuda.txt b/requirements/cuda.txt index 766acb0204..3b8ff816f1 100644 --- a/requirements/cuda.txt +++ b/requirements/cuda.txt @@ -16,7 +16,7 @@ PyNvVideoCodec==2.0.4 --extra-index-url https://flashinfer.ai/whl/ flashinfer-python==0.6.14 flashinfer-cubin==0.6.14 -apache-tvm-ffi==0.1.9 +apache-tvm-ffi==0.1.10 tilelang==0.1.9 nvidia-cudnn-frontend>=1.19.1 # Required for LLM_NVTX_SCOPES_FOR_PROFILING=1 @@ -25,7 +25,7 @@ nvtx==0.2.15 fastsafetensors >= 0.3.2 # QuACK and Cutlass DSL for FA4 (cute-DSL implementation) -nvidia-cutlass-dsl[cu13]==4.5.2 +nvidia-cutlass-dsl[cu13]==4.6.0 quack-kernels>=0.4.0 # Required for tml-fa4 # Tokenspeed_MLA for faster mla with spec decode diff --git a/requirements/test/cuda.txt b/requirements/test/cuda.txt index 7fa6378743..e4da9ac06e 100644 --- a/requirements/test/cuda.txt +++ b/requirements/test/cuda.txt @@ -43,7 +43,7 @@ anyio==4.14.1 # sse-starlette # starlette # watchfiles -apache-tvm-ffi==0.1.9 +apache-tvm-ffi==0.1.10 # via # -c requirements/cuda.txt # xgrammar diff --git a/tests/entrypoints/pooling/scoring/test_cross_encoder_online_vision.py b/tests/entrypoints/pooling/scoring/test_cross_encoder_online_vision.py index c6663dbdff..8c1d75cd76 100644 --- a/tests/entrypoints/pooling/scoring/test_cross_encoder_online_vision.py +++ b/tests/entrypoints/pooling/scoring/test_cross_encoder_online_vision.py @@ -47,6 +47,7 @@ BACKEND_TOL: dict[str, float] = { # TRITON_ATTN: gfx942/ROCm 7.2 drifts ~0.008 abs on text-vs-text (~7.9% rel). BACKEND_ABS_TOL: dict[str, float] = { "default": 0.0, + "auto": 0.007, "ROCM_AITER_FA": 0.005, "TRITON_ATTN": 0.009, "FLEX_ATTENTION": 0.006, diff --git a/vllm/model_executor/layers/mamba/gdn/qwen_gdn_linear_attn.py b/vllm/model_executor/layers/mamba/gdn/qwen_gdn_linear_attn.py index 9e286c692d..2f08d9a67b 100644 --- a/vllm/model_executor/layers/mamba/gdn/qwen_gdn_linear_attn.py +++ b/vllm/model_executor/layers/mamba/gdn/qwen_gdn_linear_attn.py @@ -2,7 +2,6 @@ # SPDX-FileCopyrightText: Copyright contributors to the vLLM project """Inference-only Qwen3-Next/Qwen3.5 model.""" -import functools from typing import Literal import torch @@ -83,70 +82,6 @@ if GDN_AITER_TRITON_AVAILABLE: logger = init_logger(__name__) -# TODO(arpera): remove ``_is_libs_cu13_install_intact`` and its caller in -# ``_resolve_gdn_prefill_backend`` once the upstream packaging bug is -# fixed and the broken wheels are yanked / superseded on PyPI: -# https://github.com/NVIDIA/cutlass/issues/3170 -# https://github.com/NVIDIA/cutlass/issues/3259 -@functools.cache -def _is_libs_cu13_install_intact() -> bool: - """Return True if every file installed by ``nvidia-cutlass-dsl-libs-cu13`` - matches the SHA-256 declared in its wheel ``RECORD``. - - ``nvidia-cutlass-dsl-libs-base`` and ``nvidia-cutlass-dsl-libs-cu13`` - both ship into the shared ``nvidia_cutlass_dsl/`` namespace and - write many of the same on-disk paths (the runtime ``.so``, the MLIR - Python bindings, cuTe-DSL Python sources, ...) with different - content. Whichever wheel extracts last wins; with a parallel - installer (e.g. ``uv``) the order is racy and the resulting venv - can end up with a mix of files from both variants. The - ``-libs-base`` variant fails MLIR legalization when JIT-compiling - the FlashInfer Blackwell GDN prefill kernel, and any other - cuTe-DSL-based kernel can break too if on-disk files diverge from - what ``-libs-cu13``'s wheel expects. Tracked upstream at: - - * https://github.com/NVIDIA/cutlass/issues/3170 - * https://github.com/NVIDIA/cutlass/issues/3259 - - This helper re-hashes every file the ``-libs-cu13`` wheel claims to - own and compares against its declared SHA-256. Returns False on any - error (uninstalled, missing RECORD, missing file, hash mismatch). - Result is cached per-process. - """ - import hashlib - import importlib.metadata - - import pybase64 as base64 - - try: - dist = importlib.metadata.distribution("nvidia-cutlass-dsl-libs-cu13") - except importlib.metadata.PackageNotFoundError: - return False - - files = dist.files - if not files: - return False - - for pkg_path in files: - file_hash = pkg_path.hash - # Skip RECORD rows without a hash (RECORD itself, generated - # ``.pyc`` files, ...) and any non-SHA-256 hash modes. - if file_hash is None or not file_hash.value: - continue - if file_hash.mode != "sha256": - continue - try: - with open(pkg_path.locate(), "rb") as f: - digest = hashlib.sha256(f.read()).digest() - except OSError: - return False - actual = base64.urlsafe_b64encode(digest).decode().rstrip("=") - if actual != file_hash.value: - return False - - return True - - def _resolve_gdn_prefill_backend( vllm_config: VllmConfig, ) -> tuple[str, Literal["triton", "flashinfer", "cutedsl"]]: @@ -157,9 +92,7 @@ def _resolve_gdn_prefill_backend( * ``platform == cuda``; * one of the following: - Hopper (SM90) — no further constraints; - - Blackwell (SM10.x) with ``head_k_dim == 128``, ``cuda_runtime >= 13``, - and an intact ``nvidia-cutlass-dsl-libs-cu13`` install on disk - (see :func:`_is_libs_cu13_install_intact`). + - Blackwell (SM10.x) with ``head_k_dim == 128``, ``cuda_runtime >= 13``. In-tree CuteDSL GDN prefill kernel is chosen when: * "cutedsl" is requested; (opt-in only) @@ -190,19 +123,8 @@ def _resolve_gdn_prefill_backend( and head_k_dim == 128 and current_platform.get_cuda_runtime_major() >= 13 ): - supports_flashinfer = _is_libs_cu13_install_intact() + supports_flashinfer = True supports_cutedsl = True - if not supports_flashinfer: - logger.warning_once( - "FlashInfer Blackwell GDN requires an intact nvidia-cutlass-dsl" - "-libs-cu13 install, but some on-disk files do not match the " - "SHA-256 declared in its RECORD (install-order race in " - "nvidia-cutlass-dsl packaging -- see " - "https://github.com/NVIDIA/cutlass/issues/3170 and " - "https://github.com/NVIDIA/cutlass/issues/3259). Falling back " - "to Triton/FLA. Repair with: pip install --force-reinstall " - "--no-deps nvidia-cutlass-dsl-libs-cu13" - ) if backend in ["flashinfer", "auto"] and supports_flashinfer: return backend, "flashinfer" diff --git a/vllm/vllm_flash_attn/flash_attn_interface.py b/vllm/vllm_flash_attn/flash_attn_interface.py index bef811f6e6..fc4948d0e9 100644 --- a/vllm/vllm_flash_attn/flash_attn_interface.py +++ b/vllm/vllm_flash_attn/flash_attn_interface.py @@ -388,7 +388,7 @@ def flash_attn_varlen_func( from vllm.vllm_flash_attn.cute.interface import _flash_attn_fwd - out, softmax_lse = _flash_attn_fwd( + out, softmax_lse, _, _ = _flash_attn_fwd( q, k, v,