Performance Analyzer¶

tilelang.tools.Analyzer inspects TileLang TIR and estimates its floating-point work, global-memory traffic, and roofline execution time. It is useful for checking a kernel’s expected arithmetic cost before compiling and benchmarking it. The result is a static estimate, not a measurement from the GPU.

Requirements¶

The analyzer and its NumPy dependency are included in the standard TileLang installation. No optional package is required.

The device model passed to the analyzer supplies architecture and bandwidth information. For example, constructing CUDA("cuda") queries CUDA device 0, so that workflow requires a working CUDA runtime and a visible GPU even though the analyzer does not execute the kernel.

Quick Start¶

Create specialized TIR with get_tir(), then analyze the resulting PrimFunc:

import tilelang
import tilelang.language as T
from tilelang.carver.arch import CUDA
from tilelang.tools import Analyzer

M = N = K = 1024


@tilelang.jit
def matmul(A, B, block_M, block_N, block_K):
    A: T.Tensor((M, K), T.float16)
    B: T.Tensor((N, K), T.float16)
    C = T.empty((M, N), T.float16)

    with T.Kernel(
        T.ceildiv(N, block_N),
        T.ceildiv(M, block_M),
        threads=128,
    ) as (bx, by):
        A_shared = T.alloc_shared((block_M, block_K), T.float16)
        B_shared = T.alloc_shared((block_N, block_K), T.float16)
        C_local = T.alloc_fragment((block_M, block_N), T.float32)

        T.clear(C_local)
        for k in T.serial(T.ceildiv(K, block_K)):
            T.copy(A[by * block_M, k * block_K], A_shared)
            T.copy(B[bx * block_N, k * block_K], B_shared)
            T.gemm(A_shared, B_shared, C_local, transpose_B=True)
        T.copy(C_local, C[by * block_M, bx * block_N])

    return C


tir = matmul.get_tir(block_M=128, block_N=128, block_K=32)
device = CUDA("cuda")
result = Analyzer.analysis(tir, device)

print(f"FLOPs: {result.total_flops}")
print(f"Global bytes: {result.total_global_bytes}")
print(f"Estimated seconds: {result.estimated_time}")
print(f"Model peak TFLOPS: {result.expected_tflops}")
print(f"Model bandwidth (GB/s): {result.expected_bandwidth_GBps}")

For this tiled GEMM, total_flops should equal 2 * M * N * K. Memory traffic is derived from the T.copy regions that cross the global-buffer boundary.

API and Result¶

The main entry point is:

result = Analyzer.analysis(fn, device)

  • fn is a tvm.IRModule or tvm.tirx.PrimFunc containing TileLang TIR.

  • device is a TileLang device description such as tilelang.carver.arch.CUDA("cuda").

  • The method returns an immutable AnalysisResult.

AnalysisResult exposes these fields:

Field

Meaning

total_flops

FLOPs attributed to recognized T.gemm calls.

total_global_bytes

Bytes attributed to recognized T.copy calls with a global source or destination.

estimated_time

Maximum of modeled compute time and memory time when a compute model is available; otherwise memory time.

expected_tflops

The theoretical peak selected by the built-in compute-capability table, or None for an unsupported capability.

expected_bandwidth_GBps

The bandwidth value obtained from the device model and converted to GB/s.

The analyzer multiplies each recognized operation by its enclosing non-thread- bound TIR loop extents and the blockIdx.x and blockIdx.y grid extents. Its compute-time model currently has entries for NVIDIA compute capabilities 8.0, 8.6, and 8.9.

Limitations¶

  • This is a simple roofline estimate. It does not measure latency, occupancy, instruction issue, cache behavior, memory coalescing, bank conflicts, or pipeline overlap.

  • FLOP counting recognizes T.gemm calls only. Elementwise arithmetic and other intrinsics are not included.

  • Memory counting recognizes T.copy regions that touch a function parameter buffer. Direct buffer loads and stores and other memory intrinsics are not included.

  • Grid scaling tracks blockIdx.x and blockIdx.y; a blockIdx.z extent is not part of the current calculation.

  • The built-in peak-TFLOPS table uses fixed architecture values. Treat expected_tflops and estimated_time as comparison aids, not device-specific benchmark results.

Examples¶