NVIDIA GPU 编程生态全景 — 从 PTX 到 CuTe DSL 的抽象阶梯

如果你在 2026 年从外部看 NVIDIA 的 GPU 编程生态,会觉得它越来越乱:PTX、CUDA C++、CUTLASS、CuTe、Triton、CuTe DSL、CuTile — 这一堆名字反复在论文、博客、PR description 里冒出来,新的还在不断加。但只要把它们放进一个从硬件到人类的连续阶梯里,整个版图会瞬间清晰下来 — 越往下越接近机器,性能上限越高,但写起来越痛苦;越往上越像写普通 Python,迭代越快,但你会逐渐失去对硬件的控制。

这篇文章用一根这样的阶梯作为骨架,讲清楚:这些工具各自在哪一层、它们的真实关系是什么、真实生产系统(以 vLLM 和 SGLang 为例)怎么把它们混着用,以及当你需要自己写一段 GPU 代码时,该如何在它们之间做选择。

抽象阶梯 — 从汇编到 Python DSL 的五层版图

先把这一堆名字按抽象级别从底到顶排一下。下面这张图是后面所有讨论的「速查表」 — 每一层的定位、典型代表、和它在整个生态里的角色,先记住这个骨架。

按层往下走一遍:

Layer 5 · Python DSL(Triton / CuTile / CuTe DSL)

抽象最高的一层。共同特征是:用 Python 写,有自己独立的编译器,直接把高层代码下到 PTX,不经过 CUDA C++。这是 2019 年以来最重要的范式变化 — 把「写 GPU kernel」从 C++ 的世界里拉了出来。

• Triton(OpenAI,2019) — 当下绝对的主流入口。核心理念是 block/tile 级编程而非 thread 级:你以 tile 为单位调度,内存合并、shared memory 同步、warp 分配这些底层细节统统交给编译器。
• CuTe DSL(NVIDIA,2025) — 给底层 CuTe 的 layout 代数套了个 Python 前端,目标是「用 Python 的开发体验拿到接近 CUTLASS C++ 的性能」。
• CuTile / CUDA Tile IR(NVIDIA,2025) — CUDA 13.1 引入的官方 tile 编程模型,NVIDIA「自己下场做 Triton」的回应。

Layer 4 · C++ 模板高性能层(CUTLASS / CuTe / ThunderKittens)

这一层的动机是:纯手写 CUDA C++ 写矩阵乘 / attention 这种主力算子太累、太容易出错,于是用 C++ 模板把高性能模式封装成可复用的抽象。

• CUTLASS(NVIDIA) — CUDA Templates for Linear Algebra Subroutines,命名上致敬老牌闭源库 cuBLAS,可以理解为「开源、可定制版的 BLAS」。
• CuTe(NVIDIA,CUTLASS 3.0 起) — CUDA Tensors,CUTLASS 内部那套描述「数据怎么摆 + 线程怎么映射到数据」的 layout 代数。它是 CUTLASS 的地基,不是并列的另一个库。
• ThunderKittens(斯坦福 Hazy Research,2024) — 反其道而行,用一小套「有主见的」抽象看能走多快多广。在 H100 上对 GEMM / attention 能打平甚至超过 CUTLASS 版的 FlashAttention-3,代码却短得多。

Layer 3 · CUDA C++(.cu)

地基。2012 年以来唯一的选择,至今仍是整个生态的根。模型是 SIMT — 你显式地按 grid → block(CTA)→ warp → thread 来组织并行,自己管理寄存器、shared memory、tiling、同步。所有上层框架最终都会 fallback 到它,所有库都是用它写的。

日常说的「CUDA」很多时候就是指这门语言,虽然这个词更宽泛地也指 NVIDIA 整个 GPU 计算平台(语言 + 编译器 + 驱动 + 库的总称)。对绝大多数应用开发者来说,你其实是通过 cuBLAS / cuDNN 这些预编译库间接用到它,而不是自己写。

Layer 2 · PTX

Parallel Thread eXecution — NVIDIA 的虚拟 ISA(中间表示)。所有上层入口(CUDA C++、Triton、CuTe DSL)最后都要变成 PTX,再由驱动里的 ptxas 编成特定架构的真实机器码 SASS。

