external-cannbot-ops-lab-tilelang-skills-tilelang-api-best-practices

name	external-cannbot-ops-lab-tilelang-skills-tilelang-api-best-practices
description	TileLang Ascend API 使用最佳实践。提供内存分配、数据搬运、矩阵计算、归约、元素级运算、同步、调度原语等 API 的正确用法和最佳实践。触发：使用 TileLang API 编写 Ascend NPU kernel 时或遇到 API 相关问题时。
original-name	tilelang-api-best-practices
synced-from	https://gitcode.com/cann/cannbot-skills
synced-date	2026-05-26
synced-commit	ac5bbd2b4cf427d011874e11f8d1e8b1bef66eda
license	UNKNOWN

TileLang Ascend API 最佳实践

API 文档索引

文档	涵盖内容	典型场景
api-kernel-memory.md	Kernel 定义（T.prim_func, T.Kernel, @jit）、内存分配（Developer: T.alloc_shared/fragment/var, Expert: T.alloc_ub/L1/L0x）、数据搬运（T.copy）	Kernel 编写、片上存储管理、数据搬运
api-compute.md	矩阵计算（T.gemm_v0, T.mma）、归约（T.reduce_sum/max/min）、Element-wise（T.Parallel + 符号 API）、Tile 扩展原语（T.tile.xxx）	GEMM、Softmax、逐元素计算、排序
api-schedule-sync.md	循环（T.serial, T.unroll）、流水线（T.Pipelined）、持久化调度（T.Persistent）、同步（T.set_flag/wait_flag, T.barrier_all, T.set_cross_flag）、调试（T.printf, T.dump_tensor）	流水线优化、多核均衡、同步、调试

---

场景索引

使用场景	相关文档	关键技巧
GEMM 矩阵乘	api-compute, api-kernel-memory	shared→fragment 层级搬运、init 参数、T.barrier_all
Softmax/LayerNorm	api-compute	T.reduce_max/sum、T.tile.exp/sub/div
逐元素计算	api-compute	T.Parallel + 符号 API 或 T.tile.xxx 两种范式
流水线优化	api-schedule-sync	T.Pipelined num_stages、核间/核内流水线
多核负载均衡	api-schedule-sync	T.Persistent 缓存友好调度
排序	api-compute	T.tile.sort → T.tile.merge_sort → T.tile.topk
Kernel 调试	api-schedule-sync	T.printf、T.dump_tensor、get_kernel_source()

---

API 速查表

Kernel 定义

API	说明
@T.prim_func	定义 kernel 函数
T.Tensor((M, N), dtype)	声明张量参数
T.Kernel(block_num, is_npu=True) as (cid, vid)	Kernel 启动上下文
@jit(out_idx=[-1], pass_configs={...})	JIT 编译装饰器
T.dyn['K'] / T.dynamic('K', 'int32')	动态 shape

内存分配

API	说明	模式
T.alloc_shared(shape, dtype)	shared 层级（编译器自动判断 L1/UB）	Developer
T.alloc_fragment(shape, dtype)	fragment 层级（编译器自动判断 L0A/B/C）	Developer
T.alloc_var(dtype, init=...)	标量变量	Developer
T.alloc_ub / T.alloc_L1 / T.alloc_L0A/L0B/L0C	显式指定存储层级	Expert

数据搬运与计算

API	说明
T.copy(src, dst)	GM/L1/UB/L0 之间搬运数据
T.gemm_v0(A, B, C, transpose_A, transpose_B, init)	标准 GEMM
T.mma(A, B, C, init)	NPU MMA 指令
T.reduce_sum/max/min(buffer, out, dim)	按维度归约

循环与调度

API	说明
T.serial(N) / T.unroll(N)	普通循环 / 循环展开
T.Parallel(ext0, ext1, ...)	元素级并行循环
T.Pipelined(range, num_stages=N)	流水线并行
T.Persistent(domain, wave_size, index)	持久化调度

同步与调试

API	说明
T.set_flag / T.wait_flag	核内流水线同步
T.barrier_all() / T.pipe_barrier(pipe)	管线屏障
T.set_cross_flag / T.wait_cross_flag	核间同步
T.sync_all()	全局同步
T.printf(fmt, *args)	设备端格式化打印
T.dump_tensor(tensor, desc, size, shape_info)	Tensor 转储

常用 pass_configs

配置项	说明
TL_ASCEND_AUTO_SYNC: True	自动同步插入
TL_ASCEND_MEMORY_PLANNING: True	自动内存规划
TL_ASCEND_AUTO_CV_COMBINE: True	自动 CV 分离（核间流水线）
tl.ascend_auto_cross_core_sync: True	自动核间同步（核间流水线）