算子、Kernel、函数到底谁是谁?从 Excel 的 SUM 一路讲到 GPU 汇编

算子、Kernel、函数到底谁是谁？从 Excel 的 SUM 一路讲到 GPU 汇编

这几个词天天见，但一追问就说不清：算子是不是函数？AITER 是不是注意力机制？有了 HIP C++ 为什么还要 FlyDSL？这篇用一个你天天用的东西当起点，把六层关系一次串通。

📌 更多推理优化实践在 GitHub

📦 GitHub Repo
: https://github.com/david-xinyuwei/david-share.git
📂 本系列
: 推理工程里那些「看起来是一回事」的概念

⭐ 欢迎 Star！

Author: 魏新宇 (Microsoft AI and Apps Global Black Belt (GBB) Architect)

上一篇讲 Triton、FlyDSL、CK、HIP C++ 的分层，发出去之后收到最多的追问是这几个：

算子到底是什么？是不是就是编译好的函数？
AITER 我一直以为是注意力机制，怎么变成算子库了？
既然有 HIP C++，AMD 为什么还要再造一个 FlyDSL？

这三个问题其实是同一个问题：这堆词分别站在哪一层。

这篇不堆名词，从一个你天天用的东西开始。

先把几个词说清楚

后面会反复出现这几个缩写，先一次性讲完，正文就不再打断：

缩写	英文全称	一句话解释
Kernel	—	在 GPU 上真正跑的那段已编译代码，中文叫【内核】，和操作系统内核不是一回事
GPU	Graphics Processing Unit	显卡，擅长同时干上万件小事
CUDA	Compute Unified Device Architecture	NVIDIA 的 GPU 软件平台
ROCm	Radeon Open Compute platform	AMD 的 GPU 软件平台，地位对应 CUDA
HIP	Heterogeneous-compute Interface for Portability	ROCm 上写 kernel 的原生 C++ 接口
AITER	AI Tensor Engine for ROCm	AMD 的算子库，不是某一个算法
CK	Composable Kernel	AMD 的 C++ 模板 kernel 库
DSL	Domain-Specific Language	领域专用语言，只为一类活儿设计的语言
GEMM	General Matrix Multiply	通用矩阵乘，深度学习里最吃算力的那一步
MoE	Mixture of Experts	混合专家，每个 token 只走其中几个子网络
cuBLAS	CUDA Basic Linear Algebra Subprograms	NVIDIA 的线性代数算子库

若只记一条：算子是「要算什么」，kernel 是「怎么算」。

一、先从 Excel 的 SUM 开始

打开 Excel，敲一行：

=SUM(A1:A100)

回车，出数。这中间其实发生了六件事：

层	发生了什么
① 你写的	`=SUM(A1:A100)`
② 谁接住	Excel 的公式引擎，认出「这是 SUM」
③ 去哪找实现	Excel 内置的函数库
④ 挑哪一个	100 个数用简单循环；100 万个数可能走并行
⑤ 真正跑的	一段早就编译好的机器码
⑥ 跑在哪	CPU

有两点值得停一下：

第一，第 ⑤ 层那段代码，是微软工程师多年前写好编译进去的，跟你按回车这一刻没关系。

第二，同一个 SUM，底下可能不止一段代码。 数据量不同，走的路径可能不同。

（Excel 内部实现微软没公开，这里只用它的结构讲道理。）

二、换成 GPU，结构一模一样

图 1：Excel 的调用链和 GPU 的调用链逐层对应。

你在 Python 里敲：

C = torch.matmul(A, B)

对照着看：

层	Excel	GPU
① 你写的	`=SUM(A1:A100)`	`torch.matmul(A, B)`
② 谁接住	公式引擎	PyTorch （或 vLLM / SGLang）
③ 去哪找	内置函数库	算子库：cuBLAS / AITER
④ 挑哪个	按数据量	按矩阵大小、精度、显卡型号查表挑
⑤ 真正跑的	编译好的机器码	编译好的 GPU Kernel
⑥ 跑在哪	CPU	GPU （经 CUDA / ROCm 装载启动）

一栏一栏都对得上。

第 ④ 层那张「查表」不是凭空来的——是厂商提前在真机上跑过上千组参数，把「什么情况用哪个版本」记下来的对照表。

三、算子、Kernel、算子库，到底谁是谁

图 2：一个算子对应多个 Kernel，算子库负责收纳和挑选。

先给三个定义：

词	是什么	在 Excel 里对应
算子	要算什么，一个功能的名字和规格	`SUM`
Kernel	编译好的、真正执行的那段代码	实现 SUM 的那段机器码
算子库	装着一堆 Kernel + 一张选择表的柜子	Excel 这个软件本身

