Add inductor_benchmarker as latency measurement option

[yf225]

As discussed in https://fb.workplace.com/groups/257735836456307/posts/967097558853461/?comment_id=971504168412800&reply_comment_id=971545435075340, the normal triton.testing.do_bench measurement includes torch.compile Dynamo invocation overhead and doesn't reflect real-world model use case where Dynamo overhead is usually hidden.

This PR adds an option to use Inductor benchmarker as the timing measurement tool, which uses cuda event for timing measurement and thus more accurately measuring only the CUDA kernel runtime.

[FindHao]

Besides that, I recalled @xuzhao9 mentioned somewhere we can use kineto's trace to sum only GPU time.

[FindHao]

I'm not sure about this change. Will this benchmark function change the results for pure handwritten triton kernel's results? If it matches the results of triton.testing.do_bench for other triton kernels, I feel we should just use this new function. If not, we need to understand what are the differences and see how to fix it.

Adding a new latency_measure_mode doesn't sound good because the results are diverged.

what do you think? @xuzhao9

[FindHao]

which uses cuda event for timing measurement and thus more accurately measuring only the CUDA kernel runtime.

https://github.com/triton-lang/triton/blob/37f265932b68868021b2fade6354b44e613dc124/python/triton/testing.py#L163
do_bench also use cuda event for timling. I may need to double check the implementation of benchmarker.benchmark_gpu

[BoyuanFeng]

@FindHao benchmarking.benchmark_gpu seems to be better for benchmarks than triton.testing.do_bench. Please check this scripts for comparison: P1908752761

use kineto's trace to sum only GPU time.

This sounds good! It would be great if we have a pytorch utility to benchmark only gpu time. We can converge all kernel benchmarks to use it if possible.

[FindHao]

benchmarking.benchmark_gpu seems to be better for benchmarks than triton.testing.do_bench.

Can you clarify what's the differences between these two implementations?

What I mean is: if you claim that the Dynamo overhead is completely hidden in real runs, I am willing to trust that. However, we still need to verify that the results from benchmark_gpu are consistent with those from do_bench on other handwritten Triton kernels.

Alternatively, measuring only the pure GPU time may be misleading, since it could overlook real Python or Triton runtime overhead.

[yf225]

measuring only the pure GPU time may be misleading, since it could overlook real Python or Triton runtime overhead

Would we be open to having a mode that measures only GPU time? I believe for real-world model use causes, the kernel launch or Python overhead is usually hidden by the previous kernel (assuming no bubble), so having a mode that only measures GPU time is useful and reflects that scenario too.

[FindHao]

measuring only the pure GPU time may be misleading, since it could overlook real Python or Triton runtime overhead

Would we be open to having a mode that measures only GPU time? I believe for real-world model use causes, the kernel launch or Python overhead is usually hidden by the previous kernel (assuming no bubble), so having a mode that only measures GPU time is useful and reflects that scenario too.

Try --metrics cuda_time.

tritonbench/docs/metrics.md

Line 15 in a404ea7

| `cuda_time` | Sum of all GPU-side kernels time of an operator backend, measured by Kineto and PyTorch Profiler. |

Is this what you need?

[xuzhao9]

Due to the limitation of triton's do_bench/do_bench_cudagraph, it is beneficial to have inductor benchmarker and more latency measurement modes (e.g. power-limit-aware benchmarking) in the future.

[xuzhao9]

I am wondering does inductor benchmarker use cudagraph?