1---2name: tilelang-api-best-practices3description: TileLang Ascend API 使用最佳实践。提供内存分配、数据搬运、矩阵计算、归约、元素级运算、同步、调度原语等 API 的正确用法和最佳实践。触发：使用 TileLang API 编写 Ascend NPU kernel 时或遇到 API 相关问题时。4---56# TileLang Ascend API 最佳实践78---910## API 文档索引1112| 文档 | 涵盖内容 | 典型场景 |13|------|---------|---------|14| [api-kernel-memory.md](references/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 编写、片上存储管理、数据搬运 |15| [api-compute.md](references/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、逐元素计算、排序、原子累加 |16| [api-schedule-sync.md](references/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） | 流水线优化、多核均衡、同步、调试 |1718---1920## 场景索引2122| 使用场景 | 相关文档 | 关键技巧 |23|---------|---------|---------|24| **GEMM 矩阵乘** | [api-compute](references/api-compute.md), [api-kernel-memory](references/api-kernel-memory.md) | shared→fragment 层级搬运、init 参数、T.barrier_all |25| **Softmax/LayerNorm** | [api-compute](references/api-compute.md) | T.reduce_max/sum、T.tile.exp/sub/div |26| **逐元素计算** | [api-compute](references/api-compute.md) | T.Parallel + 符号 API 或 T.tile.xxx 两种范式 |27| **多 block/core 累加到 GM** | [api-compute](references/api-compute.md) | T.tile.atomic_add(dst_gm, src_local)，调用前显式清零 GM |28| **CV 融合算子** | [api-kernel-memory](references/api-kernel-memory.md), [api-schedule-sync](references/api-schedule-sync.md) | workspace 索引一致性、AUTO_CV_COMBINE、vid 并行化 |29| **流水线优化** | [api-schedule-sync](references/api-schedule-sync.md) | T.Pipelined num_stages、核间/核内流水线 |30| **多核负载均衡** | [api-schedule-sync](references/api-schedule-sync.md) | T.Persistent 缓存友好调度 |31| **排序** | [api-compute](references/api-compute.md) | T.tile.sort → T.tile.merge_sort → T.tile.topk |32| **Kernel 调试** | [api-schedule-sync](references/api-schedule-sync.md) | T.printf、T.dump_tensor、get_kernel_source() |33| **dtype 标量回退适配** | [api-compute](references/api-compute.md) | 先确认硬件支持；同宽 reinterpret / kernel 内 cast / record-aware DMA / 块 DMA + UB-local fallback；宽 dtype lane 拆分仅作已验证实验 |3435---3637## API 速查表3839### Kernel 定义4041| API | 说明 |42|-----|------|43| `@T.prim_func` | 定义 kernel 函数 |44| `T.Tensor((M, N), dtype)` | 声明张量参数 |45| `T.Kernel(block_num, is_npu=True) as (cid, vid)` | Kernel 启动上下文 |46| `@jit(out_idx=[-1], pass_configs={...})` | JIT 编译装饰器 |47| `T.dyn['K']` / `T.dynamic('K', 'int32')` | 动态 shape |4849### 内存分配5051| API | 说明 | 模式 |52|-----|------|------|53| `T.alloc_shared(shape, dtype)` | shared 层级（编译器自动判断 L1/UB） | Developer |54| `T.alloc_fragment(shape, dtype)` | fragment 层级（编译器自动判断 L0A/B/C） | Developer |55| `T.alloc_var(dtype, init=...)` | 标量变量 | Developer |56| `T.alloc_ub / T.alloc_L1 / T.alloc_L0A/L0B/L0C` | 显式指定存储层级 | Expert |5758### 数据搬运与计算5960| API | 说明 |61|-----|------|62| `T.copy(src, dst)` | GM/L1/UB/L0 之间搬运数据 |63| `T.tile.atomic_add(dst_gm, src_local)` | 将本地 tensor 原子累加到 GM；V1 支持 local/UB → GM |64| `T.gemm_v0(A, B, C, transpose_A, transpose_B, init)` | 标准 GEMM |65| `T.mma(A, B, C, init)` | NPU MMA 指令 |66| `T.reduce_sum/max/min(buffer, out, dim)` | 按维度归约 |6768### 循环与调度6970| API | 说明 |71|-----|------|72| `T.serial(N)` / `T.unroll(N)` | 普通循环 / 循环展开 |73| `T.Parallel(ext0, ext1, ...)` | 元素级并行循环 |74| `T.Pipelined(range, num_stages=N)` | 流水线并行 |75| `T.Persistent(domain, wave_size, index)` | 持久化调度 |7677### 同步与调试7879| API | 说明 |80|-----|------|81| `T.set_flag / T.wait_flag` | 核内流水线同步 |82| `T.barrier_all() / T.pipe_barrier(pipe)` | 管线屏障 |83| `T.set_cross_flag / T.wait_cross_flag` | 核间同步 |84| `T.sync_all()` | 全局同步 |85| `T.printf(fmt, *args)` | 设备端格式化打印 |86| `T.dump_tensor(tensor, desc, size, shape_info)` | Tensor 转储 |8788### 常用 pass_configs8990| 配置项 | 说明 |91|-------|------|92| `TL_ASCEND_AUTO_SYNC: True` | 自动同步插入 |93| `TL_ASCEND_MEMORY_PLANNING: True` | 自动内存规划 |94| `TL_ASCEND_AUTO_CV_COMBINE: True` | 自动 CV 分离（核间流水线） |95| `TL_ASCEND_AUTO_CV_SYNC: True` | 自动核间同步（核间流水线） |
tile-ai