实际工程里几乎没人从头用 PTX 写完整 kernel。它的真实用途是:在 .cu 里通过 asm volatile 内联几条指令,去碰那些 C++ 层暴露不出来的硬件特性 — 比如某条特定的异步拷贝、某个 cache hint、或者新硬件刚出、编译器还没支持的指令。PTX 是补丁工具,不是开发语言。

Layer 1 · SASS

Streaming ASSembly — 特定架构的真实机器码。你不会写它,通常只在做极致性能分析或逆向时才看。

把这五层在心里铺开,这条阶梯就有了。后面再补两块上下文让它能直接用:一块讲编译怎么把这些上层入口收口到 PTX,另一块讲这条阶梯之外那一组 NVIDIA 平行存在的闭源库。

编译路径 — 三条前端汇到 PTX

阶梯顶端的 Layer 3 / 4 / 5 有三种不同的源语言入口,但所有路径都汇到 PTX 这同一个虚拟 ISA,再由驱动里的 ptxas 编成 SASS:

• CUDA C++(.cu) — 包括所有用 CUTLASS / CuTe 模板写出来的代码,走 nvcc 这条传统编译链。
• Triton — 有自己基于 MLIR 的编译器,把 Python 直接 lower 到 PTX。
• CuTe DSL — 同样基于 MLIR,把 Python 写出来的 CuTe layout 代数 lower 到 PTX。

注意 Triton 和 CuTe DSL 都不会生成中间的 .cu 文件 — 它们是和 CUDA C++ 并列的源语言入口,而不是建在 CUDA C++ 之上。这也是为什么 Triton 能在 PyTorch 里跟 cuBLAS / cuDNN、CUTLASS 平行共存:它们最终都变成 GPU 能跑的同一种机器码,只是源头不同。

阶梯之外 — NVIDIA 平行存在的闭源库

上面这五层都属于「源语言」的世界 — 你写代码,然后编译。但 NVIDIA 还另外维护一套用法完全不同的东西:闭源、预编译、你只调一行 API的库。最常见的两个:

• cuBLAS(CUDA Basic Linear Algebra Subroutines)— 通用线性代数,GEMM / BLAS 这套。
• cuDNN(CUDA Deep Neural Network library)— 深度学习专用算子,卷积、池化、归一化、attention 等。

你调 PyTorch 时,矩阵乘默认走 cuBLAS,卷积默认走 cuDNN — 它们是 PyTorch 性能多年来的地基。区别于阶梯上的工具:这两个库是黑盒,你不能改、看不见内部,但只要参数对就拿到 NVIDIA 多年调优出来的性能。

CUTLASS 正是 NVIDIA 给这条闭源生态做的「开源积木版本」 — 名字本身就在致敬 cuBLAS(CUDA Templates for Linear Algebra Subroutines)。功能领域重叠(都能做 GEMM),用法完全不同:

CUTLASS / CuTe 之所以值得作为一层单独存在,是因为它们封装的不只是「写起来更顺眼的 C++ 模板」,而是 NVIDIA 多年沉淀的算法资产 — bank conflict 怎么避、TMA 怎么用、warp 怎么排、pipelining 怎么做,这些工程经验通过 CuTe 的 layout 代数和 CUTLASS 的模板参数变得可表达、可组合。也正因为如此,主力算子(GEMM、attention)真要榨极限,就绕不开这一层。一个比较贴切的类比是:CuTe / CUTLASS 之于 CUDA C++,更像 NumPy 之于手写 C 循环 — 改变的是你能多快达到多高的性能、需要懂多少。

代码对比 — 同一个 GEMM 的五种写法

讲到这里,最直观的体感是把同一件事(矩阵乘法 $C = A \times B$)用五种方式写一遍。注意重点不是看懂每一行,而是感受它们要求你操心的东西完全不在一个量级。

写法 1 · cuBLAS — 黑盒调用,你只说「做什么」

// 创建句柄,调一个函数,完事。你不关心内部怎么算的。
cublasHandle_t handle;
cublasCreate(&handle);

