tilelang-mode-guide

TileLang Ascend 编程模式与 pass_configs 指南

API 用法详情（内存分配、计算原语、同步原语等）请参考 tilelang-api-best-practices skill，本文档不再重复。

---

1. 模式对比

混合模式：Developer 主体 + 少量 Expert / Ascend 专属 T.tile.xxx。使用 Developer 的 pass_configs，不写 T.Scope 和手动同步。大多数实际算子使用混合模式。

---

2. pass_configs 详解（核心）

2.1 四个 Ascend 专用开关

import tilelang

pass_configs = {
    tilelang.PassConfigKey.TL_ASCEND_AUTO_SYNC: True,        # ① 自动核内同步
    tilelang.PassConfigKey.TL_ASCEND_MEMORY_PLANNING: True,   # ② 自动内存规划
    tilelang.PassConfigKey.TL_ASCEND_AUTO_CV_COMBINE: True,   # ③ 自动CV分离
    tilelang.PassConfigKey.TL_ASCEND_AUTO_CV_SYNC: True,      # ④ 自动核间同步
}

① TL_ASCEND_AUTO_SYNC（自动核内同步）

• 底层 key："tl.ascend_auto_sync"，默认 False
• 功能：自动在数据搬运和计算之间插入 T.barrier_all() 等同步指令
• 开启时：无需手写 T.barrier_all()、T.set_flag/T.wait_flag
• 关闭时：必须手动插入所有同步点

② TL_ASCEND_MEMORY_PLANNING（自动内存规划）

• 底层 key："tl.ascend_memory_planning"，默认 False
• 功能：自动分析 buffer 生命周期，实现片上内存复用
• 开启时：自动复用 buffer 空间，减少片上内存占用
• 关闭时：需手动通过 T.annotate_address 规划内存地址

③ TL_ASCEND_AUTO_CV_COMBINE（自动 CV 分离）

• 底层 key："tl.ascend_auto_cv_combine"，默认 False
• 功能：自动将 kernel 中的 Cube 操作和 Vector 操作分离到不同的执行核
• 开启时：无需手写 with T.Scope("C") / with T.Scope("V")，编译器根据 buffer 类型和所用原语自动识别
• 关闭时：必须手动用 T.Scope 标注每段代码的执行域

注意：避免在开启 AUTO_CV_COMBINE 同时手写 T.Scope，可能会导致编译器无法正确处理代码

④ TL_ASCEND_AUTO_CV_SYNC（自动核间同步）

• 底层 key："tl.ascend_auto_cross_core_sync"，默认 False
• 功能：自动在 Cube Scope 和 Vector Scope 之间插入 T.set_cross_flag/T.wait_cross_flag
• 开启时：无需手写核间同步
• 关闭时：必须手动管理核间同步

2.2 按场景选择 pass_configs

Developer Flash Attention / Developer CV 融合：默认消除 workspace+vid（threads=2 + 片上直连），写法见 §3.1.1 与 mode-examples.md §6。 | Expert 极致性能 | ❌ | ❌ | ❌ | ❌ | ✅ |

纯 Vector 算子（来自 Programming Guide §2.2）：

pass_configs = {
    tilelang.PassConfigKey.TL_ASCEND_AUTO_SYNC: True,
    tilelang.PassConfigKey.TL_ASCEND_MEMORY_PLANNING: True,
}

Developer GEMM / Developer CV 融合（推荐配置）：

pass_configs = {
    tilelang.PassConfigKey.TL_ASCEND_AUTO_CV_COMBINE: True,  # 自动分离 Cube/Vector
    tilelang.PassConfigKey.TL_ASCEND_AUTO_SYNC: True,        # 自动核内同步
    tilelang.PassConfigKey.TL_ASCEND_MEMORY_PLANNING: True,  # 自动内存规划
    tilelang.PassConfigKey.TL_ASCEND_AUTO_CV_SYNC: True,     # 自动核间同步
}

Expert 全手动：

pass_configs = {
    tilelang.PassConfigKey.TL_ASCEND_AUTO_CV_COMBINE: False,
    tilelang.PassConfigKey.TL_ASCEND_AUTO_CV_SYNC: False,
    tilelang.PassConfigKey.TL_ASCEND_MEMORY_PLANNING: False,
    tilelang.PassConfigKey.TL_ASCEND_AUTO_SYNC: False,
}

---

3. 模式转换规则（Expert → Developer）

3.1 转换步骤

1. 开启 pass_configs：添加完整 4 个 True 开关
2. 内存分配：T.alloc_L1 → T.alloc_shared，T.alloc_L0C → T.alloc_fragment，T.alloc_ub → T.alloc_shared
3. 删除作用域：移除 with T.Scope("C") / with T.Scope("V")
4. 删除同步：移除 T.barrier_all()、T.set_flag/T.wait_flag、T.set_cross_flag/T.wait_cross_flag
5. 计算转换（可选）：T.tile.exp(dst, src) → for i,j in T.Parallel(...): dst[i,j] = T.exp(src[i,j])
6. 删除手动内存规划：移除 T.annotate_address

3.1.1 Developer 模式 CV 交互：优先消除 workspace / vid

Developer 模式下 Cube↔Vector 交互默认不写 GM workspace、不手动二分 vid，交给编译器自动处理。前提链（按序，不可跳级）：

threads=2  ──►  vid 消除  ──►  workspace 消除

四步改造：

1. 加 threads=2：T.Kernel(block_num, is_npu=True) as (cid, vid) → T.Kernel(block_num, threads=2, is_npu=True) as (cid)（编译器自动并行 2 个 V 核，这是消 vid 的前提）。
2. 删 workspace_idx：@tilelang.jit(out_idx=[N], workspace_idx=[...], ...) → @tilelang.jit(out_idx=[N], ...)，并删除 kernel 签名里的 workspace_* 参数。
3. 去 vid 偏移：v_block 不再 // 2，循环恢复整程 range(BI)，删除全部 vid * ... 索引偏移。
4. 片上直连：原「片上 buffer ↔ workspace[cid,...] ↔ 另一片上 buffer」两跳 GM 往返，合并为片上 T.copy 一跳；中转/同步交给四个 pass。

完整模板、映射表、代码骨架、自检清单与回退条件见 mode-examples.md §6。 复杂同步/多版本流水场景可回退保留 workspace+vid 写法（见 mode-examples.md §7）。

3.2 转换对照表

---

4. 示例代码与代码对比

完整代码对比（Developer vs Expert）：

• → mode-examples.md
• 包含 GEMM、Flash Attention、Softmax、CV 融合（消除 workspace/vid 推荐写法 §6 / workspace+vid 兜底写法 §7） 等示例