| name | tilelang-programming-model-guide |
| description | TileLang Ascend Developer/Expert 模式选择与 pass_configs 配置指南。当需要确定编程模式、配置 pass_configs、或在两种模式之间转换时触发。API 详情请参考 tilelang-api-best-practices skill。 |
TileLang Ascend 编程模式与 pass_configs 指南
API 用法详情（内存分配、计算原语、同步原语等）请参考 tilelang-api-best-practices skill，本文档不再重复。
1. 模式对比
| 维度 | Developer 模式 | Expert 模式 |
|---|
| 内存分配 | T.alloc_shared / T.alloc_fragment | T.alloc_L1 / T.alloc_ub / T.alloc_L0A/L0B/L0C |
| 计算表达 | T.Parallel + 符号运算 | T.tile.xxx 扩展原语 |
| 作用域 | 编译器自动分离 Cube/Vector | 手动 with T.Scope("C"/"V") |
| 同步 | 编译器自动插入 | 手动 T.barrier_all / T.set_flag / T.wait_flag |
| CV 交互 | 默认消除 workspace+vid（threads=2 + 片上直连，见 §3.1.1） | 显式 GM workspace + 手动 vid 二分 |
| pass_configs | 全部开启 | 全部关闭或不设 |
| 适用场景 | 大多数算子，跨平台兼容 | 极致性能优化，需要底层控制 |
| 示例目录 | examples/developer_mode/ | examples/flash_attention/fa_opt/flash_attn_bhsd_expert_*.py |
混合模式：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,
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
| 场景 | AUTO_SYNC | MEMORY_PLANNING | AUTO_CV_COMBINE | AUTO_CV_SYNC | 手动 Scope |
|---|
| 纯 Vector 算子（elementwise, softmax） | ✅ | ✅ | ❌ | ❌ | ❌ |
| Developer GEMM（完全自动） | ✅ | ✅ | ✅ | ✅ | ❌ |
| Developer Flash Attention（核间流水线） | ✅ | ✅ | ✅ | ✅ | ❌ |
| Developer CV 融合（Vector计算+Cube GEMM） | ✅ | ✅ | ✅ | ✅ | ❌ |
| 混合模式 CV 融合 | ✅ | ✅ | ❌ | ❌ | ✅ |
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,
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 转换步骤
- 开启 pass_configs：添加完整 4 个 True 开关
- 内存分配：
T.alloc_L1 → T.alloc_shared，T.alloc_L0C → T.alloc_fragment，T.alloc_ub → T.alloc_shared
- 删除作用域：移除
with T.Scope("C") / with T.Scope("V")
- 删除同步：移除
T.barrier_all()、T.set_flag/T.wait_flag、T.set_cross_flag/T.wait_cross_flag
- 计算转换（可选）：
T.tile.exp(dst, src) → for i,j in T.Parallel(...): dst[i,j] = T.exp(src[i,j])
- 删除手动内存规划：移除
T.annotate_address
3.1.1 Developer 模式 CV 交互：优先消除 workspace / vid
Developer 模式下 Cube↔Vector 交互默认不写 GM workspace、不手动二分 vid，交给编译器自动处理。前提链（按序，不可跳级）：
threads=2 ──► vid 消除 ──► workspace 消除
四步改造：
- 加
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 的前提）。
- 删
workspace_idx：@tilelang.jit(out_idx=[N], workspace_idx=[...], ...) → @tilelang.jit(out_idx=[N], ...)，并删除 kernel 签名里的 workspace_* 参数。
- 去 vid 偏移：
v_block 不再 // 2，循环恢复整程 range(BI)，删除全部 vid * ... 索引偏移。
- 片上直连：原「片上 buffer ↔
workspace[cid,...] ↔ 另一片上 buffer」两跳 GM 往返，合并为片上 T.copy 一跳；中转/同步交给四个 pass。
完整模板、映射表、代码骨架、自检清单与回退条件见 mode-examples.md §6。
复杂同步/多版本流水场景可回退保留 workspace+vid 写法（见 mode-examples.md §7）。
3.2 转换对照表
| Expert 写法 | Developer 写法 |
|---|
T.alloc_L1(shape, dtype) | T.alloc_shared(shape, dtype) |
T.alloc_ub(shape, dtype) | T.alloc_shared(shape, dtype) |
T.alloc_L0A/L0B(shape, dtype) | 删除（gemm_v0 内部处理） |
T.alloc_L0C(shape, dtype) | T.alloc_fragment(shape, dtype) |
with T.Scope("C"): ... | 直接写代码（编译器自动分离） |
T.barrier_all() | 删除（编译器自动插入） |
T.set_flag/T.wait_flag(...) | 删除 |
T.set_cross_flag/T.wait_cross_flag(...) | 删除 |
T.tile.exp(dst, src) | for i,j in T.Parallel(...): dst[i,j] = T.exp(src[i,j]) 或保留 |
T.annotate_address({...}) | 删除（开启 MEMORY_PLANNING） |
@jit(..., workspace_idx=[...]) + 签名 workspace_* 参数 | 删除（CV 交互改片上直连，见 §3.1.1） |
T.Kernel(..., is_npu=True) as (cid, vid) | T.Kernel(..., threads=2, is_npu=True) as (cid)（消 vid 前提） |
T.copy(buf, ws[cid,...]) + T.copy(ws[cid,vid*..], buf2) 两跳 | T.copy(buf, buf2) 片上一跳直连 |
4. 示例代码与代码对比
| 模式 | 目录 | 说明 |
|---|
| Developer | examples/developer_mode/ | GEMM、elementwise 等 |
| Developer（消除 workspace/vid） | examples/developer_mode/sparse_flash_attn_developer_vid_reduce.py（新）vs sparse_flash_attn_developer.py（旧） | threads=2 + 片上直连，逐行对照消除范式 |
| Expert | examples/gemm/example_gemm_intrinsic.py、examples/flash_attention/fa_opt/flash_attn_bhsd_expert_*.py | 极致性能优化 |
| 混合（核间流水线） | examples/flash_attention/flash_attn_bhsd_cc_sync.py、examples/flash_attention/fa_opt/flash_attn_bhsd_auto_pipeline_*.py | FA 核间流水线 |
| 纯 Vector | examples/elementwise/、examples/softmax/ | 无 Cube 操作 |
| CV 融合 | examples/dequantize_gemm/、examples/quant_batch_matmul/ | Vector 计算 + Cube GEMM |
完整代码对比（Developer vs Expert）：
- → mode-examples.md
- 包含 GEMM、Flash Attention、Softmax、CV 融合（消除 workspace/vid 推荐写法 §6 / workspace+vid 兜底写法 §7） 等示例