float alpha = 1.0f, beta = 0.0f;
// C = alpha * A * B + beta * C
cublasSgemm(handle,
            CUBLAS_OP_N, CUBLAS_OP_N,   // A、B 不转置
            M, N, K,                     // 矩阵尺寸
            &alpha, dA, M,               // 输入 A
            dB, K,                       // 输入 B
            &beta, dC, M);               // 输出 C

cublasDestroy(handle);

你操心的:矩阵尺寸、转不转置、放哪个指针。就这些。怎么 tile、用不用 Tensor Core、线程怎么分 — NVIDIA 全替你定了,你看不见也改不了。这就是「黑盒」。cuDNN 几乎一样,只是函数换成 cudnnConvolutionForward(...) 之类。

写法 2 · 纯 CUDA C++ — 你自己写,但写得朴素

// 每个线程负责算 C 里的一个元素
__global__ void gemm_naive(float* A, float* B, float* C, int M, int N, int K) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < M && col < N) {
        float sum = 0.0f;
        for (int k = 0; k < K; ++k)        // 自己写内积循环
            sum += A[row * K + k] * B[k * N + col];
        C[row * N + col] = sum;
    }
}
// 启动:gemm_naive<<<grid, block>>>(dA, dB, dC, M, N, K);

你操心的:线程怎么映射到元素、循环怎么写。注意:这段能跑,但慢得要命 — 没用 shared memory、没用 Tensor Core、内存访问也不优化。要写快,得手动加 tiling、SMEM、bank conflict 处理…… 几百上千行,而且换代硬件要重写。这就是为什么有了 CUTLASS / CuTe。

写法 3 · CuTe — 用 layout 代数描述「数据怎么摆 + 谁算哪块」

CuTe 的精髓是:它给你一套词汇,让你声明式地描述数据布局和线程映射,而不用手写一堆索引计算。核心概念是 Layout(形状 + 步长)和 Tensor(数据 + Layout)。

using namespace cute;

// 把一块裸内存「包装」成带 layout 的 Tensor
// make_shape(M,K) 是形状,make_stride 描述内存里怎么摆
Tensor mA = make_tensor(make_gmem_ptr(A), make_shape(M, K), make_stride(_1{}, M));
Tensor mB = make_tensor(make_gmem_ptr(B), make_shape(N, K), make_stride(_1{}, N));
Tensor mC = make_tensor(make_gmem_ptr(C), make_shape(M, N), make_stride(_1{}, M));

// 声明「怎么切块」:每个 block 处理 128×128×8 的一小块
auto block_tile = make_shape(Int<128>{}, Int<128>{}, Int<8>{});

// 用 layout 把全局矩阵「切」成属于当前 block 的那一块
Tensor gA = local_tile(mA, block_tile, ...);
Tensor gB = local_tile(mB, block_tile, ...);
Tensor gC = local_tile(mC, block_tile, ...);

// 声明一个「tiled MMA」:用哪条 Tensor Core 指令、warp 怎么排
TiledMMA mma = make_tiled_mma(SM80_16x8x8_F32F16F16F32_TN{}, ...);

// 然后在被 layout 描述好的 tile 上做乘累加
cute::gemm(mma, gA, gB, gC);

看出区别了吗?没有手写 A[row*K+k] 这种索引,而是用 make_shape / make_stride / local_tile 这套代数去声明「数据长什么样、怎么切」,再让 cute::gemm 按声明的 layout 去算。这就是 layout 代数 — 它管的是数据编排,不是帮你调一个现成函数。

写法 4 · CUTLASS — 用模板把 CuTe 零件拼成完整 kernel

CUTLASS 在 CuTe 之上,让你通过填模板参数来组装一个生产级 kernel,不用自己从 CuTe 原子一点点搭。

using namespace cutlass::gemm;

// 不写循环,而是「声明」这个 GEMM 的每一个维度该用什么配置
using Gemm = device::GemmUniversal<
    cutlass::half_t, cutlass::layout::RowMajor,    // A: 类型 + 布局
    cutlass::half_t, cutlass::layout::ColumnMajor, // B
    float,           cutlass::layout::RowMajor,    // C
    float,                                         // 累加用 float
    cutlass::arch::OpClassTensorOp,                // 用 Tensor Core
    cutlass::arch::Sm90,                           // 目标架构 Hopper
    Shape<_128,_128,_64>,                          // block tile 大小
    Shape<_64, _64, _64>                           // warp tile 大小
    /* 还有 epilogue、pipeline 级数等一堆可填参数 */ >;

