GPU 编程 DSL 全景对比：TileLang、Triton、Mojo、cuTile、CuTe、ThunderKittens 与 CuTe DSL

1. 引言：为什么 GPU 编程语言突然"百花齐放"

过去十年，写 GPU 高性能算子基本只有两条路：要么用 CUDA C++ 手写，自己管线程、共享内存、寄存器、流水线；要么用 cuBLAS/cuDNN 这类库，牺牲灵活性换性能。前者门槛极高，后者覆盖不到新算子。

随着大模型把注意力机制、MoE、量化、长上下文等新算子推到台前，"手写 CUDA 太慢、调库不够用"的矛盾被放大。于是出现了一整批面向 GPU 的领域特定语言（DSL）与编程模型，它们各自从不同角度切入：有的想让你用 Python 写接近手写 CUDA 的代码，有的想用 tile 抽象屏蔽硬件细节，有的干脆重造一门语言。

本文对比 TileLang、Triton、Mojo、cuTile、CuTe（C++ 模板）、ThunderKittens、CuTe DSL 这七种方案，先逐个讲诞生动机与核心特性，最后梳理一条时间线与它们之间的交互关系。

2. 逐个拆解：动机与核心特性

2.1 Triton：用 Python 写"块级"GPU 程序

诞生动机：OpenAI 在 2021 年前后推出 Triton，目标是让不熟悉 CUDA 的研究者也能写出接近手写性能的算子。传统 CUDA 要求你以"线程"为单位思考，而 Triton 让你以"块（block/tile）"为单位思考。

核心特性：

• 基于 Python 的 DSL，用 @triton.jit 装饰函数，编译为 PTX。
• 以 tile 为基本操作单位，tl.load / tl.store / tl.dot 等原语自动处理向量化与内存合并。
• 编译器负责共享内存分配、线程映射、流水线调度，用户不直接接触 __shared__。
• 生态成熟：PyTorch 的 torch.compile 后端之一、vLLM/FlashAttention 等大量项目使用。

import triton
import triton.language as tl

@triton.jit
def add_kernel(x_ptr, y_ptr, out_ptr, n, BLOCK: tl.constexpr):
pid = tl.program_id(0)
offs = pid * BLOCK + tl.arange(0, BLOCK)
mask = offs < n
x = tl.load(x_ptr + offs, mask=mask)
y = tl.load(y_ptr + offs, mask=mask)
tl.store(out_ptr + offs, x + y, mask=mask)

定位：目前最主流的"Python 写 GPU kernel"方案，抽象层次适中，性能与易用性平衡最好。

2.2 TileLang：把 tile 抽象做到极致

诞生动机：Triton 虽然好用，但在需要精细控制流水线、warp 特化、异步拷贝（TMA）等场景时，抽象仍不够。TileLang 由微软等团队推动，目标是提供更贴近硬件的 tile 级 DSL，让用户能显式表达数据搬运与计算的重叠。

核心特性：

• 以 tile 为核心抽象，显式区分全局内存、共享内存、寄存器层级。
• 支持显式流水线（pipeline）、异步拷贝、warp 特化等高级调度。
• 语法接近 Python，但提供更细粒度的控制原语。
• 面向高性能算子（GEMM、Attention）的自动调优与代码生成。

定位：介于 Triton 与手写 CUDA 之间，追求"可控性 + 生产力"的平衡，适合需要榨干硬件性能的算子开发者。

2.3 Mojo：为 AI 重造一门语言

诞生动机：Chris Lattner（LLVM、Swift、Clang 之父）创立 Modular，推出 Mojo。动机是 Python 生态好用但性能不足，C++ 性能好但开发效率低，于是想造一门兼具 Python 语法与 C 级性能的系统级语言。

核心特性：

• 语法高度兼容 Python，但引入强类型、所有权（ownership）、值语义。
• 通过 MLIR 编译，可针对 CPU、GPU（NVIDIA/AMD）等多种后端。
• 支持 fn（严格类型、高性能）与 def（动态）两种函数。
• 可直接调用 Python 生态，逐步迁移。

fn vector_add(a: DTypePointer[DType.float32], b: DTypePointer[DType.float32], c: DTypePointer[DType.float32], n: Int):
for i in range(n):
c[i] = a[i] + b[i]

定位：不是单纯的 GPU DSL，而是"AI 时代的系统编程语言"，GPU 编程只是其能力之一。

2.4 cuTile：NVIDIA 官方的 tile 级 Python DSL

诞生动机：面对 Triton 在 Python GPU 编程上的成功，NVIDIA 需要自己的官方方案。cuTile 是 NVIDIA 推出的基于 tile 抽象的 Python DSL，深度集成 CUDA 生态，目标是让 Python 开发者以官方支持的方式写高性能 kernel。

核心特性：

