# TileLang 完整解析：當 80 行 Python 打敗 500 行 CUDA，GPU 編程的新革命
80 行 Python，實現了跟 FlashMLA 手寫 CUDA 版本同等效能的 GPU kernel。比 PyTorch 快 3.76 倍，比 Triton 快將近 2 倍。而且同一份程式碼，能同時跑在 NVIDIA H100、AMD MI300X、華為 Ascend NPU 上面。
這不是某個 AI 生成的幻覺數據。這是 [TileLang](https://github.com/tile-ai/tilelang) 在 ICLR 2026 以 Oral paper 身份發表的實測結果 —— 在近兩萬篇投稿中，只有 1.18% 的論文拿到 Oral。
你可能已經隱約聽過這個名字。2025 年 9 月 DeepSeek 發布 V3.2-Exp 模型時，技術報告裡藏了一段耐人尋味的話：「我們使用高階語言 TileLang 進行快速原型設計......建議社群使用 TileLang 版本進行研究實驗。」同一天，華為、寒武紀、海光同步宣布支持。
這到底是一個純粹的技術突破，還是 GPU 編程生態的板塊位移？我花了相當長的時間研究這個項目的每一個面向。這篇文章，就是我的完整拆解。
---
## TileLang 到底是什麼？一句話說清楚
先釐清一個容易混淆的點：TileLang 不是一個新的程式語言，也不是 CUDA 的替代品。它是一套 **Python 嵌入式的領域特定語言（DSL）**，專門用來寫 GPU/CPU/加速器上的高效能運算 kernel。
什麼意思呢？你用 Python 的語法寫程式碼，但 TileLang 的編譯器會把你寫的東西轉換成跟手寫 CUDA 同等效能的底層指令。
它的核心抽象是 **Tile**（方塊/磁磚）。在 GPU 的世界裡，效能瓶頸往往不在運算本身，而在資料搬運。資料要從全域記憶體搬到共享記憶體，再搬到暫存器，每一層的頻寬和延遲都天差地別。TileLang 讓你用「方塊」的視角來思考這些資料搬運 —— 你告訴編譯器「把這塊資料放到共享記憶體」、「在暫存器裡做矩陣乘法」，至於底下怎麼分配線程、怎麼做記憶體排列，編譯器幫你搞定。
這跟 CUDA、Triton 有什麼不一樣？看一個直覺的對比：
| 框架 | 你需要管什麼 | 編譯器幫你做什麼 |
|------|------------|----------------|
| **CUDA** | 線程、warp、共享記憶體、同步、指令排程 | 幾乎不幫你 |
| **Triton** | Block 級運算邏輯 | 記憶體管理、部分排程 |
| **TileLang** | Tile 級資料流和記憶體放置 | Layout 推導、流水線、warp 特化 |
簡單說，CUDA 是手動擋，Triton 是半自動，TileLang 是帶主動安全輔助的自動擋 —— 但你仍然可以隨時切回手動模式。
---
## 從北大實驗室到 ICLR Oral：背後的故事
TileLang 的故事要從北京大學說起。主要開發者 [LeiWang1999（王磊）](https://github.com/LeiWang1999)在楊智教授的指導下，和團隊成員 chengyupku、nox-410 一起打造了這個專案。部分工作是在微軟研究院實習期間完成的，微軟的 Lingxiao Ma、Yuqing Xia 等研究員也提供了重要指導。
時間線大致是這樣的：
- **2024 年**：核心開發期，同時開發了 BitBLAS（混合精度計算庫）作為早期驗證
- **2025 年 1 月 20 日**：正式開源，首個版本發布
- **2025 年 2 月**：v0.1.0 發布，加入 debug 工具和 WebGPU 支持
- **2025 年 3 月**：80 行 Python 實現 FlashMLA，效能打平手寫 CUDA
- **2025 年 4 月**：論文提交 arXiv（[2504.17577](https://arxiv.org/abs/2504.17577)）
- **2025 年 9 月**：DeepSeek V3.2-Exp 正式採用 TileLang operators
- **2025 年 12 月**：加入 CuTe DSL 後端；同月 NVIDIA 推出 CUDA Tile
- **2026 年 1 月**：被 ICLR 2026 接受為 Oral paper
- **2026 年 2 月**：TileLang Puzzles 互動學習工具上線
一年出頭的時間，從零到 5,200+ GitHub stars，被頂級會議 Oral 接受，被 DeepSeek 這種量級的項目採用。這個速度在開源基礎設施項目裡相當罕見。
---
## 編譯器深度拆解：五階段管線如何運作
TileLang 真正的技術含量藏在它的編譯器裡。你寫的那 80 行 Python，要經過五個階段才能變成 GPU 上飛速運行的機器碼。
### 管線全貌
```
Python Code → Parser → IR Builder → Optimization → Codegen → GPU Execution
↓ ↓ ↓ ↓
Python AST TVM IR Layout/Pipeline CUDA C/HIP C/
→ TileLang Inference LLVM IR
AST
```
**第一步 Parser**：你用 `@T.prim_func` 裝飾器寫的 kernel 函式，先被解析成 Python AST，再轉成 TileLang 自己的 AST。
**第二步 IR Builder**：AST 被轉換成 TVM 的中間表示（IR）。TileLang 建構在 [Apache TVM](https://tvm.apache.org/) 之上，這讓它能直接利用 TVM 成熟的語法樹基礎設施。
**第三步 Optimization**：這是魔法發生的地方。三大核心優化在這裡執行。
**第四步 Codegen**：優化後的 IR 被翻譯成目標平台的程式碼 —— CUDA C/C++、HIP C/C++、或 LLVM IR。
**第五步 Execution**：JIT 編譯並在目標硬體上執行。
### 三大核心優化機制
要理解 TileLang 為什麼能用這麼少的程式碼達到這麼高的效能，關鍵在這三個編譯器 pass：
**1. Layout Inference（佈局推導）**
你在 TileLang 裡分配一塊共享記憶體或暫存器，不需要指定它該怎麼在線程之間分配。編譯器會根據後續的運算（比如 `T.gemm`）自動推導出最佳的記憶體佈局和線程綁定。
這用一個分層優先級系統實現 —— 越高優先級的運算（比如 Tensor Core 的 GEMM）對佈局有越嚴格的要求，編譯器從上往下逐層推導，直到所有 buffer 的佈局都確定。
**2. Pipeline Inference（流水線推導）**
在高效能 kernel 裡，資料搬運和計算要重疊執行。傳統做法需要開發者手動設計 software pipeline —— 在 CUDA 裡這意味著大量的 barrier、async copy、多重 buffer。TileLang 只需要你指定 `T.Pipelined(iters, num_stages=N)`，編譯器會自動分析依賴關係，決定哪些操作可以重疊。
**3. Warp Specialization（自動 warp 特化）**
在 NVIDIA Hopper（H100）架構上，最佳效能需要 warp specialization：讓一部分 warp 專門負責資料搬運（用 TMA），另一部分專門做計算。手動實現這個需要管理 `mbarrier` 同步物件、生產者/消費者邏輯 —— 這在 CUDA 裡是極其痛苦的工作。
TileLang 完全自動化了這個過程。它分析 buffer 的使用模式，自動將語句分類為生產者和消費者，然後插入正確的同步屏障。
### 實際程式碼長什麼樣？
一個用 TileLang 寫的矩陣乘法 kernel，核心邏輯大概是這樣：
```python
import tilelang
import tilelang.language as T
@tilelang.jit
def matmul(A, B, block_M=128, block_N=128, block_K=32):
M, K = A.shape
K, N = B.shape
C = T.empty((M, N), A.dtype)
with T.Kernel(T.ceildiv(N, block_N), T.ceildiv(M, block_M), threads=128) as (bx, by):
A_shared = T.alloc_shared((block_M, block_K), A.dtype)
B_shared = T.alloc_shared((block_K, block_N), B.dtype)
C_local = T.alloc_fragment((block_M, block_N), "float32")
T.clear(C_local)
for ko in T.Pipelined(T.ceildiv(K, block_K), num_stages=3):
T.copy(A[by * block_M, ko * block_K], A_shared)
T.copy(B[ko * block_K, bx * block_N], B_shared)
T.gemm(A_shared, B_shared, C_local)
T.copy(C_local, C[by * block_M, bx * block_N])
return C
```
注意看 —— 沒有 `threadIdx`，沒有手動同步，沒有記憶體佈局計算，沒有 warp 分工。你只需要說「分配共享記憶體」、「複製資料」、「做矩陣乘法」、「用 3 階段流水線」。剩下的，編譯器全包了。
---
## 效能數據不會說謊：TileLang 的硬實力
空談設計哲學沒有意義，看數字。以下數據來自 [TileLang 論文](https://arxiv.org/pdf/2504.17577)和 [AMD ROCm 技術部落格](https://rocm.blogs.amd.com/ecosystems-and-partners/rocm-tilelang-kernel/README.html)的實測結果。
### FlashMLA Benchmark（DeepSeek 的 MLA 注意力機制）
| 框架 | 相對延遲（越低越好） | 代碼量 |
|------|-------------------|--------|
| PyTorch | 1.00x（基準） | N/A |
| Triton | 0.655x | ~200+ 行 |
| TileLang | 0.371x | ~80 行 |
| FlashMLA (CUDA) | ~0.37x | 數千行 CUDA + CUTLASS |
TileLang 用 80 行 Python 達到了跟手寫 CUTLASS 模板相當的效能，這不是小幅改進，這是量級上的開發效率提升。
### 跨平台效能摘要
根據 [ICLR 2026 論文的 Abstract](https://openreview.net/forum?id=Jb1WkNSfUB)：
- **NVIDIA H100**：比 Triton 快最高 **5 倍**
- **AMD MI300X**：比 Triton 快最高 **6 倍**
在 AMD MI300X 上的 Flash Attention 測試中，TileLang 實現比原生 PyTorch 快了 2.7 倍，比 Triton 快 1.53 倍。而且自動調優框架搜尋 108 個配置只需要約 1 秒。
### 什麼場景效能最突出？
TileLang 的效能優勢在 **融合型 kernel**（fused kernels）上最為明顯 —— 也就是把多個操作合併到一個 kernel 裡的場景。FlashAttention、MLA Decoding、Linear Attention 這類需要精細記憶體管理的複雜 kernel，正是 TileLang 的主戰場。
對於簡單的 elementwise 操作或標準 GEMM，TileLang 和 cuBLAS/Triton 的差距就沒這麼大。Tile 級抽象的優勢在複雜度越高的 kernel 中越能體現。
---
## 四方混戰：TileLang vs Triton vs CUDA vs cuTile
GPU kernel 編程領域正在經歷一場前所未有的混戰。讓我把幾個主要玩家攤開來比較。
### 完整對比表
| 維度 | TileLang | Triton | CUDA | NVIDIA cuTile |
|------|----------|--------|------|---------------|
| 開發方 | 北大 / tile-ai | OpenAI | NVIDIA | NVIDIA |
| 語言 | Python DSL | Python DSL | C/C++ | Python DSL |
| 抽象層級 | Tile-level | Block-level | Thread-level | Tile-level |
| 記憶體控制 | 明確放置 | 隱式管理 | 完全手動 | 隱式 |
| 跨平台 | NVIDIA + AMD + 華為 + 摩爾 | NVIDIA + AMD | 僅 NVIDIA | 僅 NVIDIA |
| 自動 Warp Specialization | 有 | 開發中 | 手動 | 有 |
| 自動 Software Pipeline | 有 | 有 | 手動 | 有 |
| FP8 支持 | 部分 | 完整 | 完整 | 完整 |
| 生態成熟度 | 早期 | 中等 | 非常成熟 | 初始 |
| 開源 | Apache 2.0 | MIT | 部分 | 開源 |
### 關鍵差異解讀
**TileLang vs Triton**：最核心的差異在於記憶體控制的粒度。Triton 隱藏了共享記憶體的細節，讓編譯器自動決定。TileLang 則明確暴露記憶體層級，讓開發者可以精確控制資料放在哪裡。這在簡單 kernel 上差異不大，但在複雜融合 kernel（如 FlashAttention）上，TileLang 的明確控制帶來顯著的效能優勢。
**TileLang vs cuTile**：NVIDIA 在 2025 年 12 月隨 CUDA 13.1 推出了 cuTile，這被廣泛認為是對 TileLang 和 tile 編程趨勢的直接回應。cuTile 的定位跟 TileLang 很像 —— Python DSL、tile 級抽象。但關鍵差異是：cuTile 只支持 NVIDIA GPU，而且目前僅限 Blackwell 架構。TileLang 則是跨平台的。
**為什麼 NVIDIA 也跟著做？** Hacker News 上一位用戶一針見血地指出：「NVIDIA 不希望 CUDA 的開發（如 FlashAttention）遷移到 Triton，因為 Triton 也支持 AMD。如果生態從純 CUDA 遷移到 Triton，那對 NVIDIA 的鎖定效應不利。」cuTile 的出現，某種程度上是 NVIDIA 在 tile 編程趨勢下的防禦性佈局。
---
## 不只是技術：GPU 編程的地緣棋局
如果你只把 TileLang 當成一個技術項目來看，你會錯過故事最重要的一章。
2025 年 9 月 29 日，DeepSeek 發布 V3.2-Exp。同一天，華為 Ascend 團隊更新了 CANN 對 V3.2-Exp 的支持，寒武紀更新了 vLLM-MLU，SGLang 確認了多後端支持。[Tom's Hardware 的報導](https://www.tomshardware.com/tech-industry/deepseek-new-model-supports-huawei-cann)指出，這種同步性暗示著預先協調。
分析機構 [HelloChinaTech 的評論](https://hellochinatech.com/p/deepseek-tilelang-china-ai-stack-nvidia)更直接：「TileLang 的推出標誌著建構中國自主 AI 軟體棧的第二階段。」第一階段是 FP8 精度標準的建立（DeepSeek V3 的訓練方法論），第二階段就是透過 TileLang 建立跨硬體的編程抽象層。
這個邏輯鏈條是這樣的：
1. 中國的 AI 晶片廠商（華為、寒武紀、摩爾線程、海光）各自有不同的硬體架構
2. 如果每家都要開發者學一套新的編程模型，生態碎片化會極為嚴重
3. TileLang 提供了一個統一的抽象層 —— 同一份程式碼可以編譯到不同後端
4. 這降低了從 CUDA 遷移的門檻，也減少了對 NVIDIA 生態的依賴
但也要保持冷靜。正如 [IEEE Spectrum 的分析](https://spectrum.ieee.org/china-ai-chip)所指出的，中國晶片廠商的硬體能力與 NVIDIA 仍有代差。軟體棧的統一是必要條件，但不是充分條件。CUDA 的護城河不僅是語言本身，更是 15 年積累的函式庫、工具鏈、開發者社群和 debug 生態。
---
## 冷靜看待：TileLang 的局限與挑戰
技術文章如果只說好的，不說問題，那就是廣告。TileLang 確實有幾個需要正視的局限。
**FP8 支持仍不完善。** 學術論文 [Tawa（arXiv: 2510.14719）](https://arxiv.org/html/2510.14719v1)的實測數據顯示，在 FP8 精度下，TileLang 的 GEMM 效能落後 Tawa 最高 3.99 倍，FlashAttention 的 FP8 配置甚至無法成功執行。論文指出原因是 TileLang 缺乏處理 FP8 WGMMA 路徑所需的佈局管理和流水線調度。在 FP8 訓練越來越普及的當下，這是一個必須盡快補上的短板。
**生態仍處於早期階段。** 5,200 個 GitHub stars 聽起來不少，但跟 Triton 的生態規模相比仍有明顯差距。文件、教學資源、第三方整合都還在建設中。好消息是 TileLang Puzzles 等互動學習工具正在降低入門門檻。
**設計上的限制。** 社群開發者在實作 DeepSeek DSA（Dynamic Sparse Attention）時發現，TileLang 0.1.6 版本在非 per-head layout 的場景下存在限制，無法很好地支持某些進階開發思路（[相關 Issue #1199](https://github.com/tile-ai/tilelang/issues/1199)）。
**主要針對 AI workload。** TileLang 的抽象是為深度學習運算（GEMM、Attention、Convolution）量身設計的。如果你的需求是通用的 GPU 計算（圖形渲染、物理模擬、密碼學），CUDA 仍然是更好的選擇。
---
## 開發者該怎麼做？實際行動建議
說了這麼多，對於不同角色的讀者，我的建議是不一樣的。
### 如果你是 AI 研究者
TileLang 是你現在就該開始學的工具。原因很簡單：用 80 行 Python 實現一個跟 CUTLASS 同等效能的 MLA kernel，這種開發效率的提升對研究的迭代速度有根本性的影響。DeepSeek 官方推薦研究者使用 TileLang 版本，這不是客套話。
上手路徑：
1. `pip install tilelang` 安裝
2. 跑一遍 [TileLang Puzzles](https://github.com/tile-ai/tilelang)（10 個漸進式題目）
3. 閱讀 [FlashMLA 的 80 行實現](https://tilelang.com/deeplearning_operators/deepseek_mla.html)
4. 在你的研究中嘗試用 TileLang 做快速原型
### 如果你是 GPU 工程師
把 TileLang 加入你的工具箱，但不要丟掉 CUDA。TileLang 擅長的是快速原型和跨平台部署，CUDA 擅長的是極致效能調優和通用計算。一個合理的工作流程是：先用 TileLang 快速驗證想法和算法正確性，再視需要用 CUDA 做最後一哩路的效能榨取。
同時，密切關注 NVIDIA cuTile 的發展。tile 編程模型已經是確定的趨勢，無論你最終選擇哪個框架，理解 tile 級抽象的思維方式都是值得投資的。
### 如果你是技術決策者
關鍵觀察：GPU 編程正在從 thread-level 往 tile-level 遷移，這是一個跟 2012 年深度學習框架從手寫反向傳播到自動微分一樣量級的抽象層級提升。TileLang、Triton、cuTile 代表的是同一個方向的不同實現。
現在就要做的：
- 評估你的團隊對 CUDA 的依賴程度
- 在非關鍵路徑上試點 TileLang 或 Triton
- 如果你需要跨 NVIDIA/AMD 部署，TileLang 的跨平台能力是一個實質優勢
---
## Tile 編程時代已經到來
回到開頭的問題：TileLang 是真正的技術革命，還是曇花一現？
我的判斷是：**TileLang 本身還太早期，無法預測它是否會成為最終的贏家。但 tile 編程模型作為一個範式，已經是確定的趨勢。**
證據就在競爭者的反應裡。NVIDIA 推出 cuTile，OpenAI 的 Triton 也在開發 warp specialization，NVIDIA 還發表了 Tilus —— 每個人都在往同一個方向跑。當一個技術方向能讓所有主要玩家同時響應，你就知道這不是曇花一現。
TileLang 的獨特價值在於三件事：它是開源的、它是跨平台的、它已經被工業級項目（DeepSeek）驗證。在一個 NVIDIA 鎖定效應主導的世界裡，這三個特性加在一起，足以讓它成為值得長期關注的項目。
至於那 80 行 Python？那不是重點。重點是它背後代表的理念：GPU 編程不必這麼痛苦，開發者值得更好的抽象。
---
## 延伸閱讀
- [TileLang 官方文檔](https://tilelang.com/) - 最完整的入門資源
- [TileLang GitHub](https://github.com/tile-ai/tilelang) - 原始碼、範例、Issue 討論
- [TileLang 論文 (arXiv: 2504.17577)](https://arxiv.org/abs/2504.17577) - 完整技術細節
- [ICLR 2026 Oral 頁面](https://openreview.net/forum?id=Jb1WkNSfUB) - 審稿意見和討論
- [AMD ROCm Blog: TileLang on MI300X](https://rocm.blogs.amd.com/ecosystems-and-partners/rocm-tilelang-kernel/README.html) - AMD 平台實測
- [FlashMLA 80 行實現詳解](https://tilelang.com/deeplearning_operators/deepseek_mla.html) - 最佳入門案例
- [TileScale 分散式擴展](https://github.com/tile-ai/tilescale) - 多 GPU/多節點編程
- [NVIDIA CUDA Tile 介紹](https://developer.nvidia.com/blog/focus-on-your-algorithm-nvidia-cuda-tile-handles-the-hardware/) - NVIDIA 的回應方案
Sign in via Google
Sign in via Facebook
Sign in via X(Twitter)
Sign in via GitHub
Sign in via Dropbox
Sign in with Wallet
Connect another wallet