| 1 | # SPDX-License-Identifier: Apache-2.0 |
| 2 | # SPDX-FileCopyrightText: Copyright contributors to the vLLM project |
| 3 | |
| 4 | import itertools |
| 5 | |
| 6 | import torch |
| 7 | |
| 8 | from vllm.benchmarks.lib.utils import default_vllm_config |
| 9 | from vllm.model_executor.layers.rotary_embedding import get_rope |
| 10 | from vllm.triton_utils import triton |
| 11 | from vllm.utils.argparse_utils import FlexibleArgumentParser |
| 12 | |
| 13 | batch_size_range = [2**i for i in range(0, 8, 2)] |
| 14 | seq_len_range = [2**i for i in range(6, 10, 1)] |
| 15 | num_heads_range = [32, 48] |
| 16 | configs = list(itertools.product(batch_size_range, seq_len_range, num_heads_range)) |
| 17 | |
| 18 | |
| 19 | def get_benchmark(head_size, rotary_dim, is_neox_style, device): |
| 20 | @triton.testing.perf_report( |
| 21 | triton.testing.Benchmark( |
| 22 | x_names=["batch_size", "seq_len", "num_heads"], |
| 23 | x_vals=[list(_) for _ in configs], |
| 24 | line_arg="provider", |
| 25 | line_vals=["torch", "flashinfer", "vllm"], |
| 26 | line_names=["PyTorch", "FlashInfer", "vLLM"], |
| 27 | styles=[("blue", "-"), ("green", "-"), ("red", "-")], |
| 28 | ylabel="us", |
| 29 | plot_name=f"rope-perf{'-neox-style' if is_neox_style else ''}", |
| 30 | args={}, |
| 31 | ) |
| 32 | ) |
| 33 | @default_vllm_config() |
| 34 | def benchmark(batch_size, seq_len, num_heads, provider): |
| 35 | dtype = torch.bfloat16 |
| 36 | max_position = 8192 |
| 37 | rope_parameters = {"partial_rotary_factor": rotary_dim / head_size} |
| 38 | rope = get_rope(head_size, max_position, is_neox_style, rope_parameters) |
| 39 | rope = rope.to(dtype=dtype, device=device) |
| 40 | cos_sin_cache = rope.cos_sin_cache.to(dtype=torch.float, device=device) |
| 41 | |
| 42 | positions = torch.randint(0, max_position, (batch_size, seq_len), device=device) |
| 43 | query = torch.randn( |
| 44 | (batch_size, seq_len, num_heads * head_size), dtype=dtype, device=device |
| 45 | ) |
| 46 | key = torch.randn_like(query) |
| 47 | |
| 48 | quantiles = [0.5, 0.2, 0.8] |
| 49 | |
| 50 | if provider == "torch": |
| 51 | ms, min_ms, max_ms = triton.testing.do_bench( |
| 52 | lambda: rope.forward_native(positions, query.clone(), key.clone()), |
| 53 | quantiles=quantiles, |
| 54 | ) |
| 55 | elif provider == "flashinfer": |
| 56 | ms, min_ms, max_ms = triton.testing.do_bench( |
| 57 | lambda: torch.ops.vllm.flashinfer_rotary_embedding( |
| 58 | positions, |
| 59 | query.clone(), |
| 60 | key.clone(), |
| 61 | head_size, |
| 62 | cos_sin_cache, |
| 63 | is_neox_style, |
| 64 | ), |
| 65 | quantiles=quantiles, |
| 66 | ) |
| 67 | else: |
| 68 | ms, min_ms, max_ms = triton.testing.do_bench( |
| 69 | lambda: rope.forward_cuda(positions, query.clone(), key.clone()), |
| 70 | quantiles=quantiles, |
| 71 | ) |
| 72 | |
| 73 | return 1000 * ms, 1000 * max_ms, 1000 * min_ms |
| 74 | |
| 75 | return benchmark |
| 76 | |
| 77 | |
| 78 | if __name__ == "__main__": |
| 79 | parser = FlexibleArgumentParser( |
| 80 | description="Benchmark the rotary embedding kernels." |
| 81 | ) |
| 82 | parser.add_argument("--is-neox-style", type=bool, default=True) |
| 83 | parser.add_argument("--batch-size", type=int, default=16) |
| 84 | parser.add_argument("--seq-len", type=int, default=512) |
| 85 | parser.add_argument("--num-heads", type=int, default=8) |
| 86 | parser.add_argument( |
| 87 | "--head-size", |
| 88 | type=int, |
| 89 | choices=[64, 80, 96, 112, 120, 128, 192, 256], |
| 90 | default=128, |
| 91 | ) |
| 92 | parser.add_argument("--rotary-dim", type=int, choices=[16, 32], default=32) |
| 93 | parser.add_argument( |
| 94 | "--dtype", type=str, choices=["bfloat16", "float"], default="float" |
| 95 | ) |
| 96 | parser.add_argument("--seed", type=int, default=0) |
| 97 | parser.add_argument( |
| 98 | "--device", type=str, choices=["cuda:0", "cuda:1"], default="cuda:0" |
| 99 | ) |
| 100 | parser.add_argument("--save-path", type=str, default="./configs/rope/") |
| 101 | args = parser.parse_args() |
| 102 | |
| 103 | # Get the benchmark function |
| 104 | benchmark = get_benchmark( |
| 105 | args.head_size, args.rotary_dim, args.is_neox_style, args.device |
| 106 | ) |
| 107 | # Run performance benchmark |
| 108 | benchmark.run(print_data=True, save_path=args.save_path) |