• 官方出品，与 CUDA 工具链、Nsight 等深度集成。
• tile 级抽象，编译器负责底层映射。
• 与 CuTe/CUTLASS 生态有潜在协同。
• 面向 NVIDIA 硬件做针对性优化。

定位：NVIDIA 对 Triton 的"官方回应"，强调生态一致性与硬件亲和性。

2.5 CuTe（C++ 模板库）：CUTLASS 的布局与算法内核

诞生动机：CUTLASS 是 NVIDIA 的高性能 GEMM 模板库，但其早期代码对布局（layout）、线程映射的表达非常繁琐。CuTe 作为 CUTLASS 3.x 的核心组件被提炼出来，用编译期模板元编程统一表达"数据布局 + 算法"。

核心特性：

• 核心概念：Layout（形状与步长的组合）、Tensor（数据 + 布局）、TiledCopy / TiledMMA（分块拷贝与矩阵乘）。
• 全部在编译期完成布局推导，零运行时开销。
• 用类型系统表达线程与数据的映射关系。
• 是 CUTLASS 3.x 的底层基础。

// CuTe 风格：用 Layout 描述 (M, N) 矩阵的分块布局
using namespace cute;
auto layout = make_layout(make_shape(Int<128>{}, Int<128>{}),
make_stride(Int<128>{}, Int<1>{}));

定位：C++ 模板元编程路线的巅峰之作，性能极致但学习曲线陡峭。

2.6 ThunderKittens：为 AI 算子设计的"极简 tile 库"

诞生动机：斯坦福 Hazy Research 团队推出 ThunderKittens，动机是现有方案要么抽象太高（丢失性能），要么太底层（开发慢）。他们提出"tile 就是一切"的理念，用极简的 C++/CUDA 头文件库让开发者快速写出接近峰值的 AI 算子。

核心特性：

• 极简 API：核心就是 tile 类型与少量操作（load/store/mma/exp 等）。
• 强调"每个 SM 一个 tile"的编程模型，简化线程管理。
• 针对现代 GPU（Hopper 等）的 TMA、WGMMA 等特性优化。
• 头文件库形式，易于集成。

定位：研究友好、上手快，适合快速实现和实验新算子。

2.7 CuTe DSL：把 CuTe 带到 Python

诞生动机：CuTe 的 C++ 模板虽然强大，但编译慢、调试难、门槛高。CuTe DSL 的目标是保留 CuTe 的布局抽象思想，同时提供 Python 层的易用性，让更多人能用上 CuTe 的表达能力。

核心特性：

• 用 Python 表达 CuTe 的 Layout/Tensor 概念。
• 降低编译期模板带来的开发负担。
• 与 CUTLASS/CuTe 生态衔接。
• 面向需要精细布局控制但不想写模板元编程的用户。

定位：CuTe 思想的"Python 化"延伸，兼顾表达力与易用性。

3. 横向对比

4. 时间线与演进脉络

