tilelang.profiler.bench¶

Common entry point for GPU and wall-clock benchmarking.

Attributes¶

logger

device

Functions¶

do_bench(fn[, warmup, rep, _n_warmup, _n_repeat, ...])

Benchmark a callable with GPU timing or host wall-clock timing.

Module Contents¶

tilelang.profiler.bench.logger¶
tilelang.profiler.bench.device = 'cuda:0'¶
tilelang.profiler.bench.do_bench(fn, warmup=25, rep=100, _n_warmup=0, _n_repeat=0, quantiles=None, fast_flush=True, backend='event', return_mode='mean', device=None, cache_size=256, early_stop_baseline=None)¶

Benchmark a callable with GPU timing or host wall-clock timing.

The existing GPU timing methods provide accurate kernel timing by: - Clearing L2 cache between runs for consistent measurements - Auto-calculating warmup and repeat counts based on kernel runtime - Supporting multiple profiling backends (CUDA/MPS events, CUPTI, or CUDA graph replay) - Offering flexible result aggregation (mean/median/min/max/quantiles)

Wall timing measures host elapsed time without cache flushing. With no device or a CPU device, the callable is assumed to be synchronous. An explicit CUDA/HIP or MPS device enables synchronization before and after each wall-clock sample, including launch and synchronization overhead.

Parameters:

  • fn (collections.abc.Callable) – Function to benchmark

  • warmup (float) – Target warmup time in milliseconds (default: 25)

  • rep (float) – Target total benchmark time in milliseconds (default: 100)

  • _n_warmup (int) – Manual override for warmup iterations (default: 0 = auto)

  • _n_repeat (int) – Manual override for benchmark iterations (default: 0 = auto)

  • quantiles (list[float] | None) – Performance percentiles to compute (e.g., [0.5, 0.95])

  • fast_flush (bool) – Use faster GPU L2 cache flush with int32 vs int8 (default: True); ignored by “wall”

  • backend (Literal['event', 'cupti', 'cudagraph', 'wall']) – Timing method - “event”, “cupti”, “cudagraph”, or “wall” (default: “event”)

  • return_mode (Literal['min', 'max', 'mean', 'median']) – Result aggregation method - “mean”, “median”, “min”, or “max”

  • device (int | torch.device | None) – Optional device to benchmark on. CUDA/HIP events, streams, cache buffers, and synchronization are scoped to that device. Event timing also accepts MPS; wall timing accepts CPU and MPS.

  • cache_size (int) – GPU L2 cache flush buffer size in MB (default: 256); ignored by “wall”

  • early_stop_baseline (float | None)

Returns:

Runtime in milliseconds (float) or list of quantile values if quantiles specified

Return type:

float | list[float]