// TileLang Ascend API 使用最佳实践。提供内存分配、数据搬运、矩阵计算、归约、元素级运算、同步、调度原语等 API 的正确用法和最佳实践。触发：使用 TileLang API 编写 Ascend NPU kernel 时或遇到 API 相关问题时。
根据 pass-design.md 与 workflow 分析结果生成 TileLang-Ascend Pass 的最终实现代码（不含 UT/ST）。先输出实现骨架文档（pass-impl-skeleton.md）确认框架设计，再生成 C++ 实现、Python 封装、Pipeline 接入，并完成最小冒烟验证。测试生成由后续独立 skill 负责。触发关键词：实现 Pass、生成 Pass 代码、Pass 编码、根据设计文档实现 Pass、写 Pass 代码、落地 Pass、新增 Pass 实现。
TileLang-Ascend 算子测试设计技能。支持多种场景：(1) 从 design.md 设计测试配置 (2) 从 examples/{op}/*.py 补充测试 (3) 手动提供算子信息生成测试 (4) 测试覆盖率分析。理解算子实现逻辑后智能判断测试策略。触发：设计算子测试、生成测试用例、补充测试、测试覆盖率不足。
TileLang 算子性能调优与潜在性能劣化模式检查。提供性能数据采集、瓶颈诊断、优化实施、效果验证能力；也用于生成或评审算子时对照常见性能劣化模式示例检查当前 kernel 代码。触发：算子精度通过后需要优化性能、性能不及预期时。
TileLang Ascend Developer/Expert 模式选择与 pass_configs 配置指南。当需要确定编程模式、配置 pass_configs、或在两种模式之间转换时触发。API 详情请参考 tilelang-api-best-practices skill。
根据算子需求生成 TileLang-Ascend 算子设计文档（design.md）。涵盖编程模式选型（Developer/Expert/混合）、API 映射、内存层级规划、Tiling 策略、循环结构、同步策略、验证方案等。触发：设计算子、生成 design.md、算子方案设计、新算子开发、算子实现方案。
基于设计文档生成 TileLang-Ascend 算子实现代码与测试。从 design.md 中提取关键信息，结合 examples/ 中的参考实现生成可运行代码。触发：实现算子、写 kernel、生成代码、算子编码、根据设计文档实现。
| name | tilelang-api-best-practices |
| description | TileLang Ascend API 使用最佳实践。提供内存分配、数据搬运、矩阵计算、归约、元素级运算、同步、调度原语等 API 的正确用法和最佳实践。触发：使用 TileLang API 编写 Ascend NPU kernel 时或遇到 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，含 T.tile.atomic_add） | 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 两种范式 |
| 多 block/core 累加到 GM | api-compute | T.tile.atomic_add(dst_gm, src_local)，调用前显式清零 GM |
| CV 融合算子 | api-kernel-memory, api-schedule-sync | workspace 索引一致性、AUTO_CV_COMBINE、vid 并行化 |
| 流水线优化 | 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 | 说明 |
|---|---|
@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.tile.atomic_add(dst_gm, src_local) | 将本地 tensor 原子累加到 GM；V1 支持 local/UB → GM |
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 转储 |
| 配置项 | 说明 |
|---|---|
TL_ASCEND_AUTO_SYNC: True | 自动同步插入 |
TL_ASCEND_MEMORY_PLANNING: True | 自动内存规划 |
TL_ASCEND_AUTO_CV_COMBINE: True | 自动 CV 分离（核间流水线） |
TL_ASCEND_AUTO_CV_SYNC: True | 自动核间同步（核间流水线） |