#mermaid-svg-SStYtldIslVq7Gol{font-family:\”trebuchet ms\”,verdana,arial,sans-serif;font-size:16px;fill:#333;}@keyframes edge-animation-frame{from{stroke-dashoffset:0;}}@keyframes dash{to{stroke-dashoffset:0;}}#mermaid-svg-SStYtldIslVq7Gol .edge-animation-slow{stroke-dasharray:9,5!important;stroke-dashoffset:900;animation:dash 50s linear infinite;stroke-linecap:round;}#mermaid-svg-SStYtldIslVq7Gol .edge-animation-fast{stroke-dasharray:9,5!important;stroke-dashoffset:900;animation:dash 20s linear infinite;stroke-linecap:round;}#mermaid-svg-SStYtldIslVq7Gol .error-icon{fill:#552222;}#mermaid-svg-SStYtldIslVq7Gol .error-text{fill:#552222;stroke:#552222;}#mermaid-svg-SStYtldIslVq7Gol .edge-thickness-normal{stroke-width:1px;}#mermaid-svg-SStYtldIslVq7Gol .edge-thickness-thick{stroke-width:3.5px;}#mermaid-svg-SStYtldIslVq7Gol .edge-pattern-solid{stroke-dasharray:0;}#mermaid-svg-SStYtldIslVq7Gol .edge-thickness-invisible{stroke-width:0;fill:none;}#mermaid-svg-SStYtldIslVq7Gol .edge-pattern-dashed{stroke-dasharray:3;}#mermaid-svg-SStYtldIslVq7Gol .edge-pattern-dotted{stroke-dasharray:2;}#mermaid-svg-SStYtldIslVq7Gol .marker{fill:#333333;stroke:#333333;}#mermaid-svg-SStYtldIslVq7Gol .marker.cross{stroke:#333333;}#mermaid-svg-SStYtldIslVq7Gol svg{font-family:\”trebuchet ms\”,verdana,arial,sans-serif;font-size:16px;}#mermaid-svg-SStYtldIslVq7Gol p{margin:0;}#mermaid-svg-SStYtldIslVq7Gol .label{font-family:\”trebuchet ms\”,verdana,arial,sans-serif;color:#333;}#mermaid-svg-SStYtldIslVq7Gol .cluster-label text{fill:#333;}#mermaid-svg-SStYtldIslVq7Gol .cluster-label span{color:#333;}#mermaid-svg-SStYtldIslVq7Gol .cluster-label span p{background-color:transparent;}#mermaid-svg-SStYtldIslVq7Gol .label text,#mermaid-svg-SStYtldIslVq7Gol span{fill:#333;color:#333;}#mermaid-svg-SStYtldIslVq7Gol .node rect,#mermaid-svg-SStYtldIslVq7Gol .node circle,#mermaid-svg-SStYtldIslVq7Gol .node ellipse,#mermaid-svg-SStYtldIslVq7Gol .node polygon,#mermaid-svg-SStYtldIslVq7Gol .node path{fill:#ECECFF;stroke:#9370DB;stroke-width:1px;}#mermaid-svg-SStYtldIslVq7Gol .rough-node .label text,#mermaid-svg-SStYtldIslVq7Gol .node .label text,#mermaid-svg-SStYtldIslVq7Gol .image-shape .label,#mermaid-svg-SStYtldIslVq7Gol .icon-shape .label{text-anchor:middle;}#mermaid-svg-SStYtldIslVq7Gol .node .katex path{fill:#000;stroke:#000;stroke-width:1px;}#mermaid-svg-SStYtldIslVq7Gol .rough-node .label,#mermaid-svg-SStYtldIslVq7Gol .node .label,#mermaid-svg-SStYtldIslVq7Gol .image-shape .label,#mermaid-svg-SStYtldIslVq7Gol .icon-shape .label{text-align:center;}#mermaid-svg-SStYtldIslVq7Gol .node.clickable{cursor:pointer;}#mermaid-svg-SStYtldIslVq7Gol .root .anchor path{fill:#333333!important;stroke-width:0;stroke:#333333;}#mermaid-svg-SStYtldIslVq7Gol .arrowheadPath{fill:#333333;}#mermaid-svg-SStYtldIslVq7Gol .edgePath .path{stroke:#333333;stroke-width:2.0px;}#mermaid-svg-SStYtldIslVq7Gol .flowchart-link{stroke:#333333;fill:none;}#mermaid-svg-SStYtldIslVq7Gol .edgeLabel{background-color:rgba(232,232,232, 0.8);text-align:center;}#mermaid-svg-SStYtldIslVq7Gol .edgeLabel p{background-color:rgba(232,232,232, 0.8);}#mermaid-svg-SStYtldIslVq7Gol .edgeLabel rect{opacity:0.5;background-color:rgba(232,232,232, 0.8);fill:rgba(232,232,232, 0.8);}#mermaid-svg-SStYtldIslVq7Gol .labelBkg{background-color:rgba(232, 232, 232, 0.5);}#mermaid-svg-SStYtldIslVq7Gol .cluster rect{fill:#ffffde;stroke:#aaaa33;stroke-width:1px;}#mermaid-svg-SStYtldIslVq7Gol .cluster text{fill:#333;}#mermaid-svg-SStYtldIslVq7Gol .cluster span{color:#333;}#mermaid-svg-SStYtldIslVq7Gol div.mermaidTooltip{position:absolute;text-align:center;max-width:200px;padding:2px;font-family:\”trebuchet ms\”,verdana,arial,sans-serif;font-size:12px;background:hsl(80, 100%, 96.2745098039%);border:1px solid #aaaa33;border-radius:2px;pointer-events:none;z-index:100;}#mermaid-svg-SStYtldIslVq7Gol .flowchartTitleText{text-anchor:middle;font-size:18px;fill:#333;}#mermaid-svg-SStYtldIslVq7Gol rect.text{fill:none;stroke-width:0;}#mermaid-svg-SStYtldIslVq7Gol .icon-shape,#mermaid-svg-SStYtldIslVq7Gol .image-shape{background-color:rgba(232,232,232, 0.8);text-align:center;}#mermaid-svg-SStYtldIslVq7Gol .icon-shape p,#mermaid-svg-SStYtldIslVq7Gol .image-shape p{background-color:rgba(232,232,232, 0.8);padding:2px;}#mermaid-svg-SStYtldIslVq7Gol .icon-shape .label rect,#mermaid-svg-SStYtldIslVq7Gol .image-shape .label rect{opacity:0.5;background-color:rgba(232,232,232, 0.8);fill:rgba(232,232,232, 0.8);}#mermaid-svg-SStYtldIslVq7Gol .label-icon{display:inline-block;height:1em;overflow:visible;vertical-align:-0.125em;}#mermaid-svg-SStYtldIslVq7Gol .node .label-icon path{fill:currentColor;stroke:revert;stroke-width:revert;}#mermaid-svg-SStYtldIslVq7Gol :root{–mermaid-font-family:\”trebuchet ms\”,verdana,arial,sans-serif;}

