【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skillsname: tilelang-api-best-practices description: TileLang Ascend API 使用最佳实践。提供内存分配、数据搬运、矩阵计算、归约、元素级运算、同步、调度原语等 API 的正确用法和最佳实践。触发使用 TileLang API 编写 Ascend NPU kernel 时或遇到 API 相关问题时。TileLang Ascend API 最佳实践API 文档索引文档涵盖内容典型场景api-kernel-memory.mdKernel 定义T.prim_func, T.Kernel, jit、内存分配Developer: T.alloc_shared/fragment/var, Expert: T.alloc_ub/L1/L0x、数据搬运T.copyKernel 编写、片上存储管理、数据搬运api-compute.md矩阵计算T.gemm_v0, T.mma、归约T.reduce_sum/max/min、Element-wiseT.Parallel 符号 API、Tile 扩展原语T.tile.xxx含 T.tile.atomic_addGEMM、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-memoryshared→fragment 层级搬运、init 参数、T.barrier_allSoftmax/LayerNormapi-computeT.reduce_max/sum、T.tile.exp/sub/div逐元素计算api-computeT.Parallel 符号 API 或 T.tile.xxx 两种范式多 block/core 累加到 GMapi-computeT.tile.atomic_add(dst_gm, src_local)调用前显式清零 GMCV 融合算子api-kernel-memory, api-schedule-syncworkspace 索引一致性、AUTO_CV_COMBINE、vid 并行化流水线优化api-schedule-syncT.Pipelined num_stages、核间/核内流水线多核负载均衡api-schedule-syncT.Persistent 缓存友好调度排序api-computeT.tile.sort → T.tile.merge_sort → T.tile.topkKernel 调试api-schedule-syncT.printf、T.dump_tensor、get_kernel_source()API 速查表Kernel 定义API说明T.prim_func定义 kernel 函数T.Tensor((M, N), dtype)声明张量参数T.Kernel(block_num, is_npuTrue) 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/UBDeveloperT.alloc_fragment(shape, dtype)fragment 层级编译器自动判断 L0A/B/CDeveloperT.alloc_var(dtype, init...)标量变量DeveloperT.alloc_ub / T.alloc_L1 / T.alloc_L0A/L0B/L0C显式指定存储层级Expert数据搬运与计算API说明T.copy(src, dst)GM/L1/UB/L0 之间搬运数据T.tile.atomic_add(dst_gm, src_local)将本地 tensor 原子累加到 GMV1 支持 local/UB → GMT.gemm_v0(A, B, C, transpose_A, transpose_B, init)标准 GEMMT.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_stagesN)流水线并行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_CV_SYNC: True自动核间同步核间流水线【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
CANN TileLang API最佳实践
【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skillsname: tilelang-api-best-practices description: TileLang Ascend API 使用最佳实践。提供内存分配、数据搬运、矩阵计算、归约、元素级运算、同步、调度原语等 API 的正确用法和最佳实践。触发使用 TileLang API 编写 Ascend NPU kernel 时或遇到 API 相关问题时。TileLang Ascend API 最佳实践API 文档索引文档涵盖内容典型场景api-kernel-memory.mdKernel 定义T.prim_func, T.Kernel, jit、内存分配Developer: T.alloc_shared/fragment/var, Expert: T.alloc_ub/L1/L0x、数据搬运T.copyKernel 编写、片上存储管理、数据搬运api-compute.md矩阵计算T.gemm_v0, T.mma、归约T.reduce_sum/max/min、Element-wiseT.Parallel 符号 API、Tile 扩展原语T.tile.xxx含 T.tile.atomic_addGEMM、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-memoryshared→fragment 层级搬运、init 参数、T.barrier_allSoftmax/LayerNormapi-computeT.reduce_max/sum、T.tile.exp/sub/div逐元素计算api-computeT.Parallel 符号 API 或 T.tile.xxx 两种范式多 block/core 累加到 GMapi-computeT.tile.atomic_add(dst_gm, src_local)调用前显式清零 GMCV 融合算子api-kernel-memory, api-schedule-syncworkspace 索引一致性、AUTO_CV_COMBINE、vid 并行化流水线优化api-schedule-syncT.Pipelined num_stages、核间/核内流水线多核负载均衡api-schedule-syncT.Persistent 缓存友好调度排序api-computeT.tile.sort → T.tile.merge_sort → T.tile.topkKernel 调试api-schedule-syncT.printf、T.dump_tensor、get_kernel_source()API 速查表Kernel 定义API说明T.prim_func定义 kernel 函数T.Tensor((M, N), dtype)声明张量参数T.Kernel(block_num, is_npuTrue) 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/UBDeveloperT.alloc_fragment(shape, dtype)fragment 层级编译器自动判断 L0A/B/CDeveloperT.alloc_var(dtype, init...)标量变量DeveloperT.alloc_ub / T.alloc_L1 / T.alloc_L0A/L0B/L0C显式指定存储层级Expert数据搬运与计算API说明T.copy(src, dst)GM/L1/UB/L0 之间搬运数据T.tile.atomic_add(dst_gm, src_local)将本地 tensor 原子累加到 GMV1 支持 local/UB → GMT.gemm_v0(A, B, C, transpose_A, transpose_B, init)标准 GEMMT.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_stagesN)流水线并行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_CV_SYNC: True自动核间同步核间流水线【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
