| 1 | # SPDX-License-Identifier: Apache-2.0 |
| 2 | # SPDX-FileCopyrightText: Copyright contributors to the vLLM project |
| 3 | """Tests for the full FP8 ViT attention path (quantize -> cuDNN -> un-pad).""" |
| 4 | |
| 5 | import contextlib |
| 6 | |
| 7 | import pytest |
| 8 | import torch |
| 9 | |
| 10 | from vllm.triton_utils import HAS_TRITON |
| 11 | from vllm.utils.flashinfer import ( |
| 12 | is_flashinfer_cudnn_fp8_prefill_attn_supported, |
| 13 | ) |
| 14 | from vllm.v1.attention.backends.registry import AttentionBackendEnum |
| 15 | |
| 16 | |
| 17 | def _has_flashinfer_cudnn() -> bool: |
| 18 | """Check if FlashInfer cuDNN backend is available.""" |
| 19 | try: |
| 20 | from flashinfer.prefill import ( |
| 21 | cudnn_batch_prefill_with_kv_cache, # noqa: F401 |
| 22 | ) |
| 23 | |
| 24 | return True |
| 25 | except ImportError: |
| 26 | return False |
| 27 | |
| 28 | |
| 29 | HEAD_DIMS = [72, 80] |
| 30 | SEQ_LENS = [256] |
| 31 | NUM_HEADS = [16] |
| 32 | |
| 33 | |
| 34 | @pytest.fixture |
| 35 | def _fp8_attention(): |
| 36 | """Create FP8-enabled MMEncoderAttention via config.""" |
| 37 | from unittest.mock import MagicMock, patch |
| 38 | |
| 39 | from vllm.config import ModelConfig, VllmConfig, set_current_vllm_config |
| 40 | from vllm.config.multimodal import MultiModalConfig |
| 41 | |
| 42 | if not is_flashinfer_cudnn_fp8_prefill_attn_supported(): |
| 43 | pytest.skip("FlashInfer cuDNN FP8 prefill attention not supported") |
| 44 | |
| 45 | mm_config = MultiModalConfig(mm_encoder_attn_dtype="fp8") |
| 46 | vllm_config = VllmConfig() |
| 47 | vllm_config.model_config = MagicMock(spec=ModelConfig, multimodal_config=mm_config) |
| 48 | |
| 49 | # MMEncoderAttention reads torch.get_default_dtype() during init |
| 50 | # to determine the output dtype. In real model loading this is bf16. |
| 51 | old_dtype = torch.get_default_dtype() |
| 52 | torch.set_default_dtype(torch.bfloat16) |
| 53 | |
| 54 | with ( |
| 55 | set_current_vllm_config(vllm_config), |
| 56 | patch( |
| 57 | "vllm.model_executor.layers.attention.mm_encoder_attention" |
| 58 | ".get_vit_attn_backend", |
| 59 | return_value=AttentionBackendEnum.FLASHINFER, |
| 60 | ), |
| 61 | ): |
| 62 | yield |
| 63 | |
| 64 | torch.set_default_dtype(old_dtype) |
| 65 | |
| 66 | |
| 67 | def _build_cu_seqlens_and_meta( |
| 68 | seq_len: int, |
| 69 | num_heads: int, |
| 70 | head_dim: int, |
| 71 | fp8_padded_hidden_size: int | None = None, |
| 72 | ) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor]: |
| 73 | """Build cu_seqlens, max_seqlen, sequence_lengths for a single sequence.""" |
| 74 | import numpy as np |
| 75 | |
| 76 | from vllm.model_executor.layers.attention.mm_encoder_attention import ( |
| 77 | MMEncoderAttention, |
| 78 | ) |
| 79 | |
| 80 | cu_seqlens_np = np.array([0, seq_len], dtype=np.int32) |
| 81 | |
| 82 | sequence_lengths = MMEncoderAttention.maybe_compute_seq_lens( |
| 83 | AttentionBackendEnum.FLASHINFER, |
| 84 | cu_seqlens_np, |
| 85 | torch.device("cuda"), |
| 86 | ) |
| 87 | |
| 88 | max_seqlen = torch.tensor( |
| 89 | MMEncoderAttention.compute_max_seqlen( |
| 90 | AttentionBackendEnum.FLASHINFER, cu_seqlens_np |
| 91 | ), |
| 92 | dtype=torch.int32, |
| 93 | ) |
| 94 | |
| 95 | cu_seqlens = MMEncoderAttention.maybe_recompute_cu_seqlens( |
| 96 | AttentionBackendEnum.FLASHINFER, |
| 97 | cu_seqlens_np, |
| 98 | num_heads * head_dim, |
| 99 | 1, # tp_size |
| 100 | torch.device("cuda"), |
| 101 | fp8_padded_hidden_size=fp8_padded_hidden_size, |
| 102 | ) |
| 103 | |
| 104 | return cu_seqlens, max_seqlen, sequence_lengths |
| 105 | |
| 106 | |
| 107 | @pytest.mark.skipif( |
| 108 | not (HAS_TRITON and _has_flashinfer_cudnn()), |
| 109 | reason="Triton and FlashInfer cuDNN required", |
| 110 | ) |
| 111 | @pytest.mark.parametrize("head_dim", HEAD_DIMS) |
| 112 | @pytest.mark.parametrize("seq_len", SEQ_LENS) |
| 113 | @pytest.mark.parametrize("num_heads", NUM_HEADS) |
| 114 | def test_fp8_attn_output_shape( |
| 115 | head_dim: int, |
| 116 | seq_len: int, |
| 117 | num_heads: int, |
| 118 | _fp8_attention, |
| 119 | ) -> None: |
| 120 | """Verify FP8 attention produces correct output shape after un-padding.""" |
| 121 | from vllm.model_executor.layers.attention.mm_encoder_attention import ( |
| 122 | MMEncoderAttention, |
| 123 | ) |
| 124 | from vllm.utils.math_utils import round_up |
| 125 | |
| 126 | attn = None |
| 127 | with contextlib.suppress(ValueError, ImportError): |
| 128 | attn = MMEncoderAttention( |
| 129 | num_heads=num_heads, |
| 130 | head_size=head_dim, |
| 131 | prefix="visual.blocks.0.attn", |
| 132 | ).to("cuda") |
| 133 | |
| 134 | if attn is None or not attn.fp8_enabled: |
| 135 | pytest.skip("FP8 MMEncoderAttention not available") |
| 136 | assert attn is not None # mypy narrowing |
| 137 | |
| 138 | # FP8 always needs fp8_padded_hidden_size for correct cu_seqlens |
| 139 | fp8_padded_hidden_size = num_heads * round_up(head_dim, 16) |
| 140 | |
| 141 | cu_seqlens, max_seqlen, sequence_lengths = _build_cu_seqlens_and_meta( |
| 142 | seq_len, num_heads, head_dim, fp8_padded_hidden_size=fp8_padded_hidden_size |
| 143 | ) |
| 144 | |
| 145 | q = torch.randn( |
| 146 | seq_len, |
| 147 | num_heads, |
| 148 | head_dim, |
| 149 | device="cuda", |
| 150 | dtype=torch.bfloat16, |
| 151 | ) |
| 152 | k = torch.randn_like(q) |
| 153 | v = torch.randn_like(q) |
| 154 | |
| 155 | output = attn._forward_flashinfer(q, k, v, cu_seqlens, max_seqlen, sequence_lengths) |
| 156 | |
| 157 | # Output should have original head_dim (un-padded) |
| 158 | assert output.shape[-1] == head_dim |
| 159 | assert output.dtype == torch.bfloat16 |
| 160 | |
| 161 | |
| 162 | @pytest.mark.skipif( |
| 163 | not (HAS_TRITON and _has_flashinfer_cudnn()), |
| 164 | reason="Triton and FlashInfer cuDNN required", |
| 165 | ) |
| 166 | @pytest.mark.parametrize("head_dim", HEAD_DIMS) |
| 167 | @pytest.mark.parametrize("seq_len", SEQ_LENS) |
| 168 | @pytest.mark.parametrize("num_heads", NUM_HEADS) |
| 169 | def test_fp8_vs_bf16_close( |
| 170 | head_dim: int, seq_len: int, num_heads: int, _fp8_attention |
| 171 | ) -> None: |
| 172 | """FP8 attention output should be reasonably close to BF16 baseline.""" |
| 173 | from vllm.model_executor.layers.attention.mm_encoder_attention import ( |
| 174 | MMEncoderAttention, |
| 175 | ) |
| 176 | from vllm.utils.math_utils import round_up |
| 177 | |
| 178 | torch.manual_seed(42) |
| 179 | q = torch.randn( |
| 180 | 1, |
| 181 | seq_len, |
| 182 | num_heads, |
| 183 | head_dim, |
| 184 | device="cuda", |
| 185 | dtype=torch.bfloat16, |
| 186 | ) |
| 187 | k = torch.randn_like(q) |
| 188 | v = torch.randn_like(q) |
| 189 | |
| 190 | # FP8 path |
| 191 | attn_fp8 = None |
| 192 | with contextlib.suppress(ValueError, ImportError): |
| 193 | attn_fp8 = MMEncoderAttention( |
| 194 | num_heads=num_heads, |
| 195 | head_size=head_dim, |
| 196 | prefix="visual.blocks.0.attn", |
| 197 | ).to("cuda") |
| 198 | |
| 199 | if attn_fp8 is None or not attn_fp8.fp8_enabled: |
| 200 | pytest.skip("FP8 MMEncoderAttention not available") |
| 201 | assert attn_fp8 is not None # mypy narrowing |
| 202 | |
| 203 | fp8_padded_hidden_size = num_heads * round_up(head_dim, 16) |
| 204 | cu_seqlens, max_seqlen, seq_lengths = _build_cu_seqlens_and_meta( |
| 205 | seq_len, |
| 206 | num_heads, |
| 207 | head_dim, |
| 208 | fp8_padded_hidden_size=fp8_padded_hidden_size, |
| 209 | ) |
| 210 | |
| 211 | out_fp8 = attn_fp8._forward_flashinfer( |
| 212 | q.clone(), |
| 213 | k.clone(), |
| 214 | v.clone(), |
| 215 | cu_seqlens, |
| 216 | max_seqlen, |
| 217 | seq_lengths, |
| 218 | ) |
| 219 | |
| 220 | # BF16 baseline (create non-FP8 attention by using scale=attn_fp8.scale |
| 221 | # and calling the wrapper directly without FP8 quantization) |
| 222 | from vllm.model_executor.layers.attention.mm_encoder_attention import ( |
| 223 | _get_flashinfer_workspace_buffer, |
| 224 | ) |
| 225 | from vllm.v1.attention.ops.vit_attn_wrappers import ( |
| 226 | vit_flashinfer_wrapper, |
| 227 | ) |
| 228 | |
| 229 | out_bf16 = vit_flashinfer_wrapper( |
| 230 | q=q.clone(), |
| 231 | k=k.clone(), |
| 232 | v=v.clone(), |
| 233 | scale=attn_fp8.scale, |
| 234 | workspace_buffer=_get_flashinfer_workspace_buffer(), |
| 235 | cu_seqlens=cu_seqlens, |
| 236 | max_seqlen=max_seqlen, |
| 237 | sequence_lengths=seq_lengths, |
| 238 | ) |
| 239 | |
| 240 | out_fp8_f = out_fp8.float() |
| 241 | out_bf16_f = out_bf16.float() |
| 242 | |
| 243 | abs_diff = (out_fp8_f - out_bf16_f).abs() |
| 244 | abs_diff_flat = abs_diff.flatten() |
| 245 | |
| 246 | # Relative diff (avoid division by zero) |
| 247 | denom = out_bf16_f.abs().clamp(min=1e-6) |
| 248 | rel_diff_flat = (abs_diff / denom).flatten() |
| 249 | |
| 250 | cosine_sim = torch.nn.functional.cosine_similarity( |
| 251 | out_fp8_f.flatten().unsqueeze(0), |
| 252 | out_bf16_f.flatten().unsqueeze(0), |
| 253 | ).item() |
| 254 | |
| 255 | pcts = [50, 90, 95, 99, 99.9] |
| 256 | abs_pct = {p: torch.quantile(abs_diff_flat, p / 100).item() for p in pcts} |
| 257 | rel_pct = {p: torch.quantile(rel_diff_flat, p / 100).item() for p in pcts} |
| 258 | |
| 259 | print(f"\nFP8 vs BF16 (head_dim={head_dim}, seq_len={seq_len}):") |
| 260 | print(f" cosine_sim={cosine_sim:.6f}") |
| 261 | print( |
| 262 | f" abs_diff: max={abs_diff_flat.max().item():.6f}, " |
| 263 | f"mean={abs_diff_flat.mean().item():.6f}, " |
| 264 | + ", ".join(f"p{p}={abs_pct[p]:.6f}" for p in pcts) |
| 265 | ) |
| 266 | print( |
| 267 | f" rel_diff: max={rel_diff_flat.max().item():.6f}, " |
| 268 | f"mean={rel_diff_flat.mean().item():.6f}, " |
| 269 | + ", ".join(f"p{p}={rel_pct[p]:.6f}" for p in pcts) |
| 270 | ) |
| 271 | |
| 272 | assert abs_diff_flat.max().item() < 0.3, ( |
| 273 | f"FP8 vs BF16 max abs diff too large: {abs_diff_flat.max().item()}" |
| 274 | ) |
| 275 | assert abs_diff_flat.mean().item() < 0.03, ( |
| 276 | f"FP8 vs BF16 mean abs diff too large: {abs_diff_flat.mean().item()}" |
| 277 | ) |
| 278 | assert cosine_sim > 0.99, f"Cosine similarity too low: {cosine_sim:.6f}" |