2017-2019CUTLASS 诞生手写模板 GEMM

2021Triton 发布Python tile 级 DSL

2022-2023CuTe 提炼CUTLASS 3.x 布局抽象

2023Mojo 发布AI 系统语言

2024ThunderKittens极简 tile 库

2024-2025TileLang精细 tile DSL

2025cuTileNVIDIA 官方 Python DSL

2025CuTe DSLCuTe 的 Python 化

每个方案解决的特殊问题：

• CUTLASS（2017-2019）：解决"高性能 GEMM 需要反复手写"的问题，用模板库固化最佳实践。
• Triton（2021）：解决"Python 用户写不了 GPU kernel"的问题，用 tile 抽象 + 编译器自动化。
• CuTe（2022-2023）：解决"CUTLASS 布局表达太繁琐"的问题，用编译期类型系统统一布局与算法。
• Mojo（2023）：解决"Python 性能不足、C++ 效率低"的问题，重造一门 AI 系统语言。
• ThunderKittens（2024）：解决"现有方案要么太抽象要么太底层"的问题，用极简 tile 库快速写算子。
• TileLang（2024-2025）：解决"Triton 抽象不够精细"的问题，提供显式流水线与硬件控制。
• cuTile（2025）：解决"NVIDIA 缺少官方 Python tile DSL"的问题，官方背书 + 生态集成。
• CuTe DSL（2025）：解决"CuTe 模板编译慢、门槛高"的问题，Python 化布局抽象。

5. 它们之间如何交互与关联

这些方案并非孤立，而是形成了一张相互借鉴、层层递进的网络：

• CuTe 是 CUTLASS 的内核：CUTLASS 3.x 的 GEMM 实现完全建立在 CuTe 的 Layout/Tensor 之上，CuTe 是"库的库"。
• CuTe DSL 是 CuTe 的 Python 门面：它复用 CuTe 的布局思想，但把编译期模板换成 Python 表达，降低门槛。
• cuTile 与 CuTe/CUTLASS 同属 NVIDIA 生态：cuTile 作为官方 Python DSL，很可能在底层复用 CuTe 的布局与 CUTLASS 的 kernel 生成能力。
• Triton 与 TileLang 是竞争与借鉴关系：TileLang 在 tile 抽象上比 Triton 更细，二者都走"Python + 编译器"路线，TileLang 可视为 Triton 思路的"精细化演进"。
• ThunderKittens 与 CuTe 思想相通：都强调 tile/布局为核心，但 ThunderKittens 用极简头文件库降低使用门槛，CuTe 用模板元编程追求极致表达。
• Mojo 是另一条独立赛道：它不直接复用上述 DSL，而是从语言层面重造，但可通过 MLIR 与 CUDA/GPU 后端对接，长期可能与这些 DSL 形成"语言 vs 库"的互补。
• Triton 是生态粘合剂：PyTorch torch.compile、vLLM、FlashAttention 等大量项目基于 Triton，使它成为事实上的"Python GPU 编程标准"，其他方案往往需要与 Triton 生态互操作。

一句话总结：CuTe 提供布局理论，CUTLASS 提供工程实践，Triton 提供 Python 生产力，TileLang/ThunderKittens 提供更细或更简的 tile 控制，cuTile/CuTe DSL 是 NVIDIA 把 CuTe 思想推向 Python 的官方动作，Mojo 则是从语言层面另起炉灶的长期赌注。

6. 选型建议

• 想快速写算子、生态成熟 → Triton。
• 需要精细控制流水线、榨干性能 → TileLang 或 CuTe。
• 想用 Python 但要好布局控制 → CuTe DSL 或 cuTile。
• 研究新算子、追求上手速度 → ThunderKittens。
• 长期押注 AI 系统语言 → Mojo。
• 深度 CUTLASS 用户 → CuTe（C++）。

7. 结语

GPU 编程 DSL 的百花齐放，本质是"性能"与"生产力"这对矛盾的持续博弈。从 CUTLASS 的模板工程，到 Triton 的 Python 抽象，再到 CuTe 的布局理论、Mojo 的语言重造，每一代方案都在试图把"手写 CUDA 的痛苦"再降低一点，同时不丢掉硬件性能。理解它们的时间线与相互关系，比记住每个 API 更重要——因为下一个方案，很可能就是这些思想的又一次重组。