mirror of
https://github.com/vllm-project/vllm.git
synced 2026-08-15 10:18:10 +00:00
[CI/Build] Bump flashinfer to v0.6.11.post2 (#41711)
Signed-off-by: Artem Perevedentsev <[email protected]> Co-authored-by: Vadim Gimpelson <[email protected]>
This commit is contained in:
co-authored by
Vadim Gimpelson
parent
36e74c9ea4
commit
0867497368
+1
-1
@@ -639,7 +639,7 @@ RUN --mount=type=cache,target=/root/.cache/uv \
|
||||
# Install FlashInfer JIT cache (requires CUDA-version-specific index URL)
|
||||
# https://docs.flashinfer.ai/installation.html
|
||||
# From versions.json: .flashinfer.version
|
||||
ARG FLASHINFER_VERSION=0.6.8.post1
|
||||
ARG FLASHINFER_VERSION=0.6.11.post2
|
||||
RUN --mount=type=cache,target=/root/.cache/uv \
|
||||
uv pip install --system flashinfer-jit-cache==${FLASHINFER_VERSION} \
|
||||
--extra-index-url https://flashinfer.ai/whl/cu$(echo $CUDA_VERSION | cut -d. -f1,2 | tr -d '.')
|
||||
|
||||
@@ -217,13 +217,13 @@ RUN pip install setuptools==75.6.0 packaging==23.2 ninja==1.11.1.3 build==1.2.2.
|
||||
|
||||
|
||||
# build flashinfer for torch nightly from source around 10 mins
|
||||
# release version: v0.6.8.post1
|
||||
# release version: v0.6.11.post2
|
||||
# todo(elainewy): cache flashinfer build result for faster build
|
||||
ENV CCACHE_DIR=/root/.cache/ccache
|
||||
RUN --mount=type=cache,target=/root/.cache/ccache \
|
||||
--mount=type=cache,target=/root/.cache/uv \
|
||||
echo "git clone flashinfer..." \
|
||||
&& git clone --depth 1 --branch v0.6.8.post1 --recursive https://github.com/flashinfer-ai/flashinfer.git \
|
||||
&& git clone --depth 1 --branch v0.6.11.post2 --recursive https://github.com/flashinfer-ai/flashinfer.git \
|
||||
&& cd flashinfer \
|
||||
&& git submodule update --init --recursive \
|
||||
&& echo "finish git clone flashinfer..." \
|
||||
|
||||
@@ -68,7 +68,7 @@
|
||||
"default": "true"
|
||||
},
|
||||
"FLASHINFER_VERSION": {
|
||||
"default": "0.6.8.post1"
|
||||
"default": "0.6.11.post2"
|
||||
},
|
||||
"GDRCOPY_CUDA_VERSION": {
|
||||
"default": "12.8"
|
||||
|
||||
@@ -9,8 +9,8 @@ torchaudio==2.11.0
|
||||
# These must be updated alongside torch
|
||||
torchvision==0.26.0 # Required for phi3v processor. See https://github.com/pytorch/vision?tab=readme-ov-file#installation for corresponding version
|
||||
# FlashInfer should be updated together with the Dockerfile
|
||||
flashinfer-python==0.6.8.post1
|
||||
flashinfer-cubin==0.6.8.post1
|
||||
flashinfer-python==0.6.11.post2
|
||||
flashinfer-cubin==0.6.11.post2
|
||||
apache-tvm-ffi==0.1.9
|
||||
tilelang==0.1.9
|
||||
# Cap nvidia-cudnn-frontend (transitive dep of flashinfer) due to
|
||||
|
||||
@@ -453,11 +453,15 @@ def test_flashinfer_cutedsl_moe_masked(
|
||||
)
|
||||
|
||||
# reference
|
||||
a_fp4, a_scale_interleaved = fp4_quantize(hidden_states, input_global_scale)
|
||||
# input_global_scale is per-expert ([num_experts]); fp4_quantize and
|
||||
# dequantize_nvfp4_to_dtype are non-grouped APIs that expect [1] or
|
||||
# [num_tokens]. Use a single element since all values are uniform here.
|
||||
a_global = input_global_scale[:1].contiguous()
|
||||
a_fp4, a_scale_interleaved = fp4_quantize(hidden_states, a_global)
|
||||
a_in_dtype = dequantize_nvfp4_to_dtype(
|
||||
a_fp4,
|
||||
a_scale_interleaved,
|
||||
input_global_scale,
|
||||
a_global,
|
||||
dtype=hidden_states.dtype,
|
||||
device=hidden_states.device,
|
||||
block_size=16,
|
||||
|
||||
@@ -907,24 +907,27 @@ def convert_gpt_oss_weight_to_mxfp4_moe_kernel_format(
|
||||
else:
|
||||
assert mxfp4_backend == Mxfp4MoeBackend.FLASHINFER_CUTLASS_MXFP4_BF16
|
||||
|
||||
def _interleave_mxfp4_cutlass_sm90(w):
|
||||
w_shape = w.shape
|
||||
w_interleaved = w.reshape(w_shape[0], w_shape[1], (w_shape[2] // 4), 4)
|
||||
w_interleaved = w_interleaved.permute(0, 2, 1, 3)
|
||||
w_interleaved = w_interleaved.reshape(
|
||||
w_shape[0], w_shape[2] // 4, w_shape[1] * 4
|
||||
)
|
||||
return w_interleaved
|
||||
from flashinfer.fused_moe import (
|
||||
interleave_moe_scales_for_sm90_mixed_gemm,
|
||||
interleave_moe_weights_for_sm90_mixed_gemm,
|
||||
)
|
||||
|
||||
w31_scales = w13_scale_swapped.to(torch.uint8)
|
||||
w31_scales_interleaved = _interleave_mxfp4_cutlass_sm90(w31_scales)
|
||||
|
||||
w2_scale = w2_weight_scale.data.to(torch.uint8)
|
||||
w2_scale_interleaved = _interleave_mxfp4_cutlass_sm90(w2_scale)
|
||||
w13_weight_interleaved = interleave_moe_weights_for_sm90_mixed_gemm(
|
||||
w13_weight_swapped.contiguous(), "fp4"
|
||||
)
|
||||
w2_weight_interleaved = interleave_moe_weights_for_sm90_mixed_gemm(
|
||||
w2_weight.contiguous(), "fp4"
|
||||
)
|
||||
w31_scales_interleaved = interleave_moe_scales_for_sm90_mixed_gemm(
|
||||
w13_scale_swapped.to(torch.uint8)
|
||||
)
|
||||
w2_scale_interleaved = interleave_moe_scales_for_sm90_mixed_gemm(
|
||||
w2_weight_scale.data.to(torch.uint8)
|
||||
)
|
||||
|
||||
return (
|
||||
w13_weight_swapped,
|
||||
w2_weight,
|
||||
w13_weight_interleaved,
|
||||
w2_weight_interleaved,
|
||||
w31_scales_interleaved,
|
||||
w2_scale_interleaved,
|
||||
w13_bias_swapped,
|
||||
|
||||
Reference in New Issue
Block a user