1. 从 tilelang 这个名字说起：它到底想解决什么问题
第一次看到 tilelang 这个词，很多人会下意识拆成 “tile + lang”，直觉上跟“瓦片”和“语言”有关。这个直觉基本是对的，但需要补上关键的一层：这里的 tile 不是建筑上的瓦片，而是高性能计算里那个被反复念叨的概念——数据分块（tiling）。而 lang 则指向一门面向分块计算的领域特定语言（DSL）。把这两个词拼在一起，tilelang 的定位就清楚了：它试图把“怎么切分数据、怎么调度计算、怎么映射到硬件”这套原本靠手写底层代码和大量调参才能搞定的事情，抽象成一套更接近数学表达、又能被编译器吃透的语言。
我在实际接触这类工具之前，长期用的是最原始的路子：手写 kernel，手动调 block size，手动管 shared memory，一个算子调一整天是常事。那种体验就像用汇编去写业务逻辑，能跑，但每换一个 shape、换一块硬件，前面的功夫基本白费。tilelang 这类东西出现的背景，正是为了解决这个痛点——让写高性能算子的人从“硬件细节的泥潭”里抬起头来，把精力放回算法结构本身。
它适合谁？我认为有三类人值得认真看：第一类是做推理引擎、训练框架底层优化的工程师，天天跟矩阵乘、卷积、attention 打交道；第二类是算法研究员，想让自己的新结构快速跑出接近硬件极限的性能，但不想陷进 CUDA 细节；第三类是想理解“现代编译器怎么把高层描述变成高效机器码”的学习者。哪怕你暂时不写算子，读懂 tilelang 的设计思路，对理解整个 AI 编译栈的演进方向也很有帮助。
需要先说明一点：tilelang 目前仍是一个相对年轻的方向，不同实现、不同版本的 API 和成熟度差异不小。下面我讲的内容，一部分来自公开资料的梳理，一部分是我基于同类 DSL（比如 TVM 的 TensorIR、Triton 等）的实操经验做的合理推演。凡是推演的部分，我都会明确标出来，避免误导。
2. 核心设计思路拆解：为什么是“分块”而不是别的
2.1 分块计算为什么成了绕不开的核心
要理解 tilelang 为什么把 tile 当成一等公民，得先回到硬件现实。现代加速器（无论是 GPU 还是各类 AI 芯片）的性能瓶颈，早就不在算力本身，而在数据搬运。算一个矩阵乘，理论 FLOPs 很高，但如果你每次都从显存里把数据捞一遍，带宽根本喂不饱计算单元。解决办法就是分块：把大矩阵切成能塞进片上高速缓存（shared memory / register file）的小块，让数据在被反复复用的过程中尽量待在离计算单元最近的地方。
这个道理说起来简单，做起来极其琐碎。块切多大？切完怎么在多个计算单元之间分配？片上缓存怎么复用？边界怎么处理？这些决策互相耦合，改一个参数可能整体性能掉一半。传统做法是靠专家经验加暴力搜索，而 tilelang 的思路是：把这些决策显式地表达在语言层面，让程序员能描述“我想怎么分块”，同时让编译器负责把描述翻译成高效的底层调度。
这背后其实是一个权衡。完全自动的调度器（比如早期的一些 auto-scheduler）理论上最省心，但实际中往往搜不出最优解，或者搜索时间长得离谱。完全手写的调度最可控，但门槛高、不可移植。tilelang 走的是中间路线——给程序员一套结构化的、带语义的分块原语，既保留人的判断力，又让编译器承担机械性的翻译工作。这个取舍我认为是务实的，因为高性能算子优化本质上是个“半结构化”问题，纯自动和纯手工都不理想。
2.2 语言抽象层级的选择：为什么不做成通用语言
有人会问，为什么不干脆用 Python 或 C++ 直接写，非要搞一门 DSL？这里的关键在于可分析性。通用语言太灵活，编译器很难从一堆指针操作和动态控制流里推断出“这是一次规整的分块矩阵乘”。而 DSL 通过限制表达形式，换来了编译器对程序结构的精确理解。比如在 tilelang 里，你描述一个 tile 的形状、布局、计算关系，编译器就能明确知道数据依赖、复用模式，从而做出合理的流水线安排和内存分配。
这跟 Triton 的设计哲学有相似之处，但 tilelang 更强调“tile 作为一等抽象”这件事。在 Triton 里，你操作的是 block 级别的张量；在 tilelang 里，tile 本身可以携带更多语义信息，比如它对应哪一级存储、以什么布局排布、在哪些维度上被复用。这种更细粒度的表达，理论上给优化留出了更大空间，代价是语言本身更复杂，学习曲线更陡。
我的判断是：如果你的场景是标准算子（GEMM、conv、attention），用成熟度更高的方案可能更省事；如果你要做的是新结构、非标准算子，或者需要精细控制数据流，tilelang 这类以 tile 为核心的 DSL 才真正体现价值。选型时别被“新”字冲昏头，先看自己的需求落在哪一档。
2.3 编译流程的整体骨架
从高层描述到可执行代码，tilelang 这类系统通常要经过几个阶段，我按自己的理解梳理一下：
- 前端解析：把 DSL 写的程序转成中间表示（IR），这一步会做类型检查、形状推断。
- 分块与布局推导：根据程序员写的 tile 描述，推导出每个 tile 的存储层级、布局方式、依赖关系。
- 调度与优化：做循环变换（tiling、fusion、pipelining）、内存分配、指令选择。这一步是性能的关键，也是各家实现拉开差距的地方。
- 代码生成：把优化后的 IR 降级成目标硬件的代码（比如 CUDA C、PTX 或某种后端 IR）。
- 运行时集成：处理 host 侧调用、内存管理、kernel 启动配置。
这个流程里，第 3 步最考验功力。循环怎么融合、流水线怎么排、寄存器怎么分配，每一个决策都直接影响最终性能。tilelang 的价值主张，就是让程序员在第 2 步用声明式的方式把意图表达清楚，把第 3 步的重活交给编译器。但现实是，编译器不可能对所有情况都做出最优决策，所以好的 DSL 一定会留出“手动干预”的口子，比如允许指定某些调度原语。这一点在实操中非常重要，后面我会展开。
3. 核心细节解析与实操要点
3.1 tile 的声明：形状、布局与存储层级
在 tilelang 的语境里，声明一个 tile 通常要交代三件事：形状（shape）、布局（layout）、存储层级（memory scope）。形状好理解，就是各维度大小。布局指的是数据在物理存储里怎么排——是行优先还是列优先，有没有做 swizzle 来避免 bank conflict。存储层级则决定这个 tile 放在寄存器、shared memory 还是全局内存。
这三者不是独立的。举个例子，你把一个 tile 声明成放在 shared memory，形状是 32x32，但如果布局没选好，访问时就会撞上 bank conflict，性能直接腰斩。我踩过的坑就是：早期只顾着调 tile 大小，忽略了布局，结果算出来的 kernel 比预期慢了三倍，排查半天才发现是 shared memory 的访问模式有问题。
实操建议是：先用编译器默认的布局推导跑通，再针对热点做手动调整。别一上来就手写所有布局参数，那样既费时又容易出错。tilelang 这类工具通常会提供布局推导的 pass，默认策略在多数情况下够用，只有当你明确知道瓶颈在哪时，再介入。
3.2 数据复用与流水线：性能的真正来源
分块的目的就是复用。一个 tile 被加载到片上之后，应该尽可能多地被计算单元读取，而不是算一次就扔。tilelang 里表达复用，通常靠的是在 tile 维度上做循环嵌套，让内层循环反复访问同一个 tile。编译器会分析这些访问模式，决定要不要做 double buffering、要不要把加载和计算重叠起来。
流水线（pipelining）是另一个关键。理想情况下，当计算单元在处理当前 tile 时，下一块数据应该已经在后台加载了。这个“加载-计算”的重叠，是隐藏内存延迟的核心手段。tilelang 一般会提供类似pipeline或async_copy的原语，让你显式地表达这种重叠意图。
注意：流水线的级数不是越多越好。级数太深会占用更多片上存储，反而挤占了本该给计算用的空间。我一般从 2 级开始试，观察性能曲线，找到拐点再定。
这里有个经验：先确认你的 kernel 是计算密集还是访存密集。如果是访存密集，流水线和复用是重点；如果是计算密集，那得看指令级并行和寄存器压力。搞错方向，优化就是白费力气。判断方法很简单，算一下 arithmetic intensity（计算量除以访存量），低于硬件平衡点的就是访存密集。
3.3 边界处理：最容易被忽视的细节
真实问题里的矩阵尺寸很少是 tile 大小的整数倍，所以边界处理几乎躲不掉。tilelang 这类 DSL 通常会提供带谓词（predicate）的访问方式，或者让你显式地写边界判断。这里有个取舍：加谓词会带来额外开销，但不加就可能越界。
我的做法是分情况：如果边界 tile 占比很小，直接用谓词保护，简单可靠；如果边界占比大（比如小矩阵），那可能得考虑专门为边界写一条路径，或者干脆 pad 到整数倍。pad 的代价是多了些无效计算和存储，但换来了规整的访问模式，有时候反而更快。这个决策没有标准答案，得实测。
还有一个坑：不同硬件对越界访问的容忍度不一样。有些平台越界读会返回垃圾值但不报错，有些直接崩。所以别指望“反正结果用不到”就放任越界，调试阶段一定要开边界检查。
4. 实操过程与核心环节实现
4.1 环境准备与依赖确认
动手之前，先把环境理清楚。tilelang 这类工具通常依赖一套编译栈，可能包括 LLVM、特定版本的 CUDA 工具链、Python 绑定等。我的习惯是先建一个干净的虚拟环境，把版本钉死，避免“昨天还能跑今天就不行”的经典问题。
python -m venv tilelang-env source tilelang-env/bin/activate pip install --upgrade pip # 具体安装命令以官方文档为准，这里只示意流程 pip install tilelang装完之后，第一件事是跑官方给的示例，确认工具链通了。别急着写自己的算子，先用最小例子验证编译、运行、结果正确性这条链路。我见过太多人跳过这步，结果后面出问题时分不清是环境问题还是代码问题。
提示：记录下你用的每个组件的版本号。高性能编译栈对版本极其敏感，出问题时版本信息是排查的第一手资料。
4.2 一个矩阵乘的完整实现思路
拿最经典的 GEMM 举例，用 tilelang 写大概会经历这么几步。首先是定义输入输出的 tile 形状，比如把 A、B、C 都按 M、N、K 三个维度切块。然后声明每个 tile 的存储层级：A 和 B 的块加载到 shared memory，累加器放在寄存器。接着写循环结构：外层遍历 K 方向的块，内层做 tile 级的乘加。
关键参数是 block 的大小。这个不能拍脑袋，得算。假设目标硬件的 shared memory 每块 SM 有 48KB 可用，A 的块是 BM x BK，B 的块是 BK x BN，每个元素 4 字节（fp32），那么占用是(BM*BK + BK*BN)*4字节。要让它小于可用容量，还得留出余量给流水线的双缓冲。比如 BM=BN=64、BK=16，占用是(64*16 + 16*64)*4 = 8KB，双缓冲就是 16KB，留有余地。
寄存器那边也要算。累加器是 BM x BN 个元素，如果 BM=BN=64，那就是 4096 个 fp32，分到每个线程（假设 256 线程）是 16 个寄存器，加上其他临时变量，寄存器压力可控。如果 BM、BN 再大，寄存器就可能溢出，导致 spill，性能暴跌。
# 伪代码示意，具体 API 以实际实现为准 @tilelang.kernel def gemm(A, B, C, M, N, K): BM, BN, BK = 64, 64, 16 # 声明 tile 及其存储层级 A_shared = alloc_shared([BM, BK]) B_shared = alloc_shared([BK, BN]) C_local = alloc_register([BM, BN]) # 主循环 for k in range(0, K, BK): copy(A[k:k+BK, :], A_shared) copy(B[:, k:k+BK], B_shared) for i, j in tile(BM, BN): for kk in range(BK): C_local[i, j] += A_shared[i, kk] * B_shared[kk, j] copy(C_local, C)这段是示意，真实 API 会有差异，但结构是通的。写完之后，先验证正确性——拿小尺寸跟参考实现（比如 numpy）对比，误差在容忍范围内才算过。别一上来就测性能，正确性没过，性能数字毫无意义。
4.3 调参与性能验证
正确性过了之后，进入调参。我的流程是：固定其他参数，一次只动一个，记录性能。先调 block 大小，再调流水线级数，最后调布局。每次改动都跑多轮取中位数，避免被单次波动误导。
性能验证不能只看总时间，要拆开看。理想情况下，工具会提供 profiling 接口，能看到计算单元利用率、内存带宽占用、occupancy 等指标。如果计算利用率低，说明数据没喂饱；如果带宽打满但算力没用上，说明复用不够。根据指标反推该调哪个参数，比盲目试要高效得多。
注意：不同 shape 的最优参数往往不同。别指望一套参数打天下，实际部署时通常要针对常见 shape 做几套配置，运行时按 shape 选。这叫“shape specialization”，是工业界的常规操作。
4.4 从单算子到整网集成
单算子跑得快，不代表整网快。集成时要注意几点：算子之间的数据布局要能衔接，避免频繁的 layout 转换；内存分配要复用，别每个算子都申请释放一遍；kernel 启动开销在小算子上可能占比很高，必要时做算子融合。
tilelang 这类 DSL 通常会和上层框架（比如 PyTorch）通过某种 bridge 对接。对接时最容易出问题的是数据布局和 dtype 的一致性。我遇到过框架传过来的是 NCHW，kernel 期望 NHWC，结果算出来数值对但性能差一大截，因为中间偷偷做了转置。集成阶段一定要把每个算子的输入输出契约写清楚。
5. 常见问题与排查技巧实录
5.1 编译期问题速查
| 现象 | 可能原因 | 排查方向 |
|---|---|---|
| 编译报形状不匹配 | tile 声明与输入维度不一致 | 打印各张量 shape，逐维核对 |
| 编译超时或卡死 | 调度搜索空间过大 | 缩小 tile 候选范围，关掉激进优化 |
| 生成的代码跑不起来 | 后端工具链版本不匹配 | 核对 CUDA/LLVM 版本，看生成日志 |
| 布局推导失败 | 访问模式过于复杂 | 简化表达式，或手动指定布局 |
编译期问题相对好查，因为错误信息通常指向具体位置。麻烦的是那些“编译过了但结果不对”的情况。
5.2 运行期问题与定位思路
结果不对，第一步永远是缩小规模。把矩阵降到 4x4、8x8，用肉眼或 numpy 对比，看错在哪个位置。如果小规模对、大规模错，多半是边界处理或累加精度问题。如果小规模就错，那是逻辑问题，逐行核对。
性能不达预期，先确认基线。你手写的 CUDA 版本跑多少？厂商库（比如 cuBLAS）跑多少？有了基线才知道差距在哪。如果连基线都没测过就谈优化，那是空中楼阁。
还有一个隐蔽的坑：测量方式本身有问题。没做 warmup、没同步、没排除首次编译开销，测出来的数字全是噪声。我一般会跑 100 次取稳定后的中位数，并且用工具自带的计时接口而不是 Python 的 time。
5.3 独家避坑经验
第一条，别迷信默认配置。默认参数是为了通用性，不是为了你的特定场景。拿到工具第一件事就是把关键参数摸一遍，知道每个参数影响什么。
第二条，保留可回退的版本。优化过程中经常出现“改了半天反而更慢”的情况，如果没有版本管理，回退都回不去。我习惯每做一组有效改动就 commit 一次，附上性能数字。
第三条，关注数值精度。高性能优化里经常用 fp16、bf16 甚至更低精度，累加时如果精度不够，结果会漂。特别是 K 维度很长的时候，累加误差会累积。必要时用 fp32 累加，或者做分块累加再合并。
第四条，别忽略 host 侧开销。kernel 再快，如果 host 侧调度、内存拷贝占了大头，整体还是慢。端到端测一遍，别只盯着 kernel 时间。
6. 我对 tilelang 这类工具的判断与使用建议
用了这么久同类工具，我的整体感受是：tilelang 代表的方向是对的，但别把它当成银弹。它解决的是“高性能算子开发效率”这个真问题，但它不能替代你对硬件和算法的理解。编译器再聪明，也不知道你的业务里哪个 shape 最重要、哪个精度可以牺牲。这些判断还得人来做。
具体到使用建议，我分三种情况说。如果你只是想让标准算子跑得快，先用厂商库，不够再考虑 DSL；如果你在做新算子、新结构，tilelang 这类工具能帮你快速迭代，值得投入学习；如果你在做编译器或系统研究，那 tilelang 的设计本身就是很好的研究对象，它的分块抽象、调度策略、代码生成路径，都值得细读。
学习路径上，我的建议是先懂硬件模型（存储层级、并行执行、内存带宽），再学 DSL 语法。反过来学，你只会照猫画虎，遇到问题不知道怎么调。硬件模型是根，DSL 是叶，根扎稳了，叶子怎么长都不慌。
最后分享一个我自己的习惯：每次用新工具写算子，我都会同时手写一版最朴素的实现作为参照。不是为了比谁快，而是为了在结果不对时有个可信的对照。这个笨办法帮我省了无数排查时间。工具会骗你，但你自己写的、能跑对的朴素版本不会。
