pass a function for autotuner via 'do_bench' param

Signed-off-by: Anatoly Myachev <[email protected]>

Author: anmyachev

benchmarks/triton_kernels_benchmark/__init__.py

@@ -1,4 +1,4 @@
-from .benchmark_testing import do_bench, assert_close, perf_report, Benchmark, BENCHMARKING_METHOD  # type: ignore # noqa: F401
+from .benchmark_testing import do_bench, make_do_bench_for_autotune, assert_close, perf_report, Benchmark, BENCHMARKING_METHOD  # type: ignore # noqa: F401
 
 if BENCHMARKING_METHOD == "UPSTREAM_PYTORCH_PROFILER":
     from triton.runtime import driver

benchmarks/triton_kernels_benchmark/benchmark_testing.py

@@ -2,7 +2,6 @@
 import itertools
 import os
 from typing import Any, Dict, List
-import triton
 
 BENCHMARKING_METHOD = os.getenv("BENCHMARKING_METHOD", "UPSTREAM_PYTORCH_PROFILER")
 
@@ -160,7 +159,13 @@ def extract_kernels(funcs):
 else:
     raise NotImplementedError(f"BENCHMARKING_METHOD: {BENCHMARKING_METHOD} isn't implemented")
 
-triton.testing.do_bench = do_bench
+
+def make_do_bench_for_autotune(kernel_name: str):
+
+    def autotuner_do_bench(*args, **kwargs):
+        return do_bench(*args, n_warmup=10, n_repeat=10, kernel_name=kernel_name, **kwargs)
+
+    return autotuner_do_bench
 
 
 def assert_close(x, y, atol=None, rtol=None, err_msg=""):

benchmarks/triton_kernels_benchmark/flash_attention_fwd_benchmark.py

@@ -161,7 +161,8 @@ def _attn_fwd(Q, K, V, sm_scale, M, Out,  #
     for w in [8, 16, 32] \
     ]
 
-tuner = triton.autotune(configs, key=['N_CTX', 'BLOCK_DMODEL'])
+tuner = triton.autotune(configs, key=['N_CTX', 'BLOCK_DMODEL'],
+                        do_bench=benchmark_suit.make_do_bench_for_autotune('_attn_fwd'))
 tune_attn_fwd = tuner(_attn_fwd)
 
 

benchmarks/triton_kernels_benchmark/fused_softmax.py

@@ -50,6 +50,7 @@ def naive_softmax(x):
         triton.Config({"threads_per_warp": 16}, num_warps=4),
     ],
     key=["BLOCK_SIZE_X", "BLOCK_SIZE_Y"],
+    do_bench=benchmark_suit.make_do_bench_for_autotune(kernel_name="softmax_kernel"),
 )
 @triton.jit
 def softmax_kernel(output_ptr, input_ptr, input_row_stride, output_row_stride, n_cols, BLOCK_SIZE_X: tl.constexpr,

benchmarks/triton_kernels_benchmark/gemm_benchmark.py

@@ -39,6 +39,7 @@
             num_stages=s, num_warps=32) for s in [2, 3]
     ],
     key=['M', 'N', 'K'],
+    do_bench=benchmark_suit.make_do_bench_for_autotune(kernel_name='matmul_kernel_with_block_pointers'),
 )
 @triton.jit
 def matmul_kernel_with_block_pointers(
@@ -112,6 +113,7 @@ def matmul_kernel_with_block_pointers(
             num_stages=s, num_warps=4) for s in [2]
     ],
     key=['M', 'N', 'K'],
+    do_bench=benchmark_suit.make_do_bench_for_autotune(kernel_name='matmul_kernel_with_block_pointers'),
 )
 @triton.jit
 def matmul_kernel_with_block_pointers_batched(

benchmarks/triton_kernels_benchmark/gemm_postop_addmatrix_benchmark.py

@@ -32,6 +32,7 @@
             num_stages=2, num_warps=32),
     ],
     key=['M', 'N', 'K'],
+    do_bench=benchmark_suit.make_do_bench_for_autotune(kernel_name='matmul_kernel_with_block_pointers'),
 )
 @triton.jit
 def matmul_kernel_with_block_pointers(
@@ -106,6 +107,7 @@ def matmul_kernel_with_block_pointers(
             num_stages=2, num_warps=4),
     ],
     key=['M', 'N', 'K'],
+    do_bench=benchmark_suit.make_do_bench_for_autotune(kernel_name='matmul_kernel_with_block_pointers'),
 )
 @triton.jit
 def matmul_kernel_with_block_pointers_batched(

benchmarks/triton_kernels_benchmark/gemm_postop_gelu_benchmark.py

@@ -51,6 +51,7 @@ def gelu(x):
             num_stages=2, num_warps=32),
     ],
     key=['M', 'N', 'K'],
+    do_bench=benchmark_suit.make_do_bench_for_autotune(kernel_name='matmul_kernel_with_block_pointers'),
 )
 @triton.jit
 def matmul_kernel_with_block_pointers(
@@ -119,6 +120,7 @@ def matmul_kernel_with_block_pointers(
             num_stages=2, num_warps=4),
     ],
     key=['M', 'N', 'K'],
+    do_bench=benchmark_suit.make_do_bench_for_autotune(kernel_name='matmul_kernel_with_block_pointers'),
 )
 @triton.jit
 def matmul_kernel_with_block_pointers_batched(

benchmarks/triton_kernels_benchmark/gemm_preop_exp_benchmark.py

@@ -33,6 +33,7 @@
             num_stages=2, num_warps=32),
     ],
     key=['M', 'N', 'K'],
+    do_bench=benchmark_suit.make_do_bench_for_autotune(kernel_name='matmul_kernel_with_block_pointers'),
 )
 @triton.jit
 def matmul_kernel_with_block_pointers(
@@ -104,6 +105,7 @@ def matmul_kernel_with_block_pointers(
             num_stages=2, num_warps=4),
     ],
     key=['M', 'N', 'K'],
+    do_bench=benchmark_suit.make_do_bench_for_autotune(kernel_name='matmul_kernel_with_block_pointers'),
 )
 @triton.jit
 def matmul_kernel_with_block_pointers_batched(

benchmarks/triton_kernels_benchmark/gemm_splitk_benchmark.py

@@ -12,6 +12,7 @@
                       num_stages=4, num_warps=32),
     ],
     key=['M', 'N', 'K'],
+    do_bench=benchmark_suit.make_do_bench_for_autotune(kernel_name='_kernel'),
 )
 @triton.jit
 def _kernel(A, B, C,  #

benchmarks/triton_kernels_benchmark/gemm_streamk_benchmark.py

@@ -104,6 +104,7 @@ def mac_loop(
             num_stages=2, num_warps=32),
     ],
     key=['M', 'N', 'K'],
+    do_bench=benchmark_suit.make_do_bench_for_autotune(kernel_name='first_wave'),
 )
 @triton.jit
 def first_wave(
@@ -140,6 +141,7 @@ def first_wave(
             num_stages=2, num_warps=32),
     ],
     key=['M', 'N', 'K'],
+    do_bench=benchmark_suit.make_do_bench_for_autotune(kernel_name='full_tiles'),
 )
 @triton.jit
 def full_tiles(