Gemm gemm_op;
gemm_op({M, N, K}, {dA, lda}, {dB, ldb}, {dC, ldc}, {alpha, beta}); // 启动

你操心的:类型、布局、目标架构、tile 切多大、warp 分多少、用不用 Tensor Core…… 全是性能旋钮,你来定。CUTLASS 负责把这些选择编译成一个高度优化的 kernel。一个最简的 Hopper WGMMA + TMA GEMM 例子大约 100 行。

写法 5 · Triton — Python tile 级,把硬件细节藏起来

import triton
import triton.language as tl

@triton.jit
def gemm_kernel(A, B, C, M, N, K,
                stride_am, stride_ak, stride_bk, stride_bn, stride_cm, stride_cn,
                BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr, BLOCK_K: tl.constexpr):
    # 当前 block 负责 C 的一小块 [BLOCK_M, BLOCK_N]
    pid_m = tl.program_id(0)
    pid_n = tl.program_id(1)

    # 生成 tile 内的索引向量(不是 thread,是 block 内的一片元素)
    offs_m = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
    offs_n = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
    offs_k = tl.arange(0, BLOCK_K)

    acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)

    # 沿 K 维度迭代,每次加载 BLOCK_K 宽度的 tile,做 tile-level matmul
    for k in range(0, K, BLOCK_K):
        a = tl.load(A + offs_m[:, None] * stride_am + (k + offs_k[None, :]) * stride_ak)
        b = tl.load(B + (k + offs_k[:, None]) * stride_bk + offs_n[None, :] * stride_bn)
        acc += tl.dot(a, b)   # 编译器自动 lower 到 Tensor Core MMA

    tl.store(C + offs_m[:, None] * stride_cm + offs_n[None, :] * stride_cn, acc)

最关键的认知:整段代码完全没有 threadIdx,没有 __shared__,没有任何同步原语。你以 tile 为单位思考 — 加载一个 tile、做一次 tile-level matmul、累加 — Triton 编译器自动决定:线程怎么映射、shared memory 怎么 staging、Tensor Core 指令怎么发、内存合并怎么做。这是它和 CUDA C++ 最大的范式差异。

把五种写法放在一起对照:

最关键的直觉是:cuBLAS/cuDNN 你说「我要矩阵乘」,内部一概不管;CuTe/CUTLASS/Triton 你在描述 kernel 内部该怎么干,只是用它们给的高级词汇(layout 代数 / 模板参数 / tile 抽象)而不是手写裸索引。

真实案例 — FlashAttention 与 vLLM / SGLang

用两个真实项目把这条阶梯走一遍。FlashAttention 让你看到「纵向」的演化 — 同一个算子,随着硬件代际推进,在阶梯上一路下移又部分上移,每一代都站在「当时性能最高那一层」。vLLM 和 SGLang 让你看到「横向」的混用 — 同一时刻、同一系统里,不同算子被分配到不同层,主力算子去最优库、缝隙用 Triton 兜底。两个角度合在一起,基本就是工业界今天 GPU 编程的真实形态。

FlashAttention — 一部 NVIDIA 编程范式演进史

如果只能选一个项目来理解整个生态怎么演化的,那一定是 FlashAttention。它每一代换硬件、就换一套编程方式,正好把主流路线走了个遍 — 可以把它当作「活化石」来读。

值得专门点出来的是 FA4 这一步:用 Python 写出和 C++ 同等性能 — 这正是 CuTe DSL 的设计目标,给「Python 生态里做主力算子 + 紧贴最新硬件 + 要榨到 95% 以上」这一类场景提供了一条不必回退到 C++ 的路径。它和 Triton 在 Python 生态里各占不同位置:Triton 把 layout / TMA / WGMMA 这些硬件细节藏起来,你不需要懂那些就能写,代价是性能上限封在 80-95%;CuTe DSL 把这些细节暴露出来,你必须懂那些才能用,换来的是 95%+ 的控制力。一个面向日常的自定义算子,一个面向顶部那 5-15% 的极限场景。FA4 选了后者,因为 attention 在 Hopper / Blackwell 上必须榨到峰值,Triton 不够,而纯 C++ 模板又太累。

