【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skillsname: 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-practicesskill本文档不再重复。1. 模式对比维度Developer 模式Expert 模式内存分配T.alloc_shared/T.alloc_fragmentT.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_flagpass_configs全部开启全部关闭或不设适用场景大多数算子跨平台兼容极致性能优化需要底层控制示例目录examples/developer_mode/examples/flash_attention/fa_opt/flash_attn_bhsd_expert_*.py混合模式Developer 主体 少量 Expert。使用 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自动核内同步底层 keytl.ascend_auto_sync默认 False功能自动在数据搬运和计算之间插入T.barrier_all()等同步指令开启时无需手写T.barrier_all()、T.set_flag/T.wait_flag关闭时必须手动插入所有同步点② TL_ASCEND_MEMORY_PLANNING自动内存规划底层 keytl.ascend_memory_planning默认 False功能自动分析 buffer 生命周期实现片上内存复用开启时自动复用 buffer 空间减少片上内存占用关闭时需手动通过T.annotate_address规划内存地址③ TL_ASCEND_AUTO_CV_COMBINE自动 CV 分离底层 keytl.ascend_auto_cv_combine默认 False功能自动将 kernel 中的 Cube 操作和 Vector 操作分离到不同的执行核开启时无需手写with T.Scope(C)/with T.Scope(V)关闭时必须手动用T.Scope标注每段代码的执行域④ TL_ASCEND_AUTO_CV_SYNC自动核间同步底层 keytl.ascend_auto_cross_core_sync默认 False功能自动在 Cube Scope 和 Vector Scope 之间插入T.set_cross_flag/T.wait_cross_flag开启时无需手写核间同步关闭时必须手动管理核间同步2.2 按场景选择 pass_configs场景AUTO_SYNCMEMORY_PLANNINGAUTO_CV_COMBINEAUTO_CV_SYNC纯 Vector 算子elementwise, softmax✅✅不需要不需要Developer GEMM✅✅✅✅Developer Flash Attention核间流水线✅视情况✅✅Expert 极致性能❌❌❌❌混合模式✅✅✅✅纯 Vector 算子来自 Programming Guide §2.2pass_configs { tilelang.PassConfigKey.TL_ASCEND_AUTO_SYNC: True, tilelang.PassConfigKey.TL_ASCEND_MEMORY_PLANNING: True, }Developer GEMMpass_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 → Developer3.1 转换步骤开启 pass_configs添加完整 4 个 True 开关内存分配T.alloc_L1→T.alloc_sharedT.alloc_L0C→T.alloc_fragmentT.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_address3.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_PLANNING4. 实际代码对比完整的 Developer vs Expert 代码对比请参考→ mode-examples.md5. 示例代码位置模式目录Developerexamples/developer_mode/Expertexamples/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纯 Vectorexamples/elementwise/、examples/softmax/【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills创作声明:本文部分内容由AI辅助生成(AIGC),仅供参考