Triton 和 CUDA 有什么区别?
什么时候用 Triton,什么时候还得手写 CUDA?

英文问法What's the difference between Triton and CUDA? · Why did you write this kernel in Triton instead of CUDA? · Is that the same Triton as the Triton Inference Server?

考过别家面经原题:遂原科技 AI Infra 27 实习一面,2026-05-01「是否考虑过使用 CUDA 替代 Triton 进行算子开发,选择 Triton 的原因」(牛客)· 快手 AI Infra 一面,面试 2026-09-28「是否用过Triton写kernel?解决了什么问题」(牛客,模板化的面经帖,可信度中等)· Unity 自己的材料:FlashVector 两个 Triton 都用到了,Case 1 改的是 Triton Inference Server,Case 2 的原话是 "fuses them into a single Triton kernel" · 没有 Unity 自己的面经原题

先弄清楚:每个概念是什么

Triton(写 kernel 的语言)
OpenAI 在 2021 年开源的 Python DSL 和 compiler,用来写 GPU kernel。这页讲的是它。
Triton Inference Server
NVIDIA 的 model server,负责接请求、组 batch、调模型,和上面那个只是重名。Unity 的 JD 里出现的 Triton 全是它。
Program
Triton 的基本单位,一个 program 处理一个 block 的数据。长度 256 的向量,BLOCK_SIZE 是 64,就有 4 个 program,各管 64 个元素。
Autotune
给几组候选参数,每组跑一遍,留下最快的。@triton.autotune 的候选里有 block 大小和 num_warps。

看图:两种抽象层次

CUDA:你写一个 thread 的逻辑 一个 thread block 你的代码是这一个 thread 的逻辑 你管 thread index 怎么对应数据 shared memory 和同步 访存是不是 coalesced thread block 怎么切 Triton:你写一个 block 的逻辑 一个 program load 一块 算 store 一块 内部的 thread 由 compiler 安排 你的代码处理一整个 block 的数据 compiler 管 program 内部的 thread 安排 coalescing 和 shared memory 你还要管 block 大小,program 之间怎么分工
图 1:CUDA 的代码是一个 thread 的逻辑,每个 thread 各跑一遍。Triton 的代码处理一整个 block 的数据,program 之间怎么分工仍然要自己写。

30 秒回答

Two things are called Triton. One is NVIDIA's Triton Inference Server, a model server. The other is OpenAI's Triton language, a Python DSL for GPU kernels, and that's the one I mean. In CUDA you write one thread's logic and manage shared memory, synchronization and coalescing yourself. In Triton you write Python for a whole block of data, and the compiler handles thread mapping, coalescing and shared memory. So Triton is much faster to write, and torch.compile generates it on GPU. I'd still drop to CUDA or CUTLASS for the newest hardware features or the last bit of performance.

对比:你管什么,compiler 管什么

CUDATriton
你管什么thread index 怎么对应数据;shared memory;同步;coalescing;thread block 怎么切block 大小;program 之间怎么分工。参数可以 autotune
compiler 管什么只照你写的编译program 内部的 thread 安排、coalescing、shared memory
可移植性NVIDIA GPUNVIDIA GPU 和 AMD GPU(ROCm);CPU 还在开发
适合最新的硬件特性(TMA、warp specialization),或最后一点性能。常借 CUTLASS 来写fusion、reduction、attention 变体,开发快。torch.compile 生成的也是它

OpenAI 发布 Triton 时的对照表也是这样分的,跨 SM 的 scheduling 两边都是 manual。

同一个 vector add,两种写法

// CUDA: 一个 thread 算一个元素
__global__ void add(const float* x, const float* y, float* out, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) out[i] = x[i] + y[i];
}

# Triton: 一个 program 算一个 block
@triton.jit
def add(x_ptr, y_ptr, out_ptr, n, BLOCK: tl.constexpr):
    offs = tl.program_id(0) * BLOCK + tl.arange(0, BLOCK)
    mask = offs < n                      # 最后一个 block 可能不满
    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)

CUDA 里的 i 是一个下标。Triton 里的 offs 是一个 block 的下标,tl.load 一次读进一整块。由几个 thread 读、怎么读才 coalesced,代码里没有写。

对 agent 来说,抽象层次高了更容易写对。μCUTLASS 论文在 59 道 KernelBench 题上用 GPT-5-mini 对比,直接写底层代码的 geomean 是 0.40x,改写 DSL 后是 1.27x。它的 DSL 是 μCUTLASS,不是 Triton。

必答点

追问

FlashVector 里两个 Triton 各出现在哪?

Case 1 改的是 Triton Inference Server 的 Python backend,Case 2 让 compiler 融出一个 Triton kernel(见 FlashVector)。

Triton 写的 kernel 能和厂商库一样快吗?

官方 tutorial 说它的 FP16 matmul "achieves performance on par with cuBLAS or rocBLAS"。同一页的 benchmark 表里各尺寸互有胜负,要在目标 shape 上自己测(见 kernel benchmark)。

面试让手写 Triton 或 CUDA,会写什么?

一场 AI infra 推理方向实习面试(2026-03-03,牛客)手撕过「triton/cuda:layer norm、大矩阵规约求和、online-softmax」。求和见 reduction,online softmax 见 FlashAttention。

Triton 里最难的是什么?

Triton 文档说这种写法的主要难点是 work scheduling,也就是一个 program 负责输出的哪一块(tiling 见 matmul tiling)。

来源OpenAI:Introducing Triton(2021-07,CUDA 和 Triton 的 manual / automatic 对照表)· Triton 文档 Introduction · Triton tutorial:vector add · Triton tutorial:matmul(@triton.autotune)· Triton README(支持的硬件)· μCUTLASS · FlashVector §5.1、§5.2 · ml-八股 08 第 11 题 · 面经出处见 面试题与 JD 1.2、3.2 节