企业官网建设流程全解析

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_stages=0表示编译器推断流水（默认不启用），而手动调度则优先只给order/stage（流水深度按max(stage) + 1推断）。

一个关键推论是：不带任何上述标记的裸T.Pipelined(n)在结构上等价于串行循环，不属于 pipeline-requested，因此不会触发流水相关的语义限制。

源表示：各循环构造在源 TIR 中的身份

"始终在预期检查阶段检查表示"是文档反复强调的原则——Python 层看起来理所当然的循环结构，经过前端展开后可能面目全非。下表是各循环构造到源 TIR 的映射：

源码 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.py：nested_continuous_parallels（双层连续并行）与nested_triple_continuous_parallels（三层连续并行）均能正常编译并通过数值对比，而nested_noncontinuous_parallels因在两层并行之间插入了B[i] = 0语句，测试用pytest.raises(ValueError)断言其必然报错。

已确立规则二：管道可包含并行，并行不可包含管道

允许经典的 tile 循环形态：

for k in T.Pipelined(K, num_stages=3): 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）与后端分派，而非全局拒绝。

审查嵌套管道时应遵循五步流程：

1. 解析选定的 target 与后端 pipeline；
2. 判断该后端是否支持分层流水规划、buffer versioning、barrier 与 warp/core specialization；
3. 后端契约支持时允许嵌套；
4. 不支持时，在 target 解析之后于后端内部诊断：Backend <name> does not support nested software pipelines；
5. 若丢弃某个请求的调度可能改变异步或多缓冲语义，绝不能静默丢弃。

源扫描器保持"仅信息"性质：它应能报告如下形状的代码并正常退出，而不做 target 能否 lowering 的裁决：

for ko in T.Pipelined(K, num_stages=3): for ki in T.Pipelined(BK, num_stages=2): ...

需要覆盖的边界矩阵包括：支持的后端接受嵌套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_stages=2)的 GEMM 内核，test_nested_pipelines编译并通过数值对比（atol=1e-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 访问可以合法；
• "没有示例使用这种形式，因此它非法"——必须确认下游存在正确性需求才能下结论。

这些"非规则"与该文档的规则设计哲学一脉相承：规则必须建立在具体的正确性机制之上，而不是建立在"没人这么写"或"当前实现做不到"之上。

规则设计检查清单：新增循环规则前的十个自问

在批准任何新的循环规则之前，文档要求回答以下十个问题，这也是读者审查、贡献循环语义时的标准模板：

1. 这属于语义非法、当前实现契约，还是优化条件？
2. 在检查器阶段，确切的 TIR 形状是什么来识别该构造？
3. 最小的非法示例是什么？
4. 最接近的合法示例是什么？
5. 兄弟循环与嵌套祖先循环是否不同？
6. 中性包装（serial、unroll、if、SeqStmt）是否会改变答案？
7. 生成的 IR 是否使用相同形状并需要豁免？
8. 检查器能证明违规，还是只能未能证明安全？
9. 哪个 pass 最先依赖该不变量？
10. 诊断信息是否给出了建议的合法改写方式？

检查器如何接入编译流程

上述检查并非孤立存在。在 semantic_check.py 中，PreLowerSemanticCheck会在 lowering 之前依次运行三个后端无关的检查器：NestedLoopChecker、ParallelLocalIndexChecker、FragmentLoopChecker，并可通过TL_DISABLE_PRELOWER_SEMANTIC_CHECK配置关闭、通过TL_AST_PRINT_ENABLE开启 AST 打印。三个检查器均为prim_func_pass（opt_level=0），在 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