算子是不是编译好的函数

不完全是。 这个直觉抓住了一半。

你说的「编译好的函数」，业界叫 Kernel。而算子是它上面一层——这件事叫什么名字、输入输出是什么规格。

区别在哪？一个算子往往对应好几个 Kernel。

矩阵乘这一个算子，底下可能有：

H100 上的 FP16 版本
MI300X 上的 BF16 版本
小矩阵专用版本
大矩阵专用版本

算子只有一个，Kernel 有一堆，跑的时候按情况挑。

反过来也成立：一个 Kernel 可以同时实现好几个算子——这就是「融合」，后面会讲到。

一个简单的判断方法

能不能「调用它」。

你能写 =SUM(...) → SUM 是算子
你没法写 =Excel(...) → Excel 不是算子，是装算子的地方

同理：

你能写 torch.matmul(...) → matmul 是算子
你没法「调用 AITER」→ AITER 是算子库，你只能从里面挑一个算子来调

四、那 AITER 是不是注意力机制

不是。AITER 是算子库。

这个误会很常见，因为旁边确实有一堆名字里带 Attention 的东西。

打开 AITER 看一眼，它的算子测试文件是这样的：

文件	对应算子
`test_mha.py`	注意力
`test_mla.py`	另一种注意力
`test_moe.py`	混合专家
`test_gemm_a8w8.py`	矩阵乘
`test_rmsnorm2d.py`	归一化

官方自己的描述是「attention、MoE、GEMM、归一化、量化、通信等算子」。注意力只占其中一部分。

用 Excel 打比方：

你不会说「WPS 就是 SUM」。
同理，AITER 不是 Attention，它是装着 Attention 实现、也装着矩阵乘和 MoE 实现的库。

五、Attention 这个词，其实有六层意思

图 3：同一个词在不同语境下指向不同层次。

这才是真正容易乱的地方。

层	在这一层它是什么	具体指
① 机制	一种建模思想	让模型处理每个词时，能看向序列里其他相关的词
② 架构变体	机制的不同实现方式	多头（MHA）、多查询（MQA）、分组查询（GQA）、潜在注意力（MLA）、滑动窗口
③ 数学式	一个公式	softmax(Q·K转置 / √d)·V
④ 框架算子	一个可调用单元	`scaled_dot_product_attention()`
⑤ 计算算法	怎么算得快、省显存	FlashAttention、PagedAttention
⑥ Kernel	编译好的那段代码	AITER 里的 `fmha_kernels`

六层共用一个词。

论文里说 Attention → 说的是 ①
模型配置里说 Attention → 说的是 ②（用 MHA 还是 GQA）
写代码时说 Attention → 说的是 ④
调性能时说 Attention → 说的是 ⑤ 和 ⑥

所以「AITER 是不是 Attention」这个问题本身不成立——AITER 在第 ⑥ 层，Attention 机制在第 ① 层，差了五个层次。

就像问「汽车工厂是不是四轮驱动」。

顺带说清楚 Attention 不是 SUM

SUM 一步就完。Attention 是五步：

Q @ K转置  →  ÷√d  →  加 mask  →  softmax  →  @ V

放回 Excel，它不是一个 SUM，而是一张有中间列的表：

A 列	B 列	C 列	D 列	E 列
原始数据	A 两两相乘	B 除以常数	C 归一化	D 乘权重

E 才是答案，B/C/D 都是中间产物。

笨办法：B、C、D 老老实实填出来。放到 GPU 上，B 那个中间矩阵在长序列时可能有几十 GB，要写进显存再读出来。

FlashAttention 的办法：把 B、C、D 全删掉，从 A 直接算出 E——五步焊成一个 Kernel，中间结果只在芯片内部的高速缓存里流转。

省的不是计算量，是反复搬运数据的时间。这是大模型推理最贵的一块。

这也回答了前面那个问题：一个 Kernel 可以同时实现好几个算子。

六、为什么 GPU 需要专门的语言

到这里会冒出一个很自然的疑问：

微软写 Excel 的 SUM，用普通 C++ 就够了。GPU 上算个矩阵乘，为什么要 Triton、FlyDSL、CK、HIP C++ 这么多花样？

答案是：GPU 和 CPU 干活的方式根本不同，普通高级语言表达不了 GPU 的干活方式。

图 4：CPU 描述步骤，GPU 必须描述分工。

	CPU	GPU
有多少核	几个到几十个	上万个
每个核	很强，能干复杂活	很弱，只能干简单活
干活方式	一件一件按顺序做	上万个同时做同一件事
编程时你要说	步骤：先做这个再做那个	分工：谁负责哪块数据

