CANN/cannbot-skills：TileLang性能优化-编程实验室

【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体，本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills

TileLang 性能优化

Step 1: 精度校验（强制前置）→ Step 2: 性能数据采集 → Step 3: 算子类型判断 ↓ Step 4: 优化实施 → Step 5: 精度再验证 → Step 6: 效果验证

运行命令：python examples/{op_name}/<script_name>.py
<script_name>获取：在examples/{op_name}/目录下查找包含@tilelang.jit装饰器的 Python 脚本。若存在多个，询问用户确认使用哪一个，后续步骤自动复用该选择。
算子必须包含与参考实现的精度对比逻辑
精度未通过 → 禁止进入后续步骤

msprof op --kernel-name="main_kernel" --output=./msprof_output ./examples/{op_name}/<script_name>.py

生成翻译后的 Ascend C 代码：

在算子脚本中，JIT 编译返回的函数对象调用get_kernel_source()可获取翻译后的 Ascend C 代码：

func = jit_func(batch=B, seq_len=S, ...) print(func.get_kernel_source())

运行脚本后，从输出中搜索关键字判断算子类型：

判断依据	类型	典型算子
`IS_ASCEND_AIC`出现	Cube 型	GEMM、MatMul、Linear
`IS_ASCEND_AIV`出现	Vector 型	RoPE、Softmax、Add
两者均出现	混合型	FlashAttention、SparseFlashAttention

根据算子类型选择优化手段（详见 optimization-guide）：

生成/改写算子前的性能关注项检查：

若正在生成、改写或评审 kernel，先阅读 performance-antipatterns，对照其中的常见性能劣化模式示例检查当前代码是否存在类似 pattern
文档中的 pattern 不是正确性错误，而是需要重点关注的性能风险点；确需临时保留时，在optimization_log.md中记录 shape、dtype、保留原因和后续替换方案

优化方向	说明	典型手段
pass_configs 调优	调整编译器 pass 行为	关闭自动同步、关闭内存规划
核内优化	提升单核内指令并行度	Double Buffer、L1 常驻、指令向量化、Split-K pipelined GEMM
核间优化	优化 Cube/Vector 核间协作	num_stages 调优、同步优化、Fixed Core 模式
流水线优化	计算与访存重叠	T.Pipelined（核内/核间流水）、T.Persistent（数据块调度）
Fixed Core	按物理核数 launch，减少冗余初始化和显存膨胀	`T.Kernel(core_num, is_npu=True)`、Workspace 按物理核分配
指令融合	减少指令下发次数	AXPY 融合指令、broadcast 向量化
稀疏访存优化	离散数据高效搬运	双 vector 核访存、Gather + 连续搬出、异步拷贝

编程模式选择：

优先使用Developer 模式（自动内存规划、自动同步、编译器自动分离 Cube/Vector），参考 tilelang-programming-model-guide。

如无法满足性能要求，再使用Expert 模式手动控制（显式指定 L1/UB/L0 层级、手动同步、细粒度调度），参考同一文档。

参考以下算子的优化实践文档，学习具体优化技巧：

算子类型	最佳实践文档	核心优化技术
Vector 型	RoPE 优化实践	NPU 内动态生成 Mask、Tile API 向量化、参数简化
Cube 型	GEMM Intrinsic 优化	多缓冲流水线、细粒度 Flag 同步、MMA intrinsic、L0 分块、负载均衡
CV 融合型	Flash Attention 优化	num_stages 流水线、批量 Softmax、Cross-core Semaphore、数据布局优化

文档内容涵盖：

运行命令：python examples/{op_name}/<script_name>.py（<script_name>为examples/{op_name}/目录下的 Python 脚本文件名）
精度失败时：在保持已实施优化手段前提下调试精度，不能为修复精度撤销优化

调试手段：

方法	用途	示例
`T.printf(fmt, *args)`	设备端打印中间值	`T.printf("val=%d\n", val)`
`T.dump_tensor(tensor, desc, size, shape_info)`	转储张量内容到文件	`T.dump_tensor(acc_o, "acc_o", 128, (128, 128))`
查看生成的 Ascend C 代码	分析编译器生成逻辑	`print(func.get_kernel_source())`