全景：算子工程在 AI 计算栈的位置

写一行 torch.matmul(a, b)，背后可能调用 cuBLAS 的某个 GEMM kernel，也可能是 cuDNN 的卷积实现，或者一段融合过的自定义 CUDA。框架把这些细节藏起来，于是大多数工程师从没想过：这一行到底快不快、还能不能更快、什么时候该自己动手。

这一篇先把 AI 计算栈摊开，确定”算子（operator/kernel）“在哪一层、它和上层框架与下层硬件怎么衔接，再回答一个实用问题：什么场景下你必须越过框架去写或调 kernel。本系列后续所有篇章都围绕这一层展开，所有性能数字来自文末声明的同一块 GPU 真实运行。

一、从一行 matmul 到 SASS：五层栈

把”一次矩阵乘”自上而下拆开，大致是五层：

flowchart TD
  A["框架算子层<br/>torch.matmul / nn.Linear"] --> B["算子库层<br/>cuBLAS / cuDNN / CUTLASS"]
  B --> C["手写 kernel 层<br/>CUDA C++ / Triton"]
  C --> D["编译层<br/>NVRTC / nvcc → PTX"]
  D --> E["硬件指令层<br/>SASS / SM 执行"]

• 框架算子层：PyTorch、JAX、TensorFlow 暴露的 matmul、softmax、LayerNorm。它负责形状推断、自动微分、设备调度，但不负责”怎么算得快”。
• 算子库层：NVIDIA 提供的高度优化库。GEMM 走 cuBLAS，卷积/注意力等走 cuDNN，模板化 GEMM 走 CUTLASS。框架优先调它们。
• 手写 kernel 层：当库覆盖不到，或需要把多个操作融合成一个 kernel 时，工程师用 CUDA C++ 或 Triton 直接写。本系列的主战场。
• 编译层：CUDA C++ 由 nvcc（离线）或 NVRTC（运行时）编译为 PTX（一种虚拟 ISA），再由 ptxas 编译为目标架构的 SASS。
• 硬件指令层：SASS 是真正在 SM（Streaming Multiprocessor）上执行的机器码。

“算子工程”主要发生在中间三层：选库、写 kernel、理解编译与硬件如何影响最终性能。

PTX 与 SASS 的区别值得记住：PTX 是向前兼容的虚拟指令集，一份 PTX 可以在更新的架构上被驱动重新编译；SASS 是与具体架构（如 Ampere 的 sm_86）绑定的真实机器码。性能分析最终要看 SASS，而不是 PTX。这一点在 Nsight 调优篇 会再展开。

二、框架算子是怎么落到 kernel 的

以 PyTorch 为例，torch.matmul 经过分派（dispatch）选择后端实现，对 CUDA 张量最终调用 cuBLAS 的 GEMM 入口（如 cublasGemmEx）。这中间没有魔法：框架的职责是把高层语义翻译成”调哪个 kernel、传什么参数”，真正决定吞吐的是被调用的那个 kernel。

这带来两个推论：

1. 框架开销和 kernel 开销是两回事。 小张量上，Python 分派、张量元数据处理、kernel launch 的固定开销可能比计算本身还大；大张量上，时间几乎全在 kernel 里。判断瓶颈在哪，是优化的第一步。
2. 库已经很强，但有边界。 cuBLAS/cuDNN 针对常见形状和精度做了大量调优。但它们是”按算子”优化的——库不知道你的 matmul 后面紧跟一个 bias + GELU，于是中间结果要写回显存再读回来。这类跨算子的访存浪费，正是手写融合 kernel 的价值所在，详见 kernel fusion 篇。

三、什么时候才需要自己写 kernel

不是所有场景都该手写。优先用库，下面几类情况才值得越过框架：

判断标准始终是收益：先用 profiler 确认目标算子确实是瓶颈、确实有优化空间（compute-bound 还是 memory-bound），再动手。盲目手写一个还不如 cuBLAS 的 GEMM 是常见的时间浪费。“先定位再优化”是 Roofline 篇 和 Nsight 篇 的主题。

四、一个直觉：大多数”简单”算子是被访存卡住的

很多人以为 GPU 的瓶颈是算力，但相当多的算子卡在显存带宽。看一个最简单的例子——向量加法 c[i] = a[i] + b[i]：每个元素做 1 次加法，却要读 2 个 float、写 1 个 float，共 12 字节访存。算术强度（每字节访存对应的浮点运算）只有 \(1/12\) FLOP/byte，远低于 GPU 算力与带宽的平衡点。这种 kernel 再怎么优化计算也没用，唯一能逼近的上限是显存带宽。

在本系列的测试卡上实测，长度 \(2^{24}\) 的单精度向量加法达到约 413 GB/s 有效带宽，是该卡理论带宽 448 GB/s 的约 92%。也就是说这个 kernel 已经基本”贴着内存墙”跑，没有进一步优化空间。能不能一眼看出一个算子是这种情况，决定了你会不会把时间花在错误的地方。算术强度与 Roofline 的完整方法在 第 07 篇。

extern "C" __global__
void vadd(const float* a, const float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) c[i] = a[i] + b[i];
}

这段 kernel 的写法、blockIdx/threadIdx 索引计算、launch 配置，会在 第 04 篇 从零讲清楚。

五、本系列的实验环境与方法

为了让性能数字可核对，本系列绝大多数测量都在同一块 GPU 上完成，环境如下：

测量方法：用 CUDA event 记录 kernel 前后时间戳，预热 10 次后取 100 次运行的中位数。这套口径在涉及性能数字的篇章会反复使用，单独的方法细节见 Nsight 调优篇。

需要说明两个边界：

• 这块卡是消费级 Ampere（GA10x），第三代 Tensor Core，支持 FP16/BF16/TF32，但不支持 FP8（FP8 需要 Hopper sm_90 及以后）、也没有 Hopper 的 TMA（Tensor Memory Accelerator）。涉及这些特性的篇章（第 11、第 21）会明确标注哪些是本卡实测、哪些是引用更高代际硬件的公开数据。
• WSL2 下的绝对延迟可能略高于裸机 Linux，但本系列关心的是同一环境下不同实现之间的相对差异，结论不依赖绝对值。

六、接下来读什么

如果你完全没有 CUDA 基础，按 推荐阅读路径 的”CUDA 入门到能调优”走：先建立硬件心智（第 02、第 03 篇），再写第一个 kernel（第 04 篇），然后学会用 Roofline 和 Nsight 定位瓶颈，最后进入 GEMM、Tensor Core、FlashAttention 这些核心算子。

核心结论先记住三条：算子工程发生在框架与硬件之间的中间三层；优先用库、有明确收益再手写；动手前先用 Roofline 判断瓶颈是算力还是访存。