| 1 | # SPDX-License-Identifier: Apache-2.0 |
| 2 | # SPDX-FileCopyrightText: Copyright contributors to the vLLM project |
| 3 | |
| 4 | """Benchmark complete AITER BF16 and FP8 ViT attention calls on ROCm. |
| 5 | |
| 6 | The FP8 timing includes Q/K/V quantization and attention. Dynamic scales are |
| 7 | calibrated once per input before timing so the measured path matches serving |
| 8 | with a static scale file. |
| 9 | |
| 10 | Example: |
| 11 | python benchmarks/kernels/benchmark_vit_aiter_fp8_attn.py \ |
| 12 | --seq-lens 2304 4096 8192 16384 --head-dim 72 |
| 13 | |
| 14 | """ |
| 15 | |
| 16 | from types import SimpleNamespace |
| 17 | |
| 18 | import torch |
| 19 | |
| 20 | from vllm.config import VllmConfig, set_current_vllm_config |
| 21 | from vllm.config.multimodal import MultiModalConfig |
| 22 | from vllm.model_executor.layers.attention.mm_encoder_attention import ( |
| 23 | MMEncoderAttention, |
| 24 | ) |
| 25 | from vllm.platforms import current_platform |
| 26 | from vllm.triton_utils import triton |
| 27 | from vllm.utils.argparse_utils import FlexibleArgumentParser |
| 28 | from vllm.v1.attention.backends.registry import AttentionBackendEnum |
| 29 | |
| 30 | |
| 31 | def make_attention(num_heads: int, head_dim: int, fp8: bool) -> MMEncoderAttention: |
| 32 | mm_config = MultiModalConfig( |
| 33 | mm_encoder_attn_backend=AttentionBackendEnum.ROCM_AITER_FA, |
| 34 | mm_encoder_attn_dtype="fp8" if fp8 else None, |
| 35 | ) |
| 36 | vllm_config = VllmConfig() |
| 37 | vllm_config.model_config = SimpleNamespace(multimodal_config=mm_config) |
| 38 | with set_current_vllm_config(vllm_config): |
| 39 | return MMEncoderAttention(num_heads, head_dim).to("cuda") |
| 40 | |
| 41 | |
| 42 | def bench( |
| 43 | seq_lens: list[int], |
| 44 | num_heads: int, |
| 45 | head_dim: int, |
| 46 | warmup_ms: int, |
| 47 | repeat_ms: int, |
| 48 | ) -> None: |
| 49 | if not current_platform.is_rocm(): |
| 50 | raise RuntimeError("This benchmark requires ROCm and AITER.") |
| 51 | |
| 52 | old_dtype = torch.get_default_dtype() |
| 53 | torch.set_default_dtype(torch.bfloat16) |
| 54 | try: |
| 55 | fp8_attention = make_attention(num_heads, head_dim, fp8=True) |
| 56 | bf16_attention = make_attention(num_heads, head_dim, fp8=False) |
| 57 | finally: |
| 58 | torch.set_default_dtype(old_dtype) |
| 59 | |
| 60 | print( |
| 61 | f"{'seq_len':>8} {'BF16 ms':>12} {'FP8 ms':>12} " |
| 62 | f"{'speedup':>10} {'FP8/BF16':>12}" |
| 63 | ) |
| 64 | print("-" * 60) |
| 65 | for seq_len in seq_lens: |
| 66 | torch.manual_seed(0) |
| 67 | qkv = torch.randn( |
| 68 | 1, |
| 69 | seq_len, |
| 70 | 3, |
| 71 | num_heads, |
| 72 | head_dim, |
| 73 | device="cuda", |
| 74 | dtype=torch.bfloat16, |
| 75 | ) |
| 76 | q, k, v = qkv.unbind(dim=2) |
| 77 | cu_seqlens = torch.tensor([0, seq_len], device="cuda", dtype=torch.int32) |
| 78 | max_seqlen = torch.tensor(seq_len, device="cuda", dtype=torch.int32) |
| 79 | |
| 80 | # Calibrate per-tensor scales once, then benchmark the static-scale path. |
| 81 | fp8_attention._fp8_dynamic_scale = True |
| 82 | fp8_attention._forward_aiter_fp8(q, k, v, cu_seqlens, max_seqlen) |
| 83 | fp8_attention._fp8_dynamic_scale = False |
| 84 | |
| 85 | bf16_ms = triton.testing.do_bench( |
| 86 | lambda q=q, k=k, v=v, cu=cu_seqlens, ms=max_seqlen: ( |
| 87 | bf16_attention._forward_fa(q, k, v, cu, ms) |
| 88 | ), |
| 89 | warmup=warmup_ms, |
| 90 | rep=repeat_ms, |
| 91 | ) |
| 92 | fp8_ms = triton.testing.do_bench( |
| 93 | lambda q=q, k=k, v=v, cu=cu_seqlens, ms=max_seqlen: ( |
| 94 | fp8_attention._forward_aiter_fp8(q, k, v, cu, ms) |
| 95 | ), |
| 96 | warmup=warmup_ms, |
| 97 | rep=repeat_ms, |
| 98 | ) |
| 99 | speedup = bf16_ms / fp8_ms |
| 100 | ratio = fp8_ms / bf16_ms |
| 101 | print( |
| 102 | f"{seq_len:>8} {bf16_ms:>12.3f} {fp8_ms:>12.3f} " |
| 103 | f"{speedup:>9.2f}x {ratio:>12.3f}" |
| 104 | ) |
| 105 | |
| 106 | |
| 107 | if __name__ == "__main__": |
| 108 | parser = FlexibleArgumentParser( |
| 109 | description="Benchmark AITER BF16 vs FP8 ViT attention." |
| 110 | ) |
| 111 | parser.add_argument( |
| 112 | "--seq-lens", type=int, nargs="+", default=[2304, 4096, 8192, 16384] |
| 113 | ) |
| 114 | parser.add_argument("--num-heads", type=int, default=16) |
| 115 | parser.add_argument("--head-dim", type=int, default=72) |
| 116 | parser.add_argument("--warmup-ms", type=int, default=100) |
| 117 | parser.add_argument("--repeat-ms", type=int, default=500) |
| 118 | args = parser.parse_args() |
| 119 | bench( |
| 120 | args.seq_lens, |
| 121 | args.num_heads, |
| 122 | args.head_dim, |
| 123 | args.warmup_ms, |
| 124 | args.repeat_ms, |
| 125 | ) |