| 1 | # SPDX-License-Identifier: Apache-2.0 |
| 2 | # SPDX-FileCopyrightText: Copyright contributors to the vLLM project |
| 3 | """Compare the CuTeDSL and Triton implementations of the DSA fused query op. |
| 4 | |
| 5 | ``fused_q`` dispatches to the CuTeDSL kernel on SM100 when the dtypes and |
| 6 | shapes are supported and falls back to Triton otherwise. This benchmark forces |
| 7 | each path so the two can be compared directly. |
| 8 | """ |
| 9 | |
| 10 | import argparse |
| 11 | from unittest.mock import patch |
| 12 | |
| 13 | import torch |
| 14 | |
| 15 | import vllm.models.deepseek_v32.common.kernels as K |
| 16 | from vllm.triton_utils import triton |
| 17 | |
| 18 | # GLM-5.2 / DeepSeek-V3.2 DSA dimensions. |
| 19 | NOPE_DIM = 512 |
| 20 | ROPE_DIM = 64 |
| 21 | INDEX_HEADS = 32 |
| 22 | INDEX_HEAD_DIM = 128 |
| 23 | MAX_POS = 16384 |
| 24 | |
| 25 | # num_heads is per-rank: 128 attention heads over TP 8/4/2. |
| 26 | NUM_HEADS = [16, 32, 64] |
| 27 | # MTP=5 makes a concurrency-N decode batch 6*N tokens. |
| 28 | NUM_TOKENS = [1, 6, 24, 96, 384, 1536, 4096] |
| 29 | |
| 30 | |
| 31 | def make_inputs(num_tokens: int, num_heads: int): |
| 32 | dev = "cuda" |
| 33 | bf16 = torch.bfloat16 |
| 34 | return dict( |
| 35 | positions=torch.randint(0, MAX_POS, (num_tokens,), device=dev), |
| 36 | q_pe=torch.randn(num_tokens, num_heads, ROPE_DIM, device=dev, dtype=bf16), |
| 37 | q_pe_cos_sin_cache=torch.randn(MAX_POS, ROPE_DIM, device=dev, dtype=bf16), |
| 38 | index_q=torch.randn( |
| 39 | num_tokens, INDEX_HEADS, INDEX_HEAD_DIM, device=dev, dtype=bf16 |
| 40 | ), |
| 41 | index_q_cos_sin_cache=torch.randn(MAX_POS, ROPE_DIM, device=dev, dtype=bf16), |
| 42 | ql_nope=torch.randn(num_tokens, num_heads, NOPE_DIM, device=dev, dtype=bf16), |
| 43 | q_scale=torch.tensor([0.5], device=dev, dtype=torch.float32), |
| 44 | index_weights=torch.randn( |
| 45 | num_tokens, INDEX_HEADS, device=dev, dtype=torch.float32 |
| 46 | ), |
| 47 | index_weights_softmax_scale=INDEX_HEAD_DIM**-0.5, |
| 48 | index_weights_head_scale=INDEX_HEADS**-0.5, |
| 49 | ) |
| 50 | |
| 51 | |
| 52 | def run(kwargs, *, cutedsl: bool): |
| 53 | # is_fused_q_cutedsl_supported gates the dispatch inside fused_q. |
| 54 | from vllm.models.deepseek_v32.nvidia.ops import fused_q_cutedsl as C |
| 55 | |
| 56 | with patch.object(C, "is_fused_q_cutedsl_supported", lambda *_, **__: cutedsl): |
| 57 | K.fused_q(**kwargs) |
| 58 | |
| 59 | |
| 60 | @triton.testing.perf_report( |
| 61 | [ |
| 62 | triton.testing.Benchmark( |
| 63 | x_names=["num_tokens"], |
| 64 | x_vals=NUM_TOKENS, |
| 65 | line_arg="provider", |
| 66 | line_vals=["triton", "cutedsl"], |
| 67 | line_names=["Triton", "CuTeDSL"], |
| 68 | styles=[("blue", "-"), ("green", "-")], |
| 69 | ylabel="us", |
| 70 | plot_name=f"fused_q-num_heads-{nh}", |
| 71 | args={"num_heads": nh}, |
| 72 | ) |
| 73 | for nh in NUM_HEADS |
| 74 | ] |
| 75 | ) |
| 76 | def benchmark(num_tokens, num_heads, provider): |
| 77 | kwargs = make_inputs(num_tokens, num_heads) |
| 78 | fn = lambda: run(kwargs, cutedsl=provider == "cutedsl") # noqa: E731 |
| 79 | fn() # warm up JIT / kernel compilation before timing |
| 80 | ms, min_ms, max_ms = triton.testing.do_bench(fn, quantiles=[0.5, 0.2, 0.8]) |
| 81 | return 1000 * ms, 1000 * max_ms, 1000 * min_ms |
| 82 | |
| 83 | |
| 84 | if __name__ == "__main__": |
| 85 | parser = argparse.ArgumentParser(description=__doc__) |
| 86 | parser.add_argument("--save-path", type=str, default=None) |
| 87 | args = parser.parse_args() |
| 88 | |
| 89 | from vllm.platforms import current_platform |
| 90 | from vllm.utils.import_utils import has_cutedsl |
| 91 | |
| 92 | if not (current_platform.has_device_capability(100) and has_cutedsl()): |
| 93 | raise SystemExit("fused_q CuTeDSL requires an SM100 device and cutlass-dsl") |
| 94 | |
| 95 | benchmark.run(print_data=True, show_plots=False, save_path=args.save_path) |