打个比方：

CPU 像一个博士生
，给他一张任务清单，他从头做到尾。GPU 像一万个小学生，你不能给清单——得说「1 号加第 1 个数，2 号加第 2 个数……」，还得安排他们最后怎么汇总。

一行代码看出差别

图 5：GPU 代码必须先算出「我是几号工人」。

CPU 版求和：

float sum = 0; for (int i = 0; i < N; i++)     sum += a[i];

一个人，从头加到尾。

GPU 版求和：

__global__ void sum_kernel(float* a, float* out) {     int i = blockIdx.x * blockDim.x + threadIdx.x;     atomicAdd(out, a[i]); }

看第一行——代码里必须先算出「我是第几号工人」。

因为上万个工人跑的是同一段代码，每个必须知道自己该处理哪个数据。

普通 C++、Python、.NET 里，压根没有「我是几号工人」这个概念。 这就是必须有 GPU 专用写法的原因。

七、那 FlyDSL 是不是绕开了 ROCm

没有。它最后还是走 ROCm。

这是上一篇之后被问得最多的一个点，值得单独说清楚。

图 6：三种写法编译到同一种机器码，由同一个运行时装载。

HIP C++、Triton、FlyDSL 这三条路，最后都变成同一种 AMD GPU 机器码，都由 ROCm 装进 GPU 跑。

FlyDSL 不是 ROCm 的替代品，它是通往 ROCm 的另一个入口。

用你熟悉的东西打比方：

Word、Markdown、LaTeX 都能生成 PDF。
你不会问「有了 Word 为什么还要 Markdown」——因为不同场景，写起来效率差很多。

那为什么不都用 HIP C++

因为 HIP C++ 什么都能写，但写高性能 Kernel 时太痛苦。

这不是能力问题，是表达效率问题。

写一个高性能矩阵乘，用 HIP C++ 你得手写这些：

// 1. 每个线程负责哪块数据 —— 手算 int row = (blockIdx.y * 128) + (threadIdx.x / 16) * 8; int col = (blockIdx.x * 128) + (threadIdx.x % 16) * 8;  // 2. 从显存搬到共享内存 —— 手写地址 lds[threadIdx.x * 4 + 0] = A[row * K + k + 0]; // ... 还有几十行  // 3. 为了避开 bank conflict，地址还要错位 —— 手写位运算 int swizzled = (idx ^ ((idx >> 5) & 7)) * 4;  // 4. 双缓冲预取、同步、矩阵指令排布 —— 再几百行

一个高性能矩阵乘，HIP C++ 写下来常常上千行。这些地址算式极易出错，改一个分块参数就得全部重算。

同样的东西用 FlyDSL：

tile = logical_divide(A, make_layout((128, 32)))   # 说清楚怎么切 frag = tiled_copy.partition_S(tile)                # 说清楚谁拿哪块

你描述「怎么切」，编译器负责把那一堆地址算式生成出来。

三者真正的差别

图 7：三种写法在控制粒度上的分工。

	分块规则谁定	地址算式谁写
HIP C++	你定	你手写
FlyDSL	你定	编译器生成
Triton	编译器定	编译器生成

HIP C++ 和 FlyDSL 控制力相当，差别在于那堆地址算式谁来写。

Triton 走到了另一个极端——连「怎么切」都替你决定了。对大部分算子这挺好，省事；但对矩阵乘、Attention 这几个最吃性能的，工程师想自己定分块，Triton 就给不了。

这就像：汇编和 C 都能写程序，能力一样。但没人拿汇编写业务代码，因为效率差太多。

那 FlyDSL 是不是比 CK 快

不是。上限一样。

上一节说 FlyDSL 写得比 HIP C++ 简洁，很容易顺手推出「所以它更快」。这一步推错了。

因为最后跑在显卡上的东西是同一种：机器码。四条路殊途同归——

写法	语言	编译成
HIP C++	C++	GPU 机器码
CK	C++ 模板	GPU 机器码
Triton	Python	GPU 机器码
FlyDSL	Python	GPU 机器码

Python 那层在运行时早就不存在了。 它只是你写的时候用的语法，编译完就没了——显卡上没有 Python 解释器，一行 Python 都不跑。

语言不进机器码，怎么可能影响速度。

那 AMD 图什么？图的是试错速度。

调 Kernel 本质上是搜索：块切多大、数据怎么摆、循环展开几层，组合成千上万种，没人能一次写对，只能一个个试。

C++ 模板改一个参数，重新编译要等几分钟
Python 改一行，马上能跑

同样一下午，一个试了 5 种，一个试了 50 种。后者当然更容易撞上那个最优解。

不是语言快，是迭代快。

所以看到一份性能对比说 FlyDSL 赢了 CK，该先问一句：

对面那份 CK 实现，调优了吗？

常见的是这两种情况：

看到的结论	可能的实际原因
FlyDSL 比 CK 快	FlyDSL 那份是为这个矩阵形状精调过的，CK 那份是通用版本
FlyDSL 比 CK 快	新硬件指令 FlyDSL 先接上了，CK 还没跟上

快的是那份实现，不是那个语言。 这两件事必须分开——不然换个矩阵形状，结论就翻过来了。

（至于「FlyDSL 比 ROCm 快」，那是把路和车上的方向盘放一起比了。ROCm 是底座，上一节开头已经说过。）

八、最后：两条链，别搞混

图 8：运行时调用链与建库链是两回事。

这是整篇最值得记住的一张图。

运行时（你每次跑模型）：

① Python 调用 ② 框架 PyTorch / vLLM ③ 算子库 cuBLAS / AITER ④ 查表选 Kernel ⑤ 编译好的 Kernel ⑥ ROCm / CUDA → GPU

建库时（几个月前，在厂商实验室）：

Kernel 工程师     ↓ 用 Triton / FlyDSL / CK / HIP C++ / 汇编 编写     ↓ 真机跑上千组参数调优     ↓ 打包进算子库

Triton 和 FlyDSL 只出现在第二条链上。你跑模型的时候，根本碰不到它们。

唯一的例外是 torch.compile——它会当场生成 Triton 代码再编译。那是「现炒」，不是「热预制菜」。

每个名词归位

名词	在哪一层	一句话
注意力机制	建模思想	不在调用链上，是架构层面的概念
Attention 算子	①②	框架里可调用的那个单元
FlashAttention	⑤ 的算法	把五步焊成一步的办法
算子	②	「要算什么」的规格
Kernel	⑤	编译好的、真正执行的那段代码
算子库	③	装着一堆 Kernel + 一张选择表的柜子
cuBLAS / AITER	③	两个具体的柜子
Triton / FlyDSL / CK	建库链	造 Kernel 的工具，运行时不出现
CUDA / ROCm	⑥	地基，把 Kernel 装进 GPU 跑起来
汇编	⑥ 附近	最底层的写法，Kernel 最终都变成机器指令

小结

你写的 torch.matmul 只是个名字
，经过框架、算子库、查表，最后落到一段早就编译好的 GPU 代码上——和 Excel 的 SUM 结构完全一样。
算子是「要算什么」，Kernel 是「编译好的那段代码」，算子库是「装着一堆 Kernel 的柜子」
。一个算子可以有多个 Kernel，一个 Kernel 也能实现多个算子。
AITER 是算子库，不是注意力机制
。Attention 这个词本身有六层意思，说之前先确认在说哪一层。
GPU 需要专用语言，是因为它要描述「上万人怎么分工」
，普通高级语言没有这个概念。
FlyDSL 没有绕开 ROCm
，它和 HIP C++、Triton 一样，最后都编译成同一种机器码。
三者差别在于：分块规则谁定、地址算式谁写
。能力都够，差的是写多少行、多容易出错。
语言不决定性能上限，只决定你多快摸到那个上限
。看到「A 比 B 快」，先问 B 那份调优了没。
运行时链和建库链是两条
。Triton、FlyDSL 只在建库那条上。

自测四个问题

为什么说「算子不等于编译好的函数」？中间差了什么？
一个 Kernel 能不能同时实现好几个算子？举个例子。
为什么 Excel 的 SUM 用普通 C++ 就够，GPU 的矩阵乘却不行？
FlyDSL 和 HIP C++ 控制力相当，那 FlyDSL 的价值到底在哪？
有人拿出一份数据说「FlyDSL 比 CK 快 30%」，你应该先问什么？

参考来源

全部为公开资料：

AITER 官方仓库：https://github.com/ROCm/aiter （MIT）
FlyDSL 官方仓库：https://github.com/ROCm/FlyDSL （Apache-2.0）
Composable Kernel 官方仓库：https://github.com/ROCm/composable_kernel （MIT）
Triton 官方仓库：https://github.com/triton-lang/triton
PyTorch 官方仓库：https://github.com/pytorch/pytorch
RoFormer / RoPE 论文：https://arxiv.org/abs/2104.09864

上一篇：《Triton、FlyDSL、CK、HIP C++ 到底差在哪？10 张图讲清显卡 Kernel 技术栈》

下一篇预告：键值缓存（KV cache）从 FP8 换回 BF16，为什么最大并发会直接往下掉一截？显存这本账到底怎么算。