vLLM 和 SGLang — 一个系统里同时用上整条阶梯

vLLM 和 SGLang 的核心策略是做一个调度层:把活分派给一堆现成的最优后端,只在现成库覆盖不到的缝隙里才自己写 kernel。它们更像指挥,而不是演奏者。

vLLM 官方 README 自己列的 kernel 清单就说明一切了:优化的 attention kernel 包括 FlashAttention、FlashInfer、TRTLLM-GEN、FlashMLA 和 Triton;优化的 GEMM/MoE kernel 用 CUTLASS、TRTLLM-GEN、CuTeDSL;并用 torch.compile 做自动 kernel 生成和图级变换。一句话里就出现了至少六七种不同的 kernel 来源。

SGLang 的策略几乎一样,内置的 attention 后端比 vLLM 还杂,光 MLA(DeepSeek 那种注意力)一项就有一堆:FlashInfer MLA、FlashMLA、Cutlass MLA、TRTLLM MLA 各自对应不同的 page_size。自动选后端的逻辑也很能说明问题 — Hopper 默认用 fa3,Blackwell 默认用 trtllm_mha,其他架构默认用 flashinfer,不可用时回退到 triton。看这个回退链你就懂了:专用库优先 → 通用库次之 → Triton 兜底。

那这两个引擎自己写了什么?分三类:

1. 调度/抽象层(它们最核心的自研价值) — 一套 attention/GEMM backend 的可插拔抽象 + 运行时自动选最优后端的逻辑。这是它们真正的工程壁垒,不是 kernel 本身。
2. 自研 Triton kernel — 集中在两处:(a) 适配自己独特数据结构的算子(paged KV cache 相关的 PagedAttention、KV cache 抓取等),这些外部库不会替你写;(b) 跨硬件的兜底后端,保证哪都能跑。
3. 主力高性能算子 — 几乎全部外包给 FlashAttention、FlashInfer、CUTLASS、CuTe DSL、TRT-LLM,绝不自己重造。

这对前面整个心智模型是一次硬核印证:主力算子(attention 峰值、GEMM)→ C++ 的 CUTLASS / CuTe、专用库(极限性能绕不开 C++ 库);周边 / 适配 / 兜底算子 → Triton(Python 里快速写、够用、可移植)。两个顶级生产系统都没有用单一方案,而是按算子的重要性和场景,把不同层的工具混着用 — 这就是真实世界 GPU 编程的样子:不是选一条路,而是在一个系统里同时用上整条阶梯。

决策地图 — 怎么选择技术路线

把前面所有内容压缩成一张可以照着用的决策地图。核心问题串起来 — 每往下一层,先问:「上一层真的不够吗?」 不够才往下走,因为每下一层,开发成本和所需硬件知识都陡增。

三条心法,帮你真正用好这张图:

1. 默认往上,不往下。永远先试最省事的那层,实测不够了再往下走一格。不要因为「听说 CUTLASS 快」就一上来写 CUTLASS — 大多数时候 Triton 甚至现成库就够了,过早优化是浪费生命。
2. 不是选一条路,而是一个系统里混用。这是 vLLM/SGLang 教给我们的最大一课:主力算子(attention/GEMM)用专用库和 CUTLASS,周边和适配算子用 Triton,兜底也用 Triton。真实系统是整条阶梯同时在用,你的目标是给每个算子选对它该在的那一层。
3. 区分「调用」和「写」、「黑盒」和「积木」。cuBLAS / cuDNN 是黑盒(调 API、不可改);CUTLASS / CuTe / Triton 是积木(写代码、可定制)。FlashAttention 这些是别人用积木搭好、又能当黑盒调的成品。想清楚你要的是「拿来用」还是「自己搭」,就不会在工具选择上纠结。

结语 — 一句话收口

NVIDIA GPU 编程的版图看起来名字一大堆 — PTX、CUDA C++、CUTLASS、CuTe、Triton、CuTe DSL、CuTile…… 但本质上它就是一条从硬件到人类的连续阶梯:底层 PTX / SASS 是出口,中层 CUDA C++ + CUTLASS / CuTe 是地基和极限工具,顶层 Triton / CuTe DSL 是 Python 时代的快速入口。所有路径最终都汇到同一条出口(PTX → SASS),区别只在你从哪一层接入。

能用现成库就别写 kernel;非写不可,Python 默认 Triton、C++ 默认先调库后 CUTLASS;要榨极限上 CUTLASS / CuTe(或其 Python 版 CuTe DSL);PTX 只是哪里需要补哪里的补丁。

逐层下探、实测驱动、按算子混用 — 把这条阶梯装进心智模型,你就能拿它去读任何项目的 kernel 目录结构,看清楚它在每一层各放了什么、为什么这么放。

参考资料 — 官方文档 · 代表项目 · 社区讨论

NVIDIA 官方文档

• CUDA Toolkit — CUDA 文档主页、PTX ISA 参考手册、CUDA C++ Programming Guide
• CUTLASS — NVIDIA/cutlass GitHub、CUTLASS 3.x 文档、CuTe quick start guide
• CUDA Tile IR / CuTile — CUDA 13.1 发布说明、CuTile 编程模型介绍

Triton 与 PyTorch 集成

• Triton 项目 — triton-lang/triton GitHub、Triton 编程模型论文(Tillet et al. 2019)
• TorchInductor — torch.compile 文档、Triton 作为 PyTorch 2.x 默认 codegen 后端的设计说明
• Gluon — Triton 团队发布的「低一层」DSL,暴露 tile layout / 内存分配等底层细节

FlashAttention 演进

• FlashAttention 1 — Dao et al.(2022)“FlashAttention: Fast and Memory-Efficient Exact Attention with IO-Awareness”
• FlashAttention 2 — Dao(2023)“FlashAttention-2: Faster Attention with Better Parallelism and Work Partitioning”
• FlashAttention 3 — Shah et al.(2024)“FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-precision”,用 Hopper WGMMA + TMA 达到 ~740 TFLOPS
• FlashAttention 4 — Dao-AILab/flash-attention 的 CuTe DSL 实现,SM90 / SM100 双覆盖

代表项目

• vLLM — vllm-project/vllm,多 attention backend 抽象与自动选择
• SGLang — sgl-project/sglang,按硬件自动选 backend 的回退链
• FlashInfer — flashinfer-ai/flashinfer,推理引擎共用的 kernel 路由层
• Unsloth — unslothai/unsloth,用 Triton 重写流行 LLM 的训练 / LoRA kernel
• Liger Kernel — linkedin/Liger-Kernel,LinkedIn 出品的训练用 Triton kernel 集
• ThunderKittens — HazyResearch/ThunderKittens,斯坦福 Hazy Research 的「warp 中心」C++ 嵌入库
• llm.c — karpathy/llm.c,Karpathy 的纯 C/CUDA 训练实现,教学价值极高

行业综述与博客

• 35 Modern GPU Kernel Frameworks — 一篇梳理 GPU kernel 框架谱系的综述博客,把 CUDA C++ → Triton → CuTe → ThunderKittens → CuTile / CuTe DSL 的演化轨迹画得很清楚
• Tri Dao 博客与演讲 — FlashAttention 系列作者关于「为什么换工具」的第一手解释
• GPU MODE 社区 — Discord 社区,Triton / CUTLASS / CuTe 的高质量讨论密度最高的地方

本博客相关

• GPU 架构十年演化与 CUDA 编程模型的同步膨胀 — Pascal → Rubin 的硬件演化路线,以及 CUDA 编程模型如何从单层 thread 扩到五层(grid → cluster → block → warp → thread)
• AI 推理芯片光谱 — 通用 GPU 到模型刻片的七档专用化 — 光谱最左侧的 NVIDIA 路线深入拆解
• LLM 推理过程逐层拆解 — 理解为什么 attention / GEMM 是「主力算子」,以及它们在推理一次 forward 里所占的位置