TileLang 循环语义规则全解析从嵌套并行、管道到向量化的合法性边界【免费下载链接】tilelangDomain-specific language designed to streamline the development of high-performance GPU/CPU/Accelerators kernels项目地址: https://gitcode.com/GitHub_Trending/ti/tilelang导读TileLang 是面向高性能 GPU/CPU/加速器内核开发的领域专用语言其循环结构T.Parallel、T.Pipelined、T.vectorized、T.serial等承担了并行分布、软件流水与向量化等关键语义。本文基于仓库内.agents/skills/tilelang-semantic/references/loop-rules.md语义规范文档系统梳理 TileLang 循环嵌套的已确立规则与非规则并结合 nested_loop_checker.py、parallel_local_index_checker.py、fragment_loop_checker.py 等检查器源码及其测试用例帮助读者在编写内核时正确使用循环结构、理解编译器报错信息并学会区分语义错误、当前实现契约与优化条件三类不同性质的限制。范围与术语先分清三类规则性质在深入任何一条循环规则之前首先要明确规则的适用范围与性质。TileLang 语义文档要求一条嵌套规则只作用于一个词法路径lexical path——即从某个函数体到某一条叶语句的祖先链两个顺序排列的兄弟循环属于不同路径不受同一条嵌套规则约束。更关键的是规则被明确区分为三个层次语义规则semantic rule在所有受支持的 lowering 下保证正确性所必需的规则违反即编译错误实现契约implementation contract当前编译 pass 所要求、但未来可能放宽的约束优化条件optimization condition仅决定某个优化是否生效绝不能被当作语义非法来报错。什么是 pipeline-requested 循环文档给出了一个精确的判定术语当源 TIR 中携带num_stages或手动的tl_pipeline_order/tl_pipeline_stage元数据时该T.Pipelined循环被称为pipeline-requested。在规划planning之后规范化标记还包括tl_pipelined_num_stages以及软件流水 stage/order 属性。而辅助性的group或sync元数据本身不会触发流水 lowering。这一点与 loop.py 中T.Pipelined的 Python 接口实现完全对应num_stages、order、stage、sync、group都会作为参数传入_ffi_api.Pipelined其中num_stages0表示编译器推断流水默认不启用而手动调度则优先只给order/stage流水深度按max(stage) 1推断。一个关键推论是不带任何上述标记的裸T.Pipelined(n)在结构上等价于串行循环不属于 pipeline-requested因此不会触发流水相关的语义限制。源表示各循环构造在源 TIR 中的身份始终在预期检查阶段检查表示是文档反复强调的原则——Python 层看起来理所当然的循环结构经过前端展开后可能面目全非。下表是各循环构造到源 TIR 的映射源构造源 TIR 身份重要后果T.serial,T.grid串行For嵌套视为普通顺序迭代T.unrollForKind::kUnrolled编译期展开意图内部可以包含更低层循环T.vectorizedForKind::kVectorizedSIMD 意图当前规划器只选择最内层循环T.Parallel一个或多个ForKind::kParallel循环连续多个循环构成一个多维并行区域T.Pipelined串行For加注释通过注释识别而非ForKindT.Persistentbinds、guards、loop_break与串行For在强制 persistent 特定结构前保留标记whileWhile在真实内核中常充当动态/persistent 调度器源码 nested_loop_checker.py 印证了通过注释识别流水这一点is_pipelined_for检查的注解键为num_stages、tl_pipeline_order、tl_pipeline_stage、tl_pipeline_group与文档中的定义完全一致而不是通过ForKind判断。已确立规则一连续并行维度构成一个并行区域规则允许严格的连续并行链for i in T.Parallel(M): for j in T.Parallel(N): B[i, j] A[i, j]而一旦可执行语句或其他循环打断了链再进入T.Parallel就是非法for i in T.Parallel(M): B[i, 0] 0 for j in T.Parallel(N): B[i, j] A[i, j]背后的原理是连续的并行循环可以融合fuse为单个多维T.Parallel区域由并行 lowering 统一做线程分布与布局推断而中间插入语句后并行链断裂无法再视为单一区域。当前实现位于 nested_loop_checker.py_NestedLoopCheckVisitor在遇到ForKind.PARALLEL循环时若其直接子节点是另一个PARALLEL循环则递归放行否则若已处于并行上下文就抛出ValueError([Tilelang Semantic Check] Nested parallel loops are not allowed. Please check your loop structure.)。对应测试见 test_tilelang_nested_loop_checker.pynested_continuous_parallels双层连续并行与nested_triple_continuous_parallels三层连续并行均能正常编译并通过数值对比而nested_noncontinuous_parallels因在两层并行之间插入了B[i] 0语句测试用pytest.raises(ValueError)断言其必然报错。已确立规则二管道可包含并行并行不可包含管道允许经典的 tile 循环形态for k in T.Pipelined(K, num_stages3): for i in T.Parallel(M): ...但禁止在并行区域内出现 pipeline-requested 的循环。原因在于两者 lowering 的职责完全不同并行 lowering 负责把逐元素工作分布到线程而软件流水规划拥有串行的生产者/消费者时间线无法被引入并行区域内部。nested_loop_checker.py 中对应逻辑为当is_pipelined_for(op)为真且当前处于并行上下文时抛出[Tilelang Semantic Check] Pipelined loop cannot be nested inside a parallel loop. Please check your loop structure.。注意该判定发生在进入循环体之前且只针对 pipeline-requested 循环——裸T.Pipelined因无注解不会被拦截。已确立规则三tile 算子不得出现在并行区域内T.copy、T.gemm等通过TLOpBuilder注册的 tile 算子禁止出现在T.Parallel内部而逐元素的原子操作、reducer 更新等内置指令intrinsics则被允许不要把所有带副作用的调用都归类为 tile 算子。实现上 nested_loop_checker.py 的is_tile_op通过op.op.get_attr(TLOpBuilder) is not None精确判定——tl.reducer_update等逐迭代内置函数是普通 builtin不带TLOpBuilder属性因此天然通过检查。visit_call_在并行上下文中遇到 tile op 时抛出[Tilelang Semantic Check] Only elementwise operations are allowed inside a parallel loop. Got a tile-op ...。已确立规则四并行索引必须遵循存储所有权并行区域内的 buffer 索引要遵循存储所有权约束禁止直接用外层并行变量索引 thread-private 的 local buffer例如T.alloc_local分配的区域允许与并行变量无关的局部访问如复制的标量读取replicated scalar reads当被索引的逻辑维度需要跨线程分布时应改用 fragment 存储T.alloc_fragment保留 fragment 索引现有的符号范围限制。这三条分别落在 parallel_local_index_checker.py 与 fragment_loop_checker.py 中parallel_local_index_checker.py维护parallel_loop_stack对BufferLoad/BufferStore检查若 buffer 是 local 且索引表达式中出现了任何并行循环变量则抛出[Tilelang Semantic Check] Local buffer ... is indexed by T.Parallel loop variable ...并给出修复建议——用T.serial/T.vectorized/T.unroll做逐线程局部索引或用T.alloc_fragment让被索引维度跨线程分布。fragment_loop_checker.py则在到达最内层循环后收集所有 fragment 访问检查路径上是否存在**符号范围min/extent 非IntImm**的并行循环若其循环变量用于索引 fragment 则报错。文档中保留符号范围限制正对应此实现。嵌套管道是后端能力而非全局语义规则文档明确禁止对单条词法路径上的 pipeline-requested 循环数量施加语言级上限。分层软件流水hierarchical software pipelines在语义上有意义——已知的 Ascend 内核会同时对外层 tile 循环与内层 GEMM 归约循环做流水。要点在于某个后端 pipeline planning 或 multi-versioning pass 里的注释只代表该后端当前的实现限制不能作为跨后端的PreLowerSemanticCheck该检查运行在目标后端选定 lowering 策略之前拒绝嵌套流水的依据。pipeline-requested分类应用于盘点inventory与后端分派而非全局拒绝。审查嵌套管道时应遵循五步流程解析选定的 target 与后端 pipeline判断该后端是否支持分层流水规划、buffer versioning、barrier 与 warp/core specialization后端契约支持时允许嵌套不支持时在 target 解析之后于后端内部诊断Backend name does not support nested software pipelines若丢弃某个请求的调度可能改变异步或多缓冲语义绝不能静默丢弃。源扫描器保持仅信息性质它应能报告如下形状的代码并正常退出而不做 target 能否 lowering 的裁决for ko in T.Pipelined(K, num_stages3): for ki in T.Pipelined(BK, num_stages2): ...需要覆盖的边界矩阵包括支持的后端接受嵌套num_stages与手动 stage/order 调度不支持的后端对同一份源码给出 target 专属诊断target 无关的 pre-lower 校验对两种情况都通过裸与 pipeline-requested 的嵌套循环在分析中保持可区分顺序兄弟管道相互独立。仓库测试验证了这一立场test_tilelang_nested_loop_checker.py 中的matmul_nested_pipelines构造了一个外层T.Pipelined(extra_pipeline_repeats)无 num_stages包裹内层T.Pipelined(T.ceildiv(K, block_K), num_stages2)的 GEMM 内核test_nested_pipelines编译并通过数值对比atol1e-2。这与文档中已知内核会同时流水外层 tile 循环与内层 GEMM 归约循环的论断互为印证。向量化循环不要求是叶节点不要引入T.vectorized内部不得包含循环的 blanket 规则。当内层循环的边界与 SIMD lane 无关lane-invariant bounds时嵌套是有意义的for i in T.vectorized(M): for k in T.serial(K): C[i] A[i, k] * B[k]串行循环会为每条 SIMD lane 执行T.unroll同理可行。但如果内层边界依赖i各 lane 的 trip count 可能不同此时需要 predication 或标量化应作为独立的规则另行处理而非一概禁止。当前的实现注意点VectorizePlanner只规划最内层循环见 loop_vectorize.cc 中Must analysis vectorization on the innermost loop的注释与 L1052 处的改写逻辑。因此外层T.vectorized包裹T.serial时可能被降级为串行并给出警告——这是优化限制不是语义非法。在修改该行为前既要测试语义接受性也要测试向量化是否真的发生。非规则六条不应被采纳的伪规则以下陈述容易在代码评审中被误当成通用语义规则但文档明确否定每个嵌套T.Pipelined都是非法的或一条词法路径最多一个 pipeline-requested 循环——嵌套管道是后端能力裸外层管道语法也呈串行特性T.vectorized必须是 AST 叶子——顺序内层循环可以是合法的T.Parallel的范围必须等于启动的线程数——循环划分partitioning与复制replication刻意支持不同范围显式 loop-layout 输入形状必须等于循环范围——带守卫的尾部guarded tails与非双射布局可以是有意为之T.Parallel内的每个局部访问都是非法的——只有破坏所有权的依赖被禁止复制的或逐迭代的 scratch 访问可以合法没有示例使用这种形式因此它非法——必须确认下游存在正确性需求才能下结论。这些非规则与该文档的规则设计哲学一脉相承规则必须建立在具体的正确性机制之上而不是建立在没人这么写或当前实现做不到之上。规则设计检查清单新增循环规则前的十个自问在批准任何新的循环规则之前文档要求回答以下十个问题这也是读者审查、贡献循环语义时的标准模板这属于语义非法、当前实现契约还是优化条件在检查器阶段确切的 TIR 形状是什么来识别该构造最小的非法示例是什么最接近的合法示例是什么兄弟循环与嵌套祖先循环是否不同中性包装serial、unroll、if、SeqStmt是否会改变答案生成的 IR 是否使用相同形状并需要豁免检查器能证明违规还是只能未能证明安全哪个 pass 最先依赖该不变量诊断信息是否给出了建议的合法改写方式检查器如何接入编译流程上述检查并非孤立存在。在 semantic_check.py 中PreLowerSemanticCheck会在 lowering 之前依次运行三个后端无关的检查器NestedLoopChecker、ParallelLocalIndexChecker、FragmentLoopChecker并可通过TL_DISABLE_PRELOWER_SEMANTIC_CHECK配置关闭、通过TL_AST_PRINT_ENABLE开启 AST 打印。三个检查器均为prim_func_passopt_level0在 analysis/init.py 中统一导出这正是target 无关的 pre-lower 校验在编译管线中的落点——也解释了为什么嵌套管道是否支持必须留给后端在 target 解析后裁决而不是在这个阶段被全局拒绝。小结用规则思维写 TileLang 内核综合本文所述TileLang 循环语义的核心可以浓缩为三句话连续T.Parallel可融合为一个并行区域但区域内禁止 tile 算子与管道T.Pipelined可包含并行且允许按后端能力嵌套但不可被并行包裹T.vectorized不必是叶子T.serial可以在其中合法迭代。写作内核时遵循文档的建议——优先用T.Parallel在T.gemm级 tile 算子内部实现 tiled 运算而非其他用途——即可避开绝大多数嵌套问题。遇到编译报错时先判断它是语义规则、实现契约还是优化条件再依据loop-rules.md的检查清单定位最小的非法示例与合法的改写路径就能快速、准确地修复内核代码。【免费下载链接】tilelangDomain-specific language designed to streamline the development of high-performance GPU/CPU/Accelerators kernels项目地址: https://gitcode.com/GitHub_Trending/ti/tilelang创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考