1. 算子编程这件事，为什么突然需要一门新语言
第一次听到 TileLang 这个名字，是在一个做推理优化的群里。有人丢了一句“国产算子编程语言”，底下立刻分成两拨：一拨问“是不是又一个 DSL 套壳”，另一拨直接问“能不能跑在国产卡上”。这种反应其实很真实——算子编程这个领域，过去十年基本被 CUDA C++ 和 Triton 轮流占着，突然冒出一个国产名字，大家第一反应是警惕，第二反应才是好奇。
TileLang 要解决的问题，说白了就一句话：让写高性能算子这件事，从“手搓汇编级优化”变成“描述计算意图”。你如果写过 CUDA 算子就知道，一个像样的 FlashAttention 或者 GEMM，核心逻辑可能就几十行，但为了把性能压到硬件极限，你得手动管 shared memory、管寄存器分配、管 bank conflict、管流水线重叠，最后代码膨胀到几百行，换一张卡就得重写一遍。TileLang 的思路是把这些脏活收进编译器，让开发者用更接近数学表达的方式描述“我要算什么”，至于怎么切 tile、怎么排布内存、怎么调度流水线，交给编译后端去决定。
它适合谁？三类人最该关注。第一类是算子库开发者，天天跟 cuBLAS、cuDNN 打交道，想自己写定制算子又不想被 CUDA 绑死；第二类是国产芯片适配团队，手里有卡但缺算子生态，急需一个能快速把模型跑起来的中间层；第三类是编译器方向的研究者，想看看 tile-level 抽象到底能做到什么程度。如果你只是调包侠，平时torch.nn够用，那 TileLang 暂时跟你关系不大，但理解它的设计思路，对你判断未来算子生态往哪走很有帮助。
我自己的判断是：TileLang 这类东西的价值不在“替代 CUDA”，而在把算子开发的门槛从“硬件专家”降到“懂计算的工程师”。这个降维一旦成立，国产开源生态里最缺的那块拼图——高性能算子供给——就有机会被补上。下面我按自己实际折腾的顺序，把这门语言的核心设计、实操要点、踩坑记录和对生态的影响，一层层拆开讲。
2. TileLang 的核心设计思路拆解
2.1 为什么是 tile-level 抽象，而不是 instruction-level
要理解 TileLang 为什么长这样，得先看它想解决什么层次的痛点。CUDA 是 instruction-level 的，你写的是线程级代码，threadIdx.x怎么走、__syncthreads()放哪，全是你的事。Triton 往上走了一层，变成 block-level，你描述的是“一个 program 处理一块数据”，但 block 内部的 layout 和 swizzle 还是得自己操心。TileLang 再往上走一层，到了tile-level：你描述的是“把 A 的这块 tile 和 B 的这块 tile 做乘加，结果写到 C 的这块 tile”，至于这块 tile 怎么映射到线程、怎么放 shared memory、怎么打流水，编译器负责。
这个抽象层级的选择不是拍脑袋。我试过用 Triton 写一个带 mask 的 attention，光是处理边界 tile 的tl.where和tl.load的 mask 参数就调了半天，因为 block 内部的 layout 你得心里有数。TileLang 把 tile 当成一等公民之后，边界处理、layout 转换这些事被收进类型系统里，写起来确实更接近“我想算什么”而不是“我怎么算”。
注意：tile-level 抽象不是银弹。当你需要极致的手工调优，比如针对某张卡的 L2 cache 做特殊分块，tile-level 的表达力可能不够，这时候还是得往下钻。TileLang 一般会留 escape hatch，但用多了就失去抽象的意义了。
2.2 计算与调度分离：这门语言最值钱的地方
TileLang 设计里我觉得最值钱的一点，是计算描述和调度策略的分离。你写算子的时候，先定义计算逻辑——输入输出是什么、做什么运算、tile 形状多大；然后单独写调度——怎么分块、怎么流水、用什么内存层级。这两部分解耦之后，同一份计算逻辑可以配不同的调度，针对不同硬件快速试。
这个思路其实借鉴了 TVM 的 schedule 概念，但 TileLang 把它做得更“算子友好”。TVM 的 schedule 原语偏底层，写起来像在写编译器 pass；TileLang 的调度描述更接近硬件工程师的直觉，比如“这个循环做 pipeline”“这个 buffer 放 shared”“这个 tile 做 swizzle”。我实测下来，同一个 GEMM 计算逻辑，换三套调度分别适配不同显存带宽的卡，改动量大概只有十几行，这在 CUDA 时代是不可想象的。
为什么这个分离重要？因为硬件碎片化是国产生态的常态。你不可能为每张卡重写一遍算子，但你可以为每张卡写一套调度。计算逻辑复用，调度按卡定制，这才是可持续的算子供给模式。
2.3 编译后端的选择逻辑
TileLang 往下编译，最终要落到具体硬件上。它的后端设计我理解是分层的：上层是 tile-level IR，中间经过若干轮 lowering，最后生成目标代码。目标可以是 CUDA C、可以是某类国产卡的编程接口、也可以是更底层的 IR。这个分层的好处是新增一个硬件后端，不用动前端语言。
我踩过的一个坑是：早期版本里某些 tile 操作在 lowering 到特定后端时会失败，报错信息很含糊。后来发现是那个后端的某个 intrinsic 没实现。这提醒我，评估这类语言时，后端成熟度比前端语法重要得多。前端再漂亮，后端跑不通就是零。选型的时候一定要拿你真正要部署的那张卡去跑 benchmark，别只看文档里支持的硬件列表。
3. 上手实操：从零写一个 TileLang 算子
3.1 环境准备与最小可运行示例
假设你已经装好了基础工具链，TileLang 一般以 Python 包的形式提供，因为它前端是嵌在 Python 里的 DSL。安装方式通常是 pip，但具体包名和版本得看官方仓库，我这里不写死，你按官方 README 来。装完之后，第一件事是跑通一个最小示例，确认后端能编译、能执行。
我建议的第一个算子不要选 GEMM，选element-wise 加法。原因很简单：加法没有 tile 间的数据依赖，能让你专注理解“怎么定义 tile、怎么读写、怎么启动”。一个典型的骨架大概是这样：
import tilelang import tilelang.language as T @tilelang.jit def add_kernel(M, N, block_M, block_N): @T.prim_func def main(A: T.Tensor((M, N), "float16"), B: T.Tensor((M, N), "float16"), C: T.Tensor((M, N), "float16")): with T.Kernel(T.ceildiv(N, block_N), T.ceildiv(M, block_M)) as (bx, by): for i, j in T.Parallel(block_M, block_N): C[by * block_M + i, bx * block_N + j] = ( A[by * block_M + i, bx * block_N + j] + B[by * block_M + i, bx * block_N + j] ) return main这段代码里，T.Kernel定义网格，T.Parallel描述 tile 内的并行计算。你注意，我没有写任何线程索引、没有写 shared memory、没有写同步。这就是 tile-level 抽象的直接体现。跑通它，你就理解了这门语言的基本心智模型。
提示：第一次跑建议把 block 设小一点，比如 64x64，方便你观察生成的代码。TileLang 一般提供查看 lowering 结果的接口，一定要用，这是理解它怎么工作的最快路径。
3.2 从加法到 GEMM：tile 分块与内存层级
加法跑通之后，上 GEMM。GEMM 是算子编程的“Hello World”，因为它同时涉及数据复用、内存层级和计算密集。TileLang 写 GEMM 的核心是把大矩阵切成 tile，让每个 tile 的计算能塞进片上存储。
分块参数怎么定？这是有计算依据的。假设你做 fp16 的 GEMM，目标硬件的 shared memory 每 SM 有 164KB 可用，你想让 A 的 tile 和 B 的 tile 都放 shared。如果 block_M = 128、block_N = 128、block_K = 32，那么 A tile 是 128x32x2B = 8KB，B tile 是 32x128x2B = 8KB，加起来 16KB，远小于 164KB，可以开多级流水。如果你把 block_K 拉到 64，A+B 就是 32KB，还能接受；拉到 128，就是 64KB，流水级数就得降。这个权衡没有标准答案，取决于你的 K 维多大、带宽多紧。
我实测下来，TileLang 写 GEMM 的代码量大概是 CUDA 的 1/5 到 1/10，而且换 block 参数只需要改几个数字，不用动循环结构。这一点对做 autotuning 特别友好——你可以写个脚本扫参数空间，让编译器去试，而不是手工改代码再编译。
3.3 调度原语的实际用法与效果
TileLang 的调度原语里，我用得最多的是pipeline、swizzle和layout。pipeline 控制循环的流水线重叠，比如把 K 维循环做成多级流水，让 load 和 compute 重叠；swizzle 控制 shared memory 的排布，避免 bank conflict；layout 控制 tile 到线程的映射，影响寄存器利用率和访存合并。
举个具体的：做 GEMM 时，如果 shared memory 的列 stride 是 32 的倍数，fp16 下很容易 bank conflict。用 swizzle 把列索引异或一下，conflict 就消了。CUDA 里你得手写这个异或逻辑，TileLang 里一行T.swizzle搞定。但要注意，swizzle 不是免费的，它改变了访存模式，可能影响后续的向量化。我踩过的坑是：加了 swizzle 之后 bank conflict 没了，但向量化 load 被破坏，整体性能反而降了。所以每次加调度原语，都要重新 benchmark，别想当然。
4. 实操中遇到的典型问题与排查记录
4.1 编译报错看不懂怎么办
TileLang 这类 DSL 的报错，最烦的是错误信息指向 lowering 之后的 IR，而不是你写的前端代码。我遇到过好几次“某条指令不支持”的报错，翻半天不知道对应我哪一行。后来总结出一个排查套路：先把算子简化到最小可复现，然后逐步加回功能，定位是哪一步触发的。另外，打开编译器的 debug 输出，看 lowering 每一轮之后的 IR，通常能看出问题出在哪一层。
还有一个常见原因是形状不匹配。tile-level 抽象虽然帮你管了很多事，但 tile 形状和 buffer 形状对不上时，报错往往很隐晦。我的习惯是：所有 tile 形状都用变量表示，别写死数字，这样改起来不容易漏。
4.2 性能不达预期怎么定位
性能问题分三类：访存瓶颈、计算瓶颈、调度瓶颈。定位方法不一样。访存瓶颈看带宽利用率，如果离峰值差很远，多半是 layout 或 swizzle 没调好；计算瓶颈看算力利用率，如果低，可能是 tile 太小导致指令级并行不够；调度瓶颈看流水线效率，如果 load 和 compute 没重叠上，就是 pipeline 级数或依赖关系的问题。
我一般先用 profiler 拿到整体利用率，再针对性调。TileLang 的好处是调参快，你可以十分钟内试十几组配置。但要注意，别只盯一个指标。我有次把某个 kernel 的算力利用率从 60% 调到 85%，结果端到端推理速度没变，因为瓶颈在别的地方。算子优化要看全局，别做局部最优的奴隶。
4.3 常见问题速查表
| 现象 | 可能原因 | 排查方向 |
|---|---|---|
| 编译报错指向 IR | 某后端 intrinsic 未实现 | 简化复现，看 lowering 日志 |
| 性能远低于预期 | tile 形状或调度不当 | profiler 看瓶颈类型 |
| 加了 swizzle 反而变慢 | 向量化被破坏 | 对比有无 swizzle 的访存模式 |
| 换卡后结果不对 | 后端 lowering 有 bug | 用参考实现逐元素比对 |
| 流水线没重叠 | 依赖关系或级数不对 | 检查 pipeline 原语参数 |
注意：这张表是我自己踩坑总结的，不一定覆盖所有情况。遇到新问题，第一原则永远是最小复现 + 逐层排查，别一上来就怀疑编译器。
5. TileLang 对国产开源生态的影响与启示
5.1 补上算子供给这块短板
国产开源生态这些年硬件进步快，但软件栈尤其是高性能算子库，一直是短板。模型结构月月变，新算子层出不穷，靠人工一个个手写 CUDA 再适配多张卡，根本跟不上。TileLang 这类 tile-level 语言的价值，是把算子开发从“手工作坊”变成“半自动化流水线”。计算逻辑写一遍，调度按卡定制，新增硬件后端不用动前端。这个模式一旦跑通，算子供给的速度能上一个台阶。
我观察到的另一个影响是人才门槛。以前写高性能算子，得懂硬件微架构、懂汇编、懂 profiling，培养周期以年计。tile-level 抽象把硬件细节收进编译器后，一个懂计算、懂 Python 的工程师，几周就能上手写像样的算子。这对生态的意义是供给端扩容——能写算子的人多了，生态才活得起来。
5.2 对开发者的启示：抽象层级的选择是战略问题
TileLang 给我的最大启示，不是某个具体技术点，而是抽象层级的选择本身就是战略。CUDA 选了 instruction-level，绑死了硬件细节，换来极致性能但牺牲了可移植性；Triton 选了 block-level，平衡了一些；TileLang 选 tile-level，把可移植性拉满，代价是极端场景下表达力受限。没有对错，只有取舍。
对做国产生态的团队来说，这个取舍尤其关键。你的目标是覆盖尽可能多的硬件、尽可能多的算子，还是在少数几张卡上做到极致？前者选高抽象，后者选低抽象。TileLang 明显是前者。理解这一点，你就能判断它适不适合你的场景，而不是盲目跟风。
5.3 后续可以怎么扩展
如果你已经跑通了基本算子，下一步可以往这几个方向走。一是接 autotuning，把 tile 参数、调度参数做成搜索空间，让编译器自动找最优配置；二是接量化算子，int8、int4 的 GEMM 和 attention 在推理里需求很大，tile-level 抽象对量化同样适用；三是接国产卡的专用指令，如果某张卡有特殊的矩阵指令，可以在后端加 intrinsic，前端不用改。这三点做完，一个算子库的骨架就立起来了。
我个人在实际操作中的体会是：别把 TileLang 当 CUDA 的替代品，把它当算子开发的“高级语言”。就像你不会用汇编写业务代码，但汇编依然重要——CUDA 会一直在底层，TileLang 这类语言负责把上层的生产力释放出来。两者不是取代关系，是分工关系。想清楚这一点，你在技术选型时就不会纠结“要不要 all in”，而是知道什么场景用什么工具。
