Skip to content

Commit 5930ffd

Browse files
zhanwei33insanecoder-zzj
authored andcommitted
Print autotune benchmark timings (triton-lang#1392)
Co-authored-by: zhanwei33 <22679270+zhanwei33@users.noreply.github.com>
1 parent c47ebc7 commit 5930ffd

2 files changed

Lines changed: 76 additions & 1 deletion

File tree

third_party/ascend/backend/runtime/autotuner.py

Lines changed: 20 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -80,6 +80,14 @@ def _inject_default_simt_stack_limit(options: Dict[str, object], stack_limit: in
8080
options["simt_stack_limit"] = stack_limit
8181

8282

83+
def _format_autotune_timing(timing) -> str:
84+
"""Format the timing returned by the active autotune benchmarker."""
85+
if isinstance(timing, (tuple, list)):
86+
labels = ("p50", "p20", "p80")
87+
return ", ".join(f"{label}={value:.4f} ms" for label, value in zip(labels, timing))
88+
return f"mean={timing:.4f} ms"
89+
90+
8391
def _get_constexpr_candidates_from_fn(fn) -> List[str]:
8492
"""
8593
Returns all constexpr parameter names from the kernel function definition.
@@ -2131,6 +2139,7 @@ def benchmark():
21312139
if self.print_autotuning and did_benchmark:
21322140
print(f"Triton autotuning for function {self.base_fn.__name__} finished after "
21332141
f"{self.bench_time:.2f}s; best config selected: {self.best_config};")
2142+
self._print_benchmark_results(self.configs_timings)
21342143

21352144
if did_benchmark and self.auto_profile_dir is not None:
21362145
self._profile(*args, config=self.best_config, **kwargs)
@@ -2173,6 +2182,15 @@ def _try_ubtuner(self, *args, config, excp, run_fns, **kwargs):
21732182
if self.print_autotuning:
21742183
print(f"[WARN] encounter exception when try ubtune, Details: {e}")
21752184

2185+
def _print_benchmark_results(self, timings) -> None:
2186+
if not self.print_autotuning:
2187+
return
2188+
2189+
print(f"Triton autotuning benchmark results for function {self.base_fn.__name__}:")
2190+
for config, timing in timings.items():
2191+
selected = " [selected]" if config == self.best_config else ""
2192+
print(f" config={config}; {_format_autotune_timing(timing)}{selected}")
2193+
21762194
def _batch_bench(self, *args, configs, **kwargs):
21772195
from triton.compiler.errors import CompileTimeAssertionFailure, MLIRCompilationError
21782196
from triton.runtime.errors import OutOfResources
@@ -2607,7 +2625,8 @@ def kernel(x_ptr, x_size, **META):
26072625
26082626
If the environment variable :code:`TRITON_PRINT_AUTOTUNING` is set to
26092627
:code:`"1"`, Triton will print a message to stdout after autotuning each
2610-
kernel, including the time spent autotuning and the best configuration.
2628+
kernel, including the benchmark timing for each valid configuration, the
2629+
time spent autotuning, and the best configuration.
26112630
26122631
:param configs: a list of :code:`triton.Config` objects
26132632
:type configs: list[triton.Config]

third_party/ascend/unittest/autotune_ut/test_do_bench_compat.py

Lines changed: 56 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -176,6 +176,62 @@ def _dummy_kernel():
176176
assert marker["called"] is False
177177

178178

179+
def test_print_benchmark_results_formats_quantiles_and_selected_config(capsys):
180+
181+
def _dummy_kernel():
182+
return None
183+
184+
selected = Config({"BLOCK_SIZE": 128}, num_warps=4)
185+
other = Config({"BLOCK_SIZE": 256}, num_warps=8)
186+
tuner = object.__new__(AutoTilingTuner)
187+
tuner.print_autotuning = True
188+
tuner.base_fn = _dummy_kernel
189+
tuner.best_config = selected
190+
191+
tuner._print_benchmark_results({
192+
selected: (1.23456, 1.0, 1.8),
193+
other: 2.34567,
194+
})
195+
196+
output = capsys.readouterr().out
197+
assert "Triton autotuning benchmark results for function _dummy_kernel:" in output
198+
assert "config=BLOCK_SIZE: 128, num_warps: 4" in output
199+
assert "p50=1.2346 ms, p20=1.0000 ms, p80=1.8000 ms [selected]" in output
200+
assert "config=BLOCK_SIZE: 256, num_warps: 8" in output
201+
assert "mean=2.3457 ms" in output
202+
203+
204+
def test_print_benchmark_results_is_disabled_without_debug_flag(capsys):
205+
tuner = object.__new__(AutoTilingTuner)
206+
tuner.print_autotuning = False
207+
208+
tuner._print_benchmark_results({Config({"BLOCK_SIZE": 128}): 1.0})
209+
210+
assert capsys.readouterr().out == ""
211+
212+
213+
def test_run_prints_benchmark_results_after_tuning(capsys):
214+
215+
def _dummy_kernel():
216+
return None
217+
218+
selected = Config({"BLOCK_SIZE": 128}, num_warps=4)
219+
other = Config({"BLOCK_SIZE": 256}, num_warps=8)
220+
tuner, _ = _make_run_tuner([selected, other])
221+
tuner.cache_results = False
222+
tuner.print_autotuning = True
223+
tuner.base_fn = _dummy_kernel
224+
tuner._batch_bench = lambda *args, configs, **kwargs: {
225+
selected: (1.0, 0.9, 1.1),
226+
other: (2.0, 1.8, 2.2),
227+
}
228+
229+
assert tuner.run() == "kernel-result"
230+
output = capsys.readouterr().out
231+
assert "Triton autotuning benchmark results for function _dummy_kernel:" in output
232+
assert "p50=1.0000 ms, p20=0.9000 ms, p80=1.1000 ms [selected]" in output
233+
234+
179235
def test_ascend_autotune_decorator_forwards_do_bench(monkeypatch):
180236
import triton.backends.ascend.runtime.autotuner as ascend_autotuner
181237

0 commit comments

Comments
 (0)