TileLang 性能分析器实战用 Analyzer 从 TIR 静态估算 FLOPs、全局内存流量与 Roofline 执行时间【免费下载链接】tilelangDomain-specific language designed to streamline the development of high-performance GPU/CPU/Accelerators kernels项目地址: https://gitcode.com/GitHub_Trending/ti/tilelangtilelang.tools.Analyzer是 TileLang 提供的静态性能分析工具它不编译、不运行内核而是直接检查 TileLang TIR中间表示估算一次内核执行将产生的浮点运算量FLOPs、全局内存流量global-memory traffic并基于内置硬件模型给出 Roofline 执行时间估计。读完本文你将掌握Analyzer.analysis()的完整用法与AnalysisResult各字段含义能像 examples/analyze/README.md 中的示例一样对 GEMM、卷积等 TileLang 内核做 FLOP 校验并能从源码层面理解其估算模型与当前局限从而在编译和实测之前先评估内核的算术成本。分析对象从 TileLang TIR 出发的静态估算Analyzer的输入是经过get_tir()特化后的 TIRtvm.tirx.PrimFunc或tvm.IRModule即一个已经填入了具体 tile 参数block_M、block_N、block_K等的PrimFunc。它的工作分两步IR 遍历统计通过 TVM 的 IR transform 回调遍历函数体识别其中的T.copy内存拷贝与T.gemm矩阵乘调用把它们分别计入全局字节数和 FLOPsRoofline 计算结合设备模型device model提供的计算能力表和内存带宽计算计算时间与内存时间取较大者作为estimated_time。官方文档 docs/tools/analyzer.md 强调结果是一个静态估计值static estimate不是 GPU 上的实测数据。它适合在编译与 benchmark 之前快速核对一个内核“理论上该做多少功”而不适合用来预测真实延迟。前置要求为什么分析内核也需要一块可见的 GPUexamples/analyze/README.md 在末尾给出了一个容易踩坑的前提示例从当前活跃的 PyTorch 运行时构造 CUDA 或 CDNA 设备模型因此即使分析器不执行内核也要求有一块可见的 GPU。原因在于设备模型tilelang.carver.arch.CUDA的构造函数tilelang/carver/arch/cuda.py会真实查询运行时device tvm.runtime.cuda(0) if not device.exist: raise RuntimeError(Cannot find cuda device 0.) self.name cuda_driver.get_device_name() self.device: tvm.runtime.Device device self.sm_version check_sm_version(self.target.attrs[arch]) self.compute_max_core device.multi_processor_count self.compute_capability device.compute_version.replace(., )也就是说CUDA(cuda)会去查询 CUDA 设备 0 的名称、SM 数量和 compute capability如80/86/89这些信息随后被分析器用于查内置峰值性能表。AMD 一侧的tilelang.carver.arch.CDNAtilelang/carver/arch/cdna.py同理会查询tvm.runtime.rocm(0)。从源码结构看分析器还依赖设备模型上的bandwidth字段CUDA 模型中它是硬编码的推荐值[750, 12080]单位 MB/s源码注释标注了 TODO希望未来能取到真实带宽CDNA 模型中为[1300, 14000]。分析器取第二个分量并换算为 GB/s见下文calculate()所以expected_bandwidth_GBps在 CUDA 模型下固定约为12080 / 1000 12.08这个数值应当理解为设备模型提供的参考带宽而非实测值。另外docs/tools/analyzer.md 说明分析器及其 NumPy 依赖包含在 TileLang 标准安装中不需要任何可选包。快速上手完整 Quick Start 示例以下示例继承自 docs/tools/analyzer.md 的 Quick Start用tilelang.jit定义一个分块 GEMM用get_tir()特化出 TIR再交给Analyzer分析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), threads128, ) 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_BTrue) T.copy(C_local, C[by * block_M, bx * block_N]) return C tir matmul.get_tir(block_M128, block_N128, block_K32) device CUDA(cuda) result Analyzer.analysis(tir, device) print(fFLOPs: {result.total_flops}) print(fGlobal bytes: {result.total_global_bytes}) print(fEstimated seconds: {result.estimated_time}) print(fModel peak TFLOPS: {result.expected_tflops}) print(fModel bandwidth (GB/s): {result.expected_bandwidth_GBps})对这个分块 GEMM文档给出的校验基准是total_flops应当等于2 * M * N * K本例即2 * 1024 * 1024 * 1024 2147483648。内存流量则由跨全局缓冲区边界的T.copy区域推导而来——注意上例中A/B的T.copy位于T.serial的 k 循环内会被循环次数放大而C_local到全局C的写出发生在 k 循环之外。GEMM 分析示例example_gemm_analyze.py仓库中的正式示例 examples/analyze/example_gemm_analyze.py 在 Quick Start 基础上更完整它使用了T.Pipelined软件流水、swizzle 调度和更真实的写回路径fragment → shared → global并自动区分 CUDA/AMD 平台import tilelang import tilelang.language as T from tilelang.tools import Analyzer from tilelang.carver.arch import CUDA from tilelang.carver.arch import CDNA import torch M N K 1024 tilelang.jit def kernel( A, B, block_MNone, block_NNone, block_KNone, num_stagesNone, thread_numNone, enable_rasterationNone, ): dtype T.float16 accum_dtype T.float32 A: T.Tensor((M, K), dtype) B: T.Tensor((N, K), dtype) C T.empty((M, N), dtype) with T.Kernel(T.ceildiv(N, block_N), T.ceildiv(M, block_M), threadsthread_num) as (bx, by): A_shared T.alloc_shared((block_M, block_K), dtype) B_shared T.alloc_shared((block_N, block_K), dtype) C_local T.alloc_fragment((block_M, block_N), accum_dtype) C_shared T.alloc_shared((block_M, block_N), dtype) T.use_swizzle(panel_size10, enableenable_rasteration) T.clear(C_local) for k in T.Pipelined(T.ceildiv(K, block_K), num_stagesnum_stages): 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_BTrue, ) T.copy(C_local, C_shared) T.copy(C_shared, C[by * block_M, bx * block_N]) return C def main(): my_func kernel.get_tir(block_M128, block_N128, block_K32, num_stages3, thread_num128, enable_rasterationTrue) cuda_device CUDA(cuda) if torch.version.hip is None else CDNA(hip) result Analyzer.analysis(my_func, cuda_device) print(fAnalyzed FLOPs: {result.total_flops}) print(fExpected FLOPs: {2 * M * N * K}) if __name__ __main__: main()参数逐项说明参数取值作用block_M/block_N/block_K128 / 128 / 32tile 分块尺寸k 循环次数为ceildiv(K, 32) 32num_stages3T.Pipelined流水级数只影响编译期调度不影响静态 FLOP 计数thread_num128每 block 线程数enable_rasterationTrue开启 swizzle 以提升 L2 命中同样不改变静态计算量示例的核心输出是实测对照Analyzed FLOPs与手算值2 * M * N * K打印在同一屏上用于验证分析器对这个标准分块 GEMM 的计数是否正确。由于每次T.gemm的调用粒度是128 × 128 × 32的子块k 循环 32 次、grid 为ceildiv(1024/128) × ceildiv(1024/128) 8 × 8 64个 block三者相乘后总量恰为2 * 1024 * 1024 * 1024——这正是分析器“逐调用计数 × 循环次数 × block 数”模型在标准 GEMM 上的自洽性验证。卷积分析示例im2col 分块 GEMM第二个示例 examples/analyze/example_conv_analyze.py 展示了分析器对非 GEMM 形态算子的适用方式把卷积表达成T.im2colim2col 展开T.gemm的分块 GEMM。完整代码如下import tilelang import tilelang.language as T from tilelang.tools import Analyzer from tilelang.carver.arch import CUDA from tilelang.carver.arch import CDNA import torch N 64 C 256 H 512 W 512 F 512 K 3 S 1 D 1 P 1 tilelang.jit def kernel(data, kernel, S, D, P, block_M, block_N, block_K, num_stages, threads, dtypeT.float16, accum_dtypeT.float32): N, C, H, W, F, K T.const(N, C, H, W, F, K) KH, KW K, K OH (H 2 * P - D * (K - 1) - 1) // S 1 OW (W 2 * P - D * (K - 1) - 1) // S 1 dtype T.float16 accum_dtype T.float32 data: T.Tensor((N, H, W, C), dtype) kernel: T.Tensor((KH, KW, C, F), dtype) out T.empty((N, OH, OW, F), dtype) with T.Kernel(T.ceildiv(F, block_N), T.ceildiv(N * OH * OW, block_M), threadsthreads) as (bx, by): data_shared T.alloc_shared((block_M, block_K), dtype) kernel_shared T.alloc_shared((block_K, block_N), dtype) out_local T.alloc_fragment((block_M, block_N), accum_dtype) out_shared T.alloc_shared((block_M, block_N), dtype) kernel_flat T.Tensor((KH * KW * C, F), dtype, kernel.data) out_flat T.Tensor((N * OH * OW, F), dtype, out.data) T.clear(out_local) for k_iter in T.Pipelined(T.ceildiv(KH * KW * C, block_K), num_stagesnum_stages): T.im2col(data, data_shared, by, k_iter, KH, S, D, P) T.copy(kernel_flat[k_iter * block_K, bx * block_N], kernel_shared) T.gemm(data_shared, kernel_shared, out_local) T.copy(out_local, out_shared) T.copy(out_shared, out_flat[by * block_M, bx * block_N]) return out def main(): my_func kernel.get_tir( NN, CC, HH, WW, FF, KK, SS, DD, PP, block_M64, block_N128, block_K32, num_stages3, threads256, ) cuda_device CUDA(cuda) if torch.version.hip is None else CDNA(hip) result Analyzer.analysis(my_func, cuda_device) print(result) print(fAnalyzed FLOPs: {result.total_flops}) if __name__ __main__: main()关键设计点卷积参数batchN64、输入通道C256、特征图HW512、输出通道F512、3×3卷积核K3、步长S1、膨胀D1、paddingP1。由公式OH (H 2P - D(K-1) - 1)//S 1得OH OW 512。im2col 化的张量视图kernel_flat把权重视作(KH*KW*C, F)out_flat把输出视作(N*OH*OW, F)于是“每个输出位置的卷积”变成“一行 im2col 矩阵乘权重矩阵”的标准 GEMM。k 维长度KH*KW*C 2304按block_K32切出 72 次迭代。T.im2col计入内存统计im2col 从全局输入data读入data_shared属于触达函数参数缓冲区的拷贝会被计入total_global_bytes而卷积的 FLOPs 完全由内层T.gemm调用累加得出。grid 维度bx覆盖ceildiv(F, block_N) 4by覆盖ceildiv(N*OH*OW, block_M)分析器会按blockIdx.x × blockIdx.y放大总量。这个示例的意义在于只要你的算子卷积、以及更多“摊平成 GEMM”的算子最终以T.gemm表达算术、以T.copy/T.im2col表达数据搬运分析器就能给出可对照的静态口径。运行方式与测试入口按 examples/analyze/README.md 的说明从仓库根目录运行python examples/analyze/example_gemm_analyze.py python examples/analyze/example_conv_analyze.py示例目录的文件分工如下继承自 README 的文件表文件覆盖内容example_gemm_analyze.py分块 GEMM 分析与 FLOP 校验example_conv_analyze.py通过T.im2col与T.gemm降级的卷积test_example_analyze.py两个示例的测试入口test_example_analyze.py 的实现非常直接——它就是把两个示例的main()各包装成一个 pytest 用例test_example_gemm_analyze与test_example_conv_analyze因此这两个示例也被纳入回归测试体系在 GPU 环境下运行该测试文件即可同时验证两个分析路径可执行且无异常。AnalysisResult 字段详解Analyzer.analysis(fn, device)返回一个不可变的AnalysisResultfrozen dataclass定义在 tilelang/tools/Analyzer.py。结合 docs/tools/analyzer.md 的字段表与源码实现各字段含义如下字段含义源码依据total_flops归因于被识别的T.gemm调用的 FLOPs_analyze_gemm()每次调用计2*M*N*Ktotal_global_bytes归因于被识别的、源或目的为全局缓冲区的T.copy调用的字节数_analyze_copy()按 region 元素数 × dtype 大小计estimated_time有计算模型时取“建模计算时间”与“内存时间”的较大者否则只取内存时间calculate()max(mem_time, compute_time)expected_tflops内置 compute capability 表选出的理论峰值不支持的架构返回Noneget_peak_tflops()expected_bandwidth_GBps从设备模型取得并换算为 GB/s 的带宽值bandwidth[1] / 1000calculate()注意estimated_time的取值逻辑tilelang/tools/Analyzer.pymem_time self.total_global_bytes / (bandwidth_GBps * 1e9) compute_time self.total_flops / (peak_tflops * 1e12) if peak_tflops else None estimated_time max(mem_time, compute_time) if peak_tflops else mem_time即设备模型有对应 compute capability 的峰值配置时走标准 roofline 取max(计算时间, 内存时间)没有时peak_tflops为None只能退化为纯内存时间估计。源码级原理分析器如何遍历 TIR理解 tilelang/tools/Analyzer.py 的实现可以清楚每个统计量的来源边界。1. 用 IR transform 做只读遍历Analyzer.ir_pass()通过tvm.tirx.transform.prim_func_pass注册了_pre_visit/_post_visit回调tilelang/tools/Analyzer.py在遍历过程中维护三样状态self.global_buffers来自f.buffer_map的函数参数缓冲区集合用于判断一次T.copy是否“触达全局内存”self.block_countsblockIdx.x/blockIdx.y两个 grid 维度的规模。它从thread_extent属性语句或带thread_binding的For节点中读取 extentself.loop_stack非线程绑定的For循环栈。进入循环时压入 extent离开时_post_visit弹出从而保证任意嵌套深度下“当前所在循环次数之积”始终正确。2. 只认两类操作遍历中只对Evaluate节点里的Call做判定tilelang/tools/Analyzer.py操作名在(tl.copy, tl.tileop.copy)中 →_analyze_copy()检查call.args中的源/目的缓冲区是否属于global_buffers然后统计 region 的总元素数乘以 dtype 字节数再乘上当前loop_stack各 extent 之积与blockIdx.x × blockIdx.y操作名在(tl.gemm, tl.tileop.gemm)中 →_analyze_gemm()从调用实参取M、N、K按2*M*N*K计 FLOPs同样乘上循环次数与 block 数。这就是官方文档“Limitations”一节所说行为的直接实现逐元素算术与其它内建指令不计入 FLOP直接对 buffer 的 load/store 与别的内存内建不计入字节数。同时blockIdx.z维度并未进入block_counts初始集合初始为{blockIdx.x: 1, blockIdx.y: 1}见init所以blockIdx.z的 grid 规模不参与当前计算。3. 内置峰值表只有 8.0 / 8.6 / 8.9 三档expected_tflops来自模块级的ARCH_CONFIGStilelang/tools/Analyzer.py# Each entry contains: (cores per SM, default clock (GHz), FLOPs per cycle, max SM count) ARCH_CONFIGS {80: (128, 1.41, 2, 108), 86: (128, 1.70, 2, 84), 89: (128, 2.52, 2, 128)}查表键是设备compute_capability的前两位80/86/89对应 Ampere 与 Ada 的固定架构参数。计算式为get_peak_tflops()total_cores compute_max_core * cores_per_sm # 最大 SM 数 × 每 SM core 数 tflops (total_cores * default_clock * flops_per_cycle) / 1e3 # GHz 下折算为 TFLOPS若架构不在表中例如 sm90 Hopper会打日志并返回None此时expected_tflops为None、estimated_time退化为内存时间。因此文档提醒expected_tflops与estimated_time应视为对比辅助comparison aids不是特定设备的 benchmark 结果——表中用的是固定架构参数而非你手上那块卡的实测频率与规格。4. 入口 API对外只有一个类方法tilelang/tools/Analyzer.pyresult Analyzer.analysis(fn, device) # 内部执行 cls(fn, device).ir_pass().calculate()fn包含 TileLang TIR 的tvm.IRModule或tvm.tirx.PrimFuncPrimFunc会被自动包进{main: fn}的 IRModule见initdeviceTileLang 设备描述对象如tilelang.carver.arch.CUDA(cuda)或CDNA(hip)。已知局限与使用建议综合 docs/tools/analyzer.md 的 Limitations 与 tilelang/tools/Analyzer.py 的实现把分析器结果用于决策前应明确以下边界它是简单的 roofline 估计不测量延迟、占用率occupancy、指令发射、缓存行为、访存合并coalescing、bank conflict 与流水重叠FLOP 口径只识别T.gemm调用逐元素算术T.reduce、、*等与其它内建不计入如果你的算子大量算力消耗在 softmax、归一化等非 GEMM 部分total_flops会系统性偏低内存口径只识别触达函数参数 buffer 的T.copy区域直接的 buffer load/store 与其它内存内建不计入grid 缩放只跟踪blockIdx.x与blockIdx.yblockIdx.z的 extent 不参与当前计算峰值表是固定架构值expected_tflops仅覆盖 8.0 / 8.6 / 8.9且用默认频率与最大 SM 数不是设备实测峰值带宽来自设备模型的推荐值CUDA 模型中为硬编码的12080 MB/s源码中留有获取真实带宽的 TODO跨设备比较时口径一致但绝对值不代表你设备的实测带宽。基于以上推荐的用法是把total_flops与2*M*N*K这类理论式子做一致性校验GEMM 示例已示范、用total_global_bytes对比不同 tiling 方案的访存量、用estimated_time在不同 block 尺寸/流水配置之间做相对比较再交给真实的编译与 benchmark 流程去确认端到端延迟。小结TileLang 的Analyzer提供了一条“先算后跑”的性能评估路径get_tir()特化 →Analyzer.analysis(tir, device)→ 读取AnalysisResult五个字段。仓库中 examples/analyze/ 下的 GEMM 与 im2col 卷积两个示例给出了可复制的最小工作流并配套了 pytest 测试入口而 docs/tools/analyzer.md 与 tilelang/tools/Analyzer.py 则完整界定了它的统计口径与建模边界。掌握它你可以在写内核的早期阶段就获得 FLOPs 与全局访存量的静态口径为 tiling 选择与 roofline 分析提供不依赖实测环境的参照。【免费下载链接】tilelangDomain-specific language designed to streamline the development of high-performance GPU/CPU/Accelerators kernels项目地址: https://gitcode.com/GitHub_Trending/ti/tilelang创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考