(D, dtype, device, is_causal, is_fp8=False)
| 88 | |
| 89 | @torch.no_grad() |
| 90 | def _test_benchmark_attn_func(D, dtype, device, is_causal, is_fp8=False): |
| 91 | import triton |
| 92 | |
| 93 | torch.manual_seed(0) |
| 94 | |
| 95 | B = 16 |
| 96 | H = 16 |
| 97 | S_Q = 8192 |
| 98 | S_KV = 8192 |
| 99 | |
| 100 | query = torch.randn(B, H, S_Q, D, dtype=dtype, device=device) |
| 101 | key = torch.randn(B, H, S_KV, D, dtype=dtype, device=device) |
| 102 | value = torch.randn(B, H, S_KV, D, dtype=dtype, device=device) |
| 103 | |
| 104 | def attention_fn(): |
| 105 | if is_fp8: |
| 106 | fp8_attention(query, key, value, is_causal) |
| 107 | else: |
| 108 | vanilla_attention(query, key, value, is_causal) |
| 109 | |
| 110 | try: |
| 111 | attention_fn() |
| 112 | except ValueError as e: |
| 113 | pytest.skip(str(e)) |
| 114 | |
| 115 | def fa_fn(): |
| 116 | flash_attention(query, key, value, is_causal) |
| 117 | |
| 118 | def cudnn_sdpa_fn(): |
| 119 | cudnn_sdpa(query, key, value, is_causal) |
| 120 | |
| 121 | flops_per_matmul = 2 * B * H * S_Q * S_KV * D |
| 122 | total_flops = 2 * flops_per_matmul |
| 123 | |
| 124 | if is_causal: |
| 125 | total_flops //= 2 |
| 126 | |
| 127 | ms_fa = triton.testing.do_bench(fa_fn) |
| 128 | tflops_fa = total_flops * 1e-12 / (ms_fa * 1e-3) |
| 129 | print(f"TFLOPS (Flash Attention): {tflops_fa:.2f}") |
| 130 | |
| 131 | if D <= 128: |
| 132 | ms_cudnn_sdpa = triton.testing.do_bench(cudnn_sdpa_fn) |
| 133 | tflops_cudnn_sdpa = total_flops * 1e-12 / (ms_cudnn_sdpa * 1e-3) |
| 134 | print(f"TFLOPS (CUDNN SDPA): {tflops_cudnn_sdpa:.2f}") |
| 135 | |
| 136 | ms_quantum_attention = triton.testing.do_bench(attention_fn) |
| 137 | tflops_quantum_attention = total_flops * 1e-12 / (ms_quantum_attention * 1e-3) |
| 138 | print(f"TFLOPS (Quantum Attention): {tflops_quantum_attention:.2f}") |
| 139 | |
| 140 | |
| 141 | @pytest.mark.parametrize("D", [64, 128, 256]) |
no test coverage detected