跳到主要内容

GPU 之上 — CUDA kernel 里训练长什么样

这一章讲三件事: 你写的每一行 PyTorch,在显卡上变成了什么(线程、块、网格); 为什么显卡上的瓶颈往往不是「算」而是「搬数据」;以及四个经典的并行翻车坑。 这是全书的最后一章拆解——到这里,从文字到 checkpoint 到显卡内部,整条链都打开了。

1. 先看现象:同一行代码,两个世界的速度

CPU 像几位教授:每个都极聪明,一次解一道难题。 GPU 像一间有一万六千名小学生的礼堂:每个人只会算一道乘法,但一万六千人同时算1。 神经网络恰好全是「几百万次互不依赖的乘加」,正是小学礼堂的主场—— 原书给的对照:CPU 一般 8-64 核,H100 超过 16,000 个 CUDA 核心2

2. 主走查:一个 kernel 的诞生

kernel(kernel,在显卡上被成千上万线程同时执行的一个函数)是 GPU 编程的最小单位。 拿原书的第一个例子走一遍——把 1024 个数各翻一倍(线程编号的计算方式照录代码)3:

CPU 视角: for i in range(1024): data[i] *= 2 ← 串行,1024 次

GPU 视角: 写一个 kernel:算出「我是几号」,只翻自己那一个
编号 = 块号 × 每块线程数 + 块内编号
data[我的编号] *= 2
启动时说一句「派 8 个块、每块 128 线程」
→ 1024 个线程同时执行,每个翻自己的格子,一步完成

图说:这就是原书说的 SIMT——单指令、多线程:
所有线程执行同一份代码,只是各自处理数据的不同一角[^4]。

原书给的规模感:1024×1024 的矩阵乘,可以派 1,048,576 个线程,每个线程算输出矩阵的一个元素—— CPU 上毫秒级的活儿,显卡上微秒级完成4

线程的组织是三层,Reporting 关系固定5:

是什么原书例子
线程最小执行单位,算一个格上例的每个「翻倍工」
一组线程,块内可通过共享内存协作每块 256-1024 个线程
网格一次 kernel 启动的全部块32×32 个块

3. 真正的瓶颈:搬数据,不是算

原书这一节推翻了新手的默认想象:现代 GPU 常常是「内存受限」的——算得快,喂不过来6。 显卡的显存层级像食堂:寄存器(自带饭盒,最快)、共享内存(块内公用的取餐台)、全局显存(大仓库,最慢)。 一个朴素的矩阵乘 kernel,每个元素都要从大仓库取一遍——大部分时间在排队等饭,不是在吃

两个解法,原书都给了代码7:

  • 共享内存+分块(tiling):把矩阵切成小方块,整块搬进取餐台,块内线程各取所需—— 每个元素从大仓库只读一次,而不是每个线程都读一遍;
  • 算子融合(fusion):把「乘法→加偏置→激活」几个 kernel 合成一个,中间结果不落仓库—— 原书的结论:省下的搬运常常比省下的计算更值钱8

第 07 章提过一句的 FlashAttention 就是同一思想的代表作:注意力矩阵根本不整体落盘, 分块在高速层内算完就丢,显存从 O(n²) 降到 O(n)9

4. 训练循环的 kernel 视角

第 10 章的训练循环,在显卡上是一串 kernel 排队10:

前向:矩阵乘 kernel(算层输出)→ 激活 kernel → 归一化 kernel → loss kernel
反向:loss 梯度 kernel → 逐层梯度 kernel → 权重梯度 kernel → 梯度汇总 kernel
更新:参数更新 kernel →(可选)梯度裁剪 kernel

图说:第 02 章「反向传播=链式法则」的那一趟,在这里就是
十几类 kernel 的固定排班;每一步都是全卡并行。

PyTorch 的角色因此清楚了:它是个排班员——你写 x @ W,它挑 cuBLAS(高度调优的矩阵乘库)的 kernel; 你写注意力,它挑 cuDNN 或 FlashAttention 的 fused kernel11。 「库先于自写」是原书的明确建议:那些库是多年的调优结晶;自写 kernel 只在三种情况值得—— 全新算子、显存关键路径、可融合的机会12

5. 四个经典翻车坑

原书列的坑,全可以用「一万六千人同时干活」的类比记住13:

人话版解法
竞态两个人同时改同一格,结果看谁手快同步点(块内屏障)或原子操作
bank 冲突取餐台分 32 个窗口,全挤一个窗口=退化成排队让相邻线程取相邻地址
分支发散同组线程走了 if 的两边,两边都得慢慢执行完少分支;同组走同一路
占用不足每个线程占的饭盒(寄存器)太多,礼堂里挤不下足够多的人限制寄存器/共享内存用量

6. 作者的判断与证据

说法书里的证据我们的标注
神经网络天然适合 GPU矩阵乘/逐元素运算天然并行的机制描述14领域共识
内存受限、融合最值钱「Modern GPUs are memory-bound」的直接论断+tiling/融合的代码演示6与 HPC 领域共识一致
库优先于自写cuBLAS/cuDNN「代表多年的优化」+自写 kernel 的三个适用条件12工程共识
四个坑逐一给了成因与解法13CUDA 编程的标准清单

判断(我们的,不是书里的): 这一章表面上讲 CUDA,实际讲的是**「性能问题要换一个抽象层才能看清」** ——第 04 章的 ε、rsqrt,第 06 章的滑窗,第 10 章的混合精度,单看都是小聪明; 换到「数据搬运是主要成本」这一层看,它们全是同一件事:少搬一次数。 这也是为什么同一份模型代码,换硬件、换实现,速度能差十倍。 如果错,会错在: 如果未来硬件把算与存的边界抹掉(行话叫存算一体:把存储和计算合成一体的硬件思路), 「搬运是主要成本」的框架就会过时——但「换抽象层看性能」这个方法论不会。

7. 边界与局限

  • 原书没讲 AMD 的 ROCm 或 TPU 这类非 CUDA 生态——标题里的「Modern Frameworks」实际只覆盖 NVIDIA 系;
  • kernel 代码示例是教学级(朴素矩阵乘),原书自己注明「真实实现会用 tiling、共享内存与寄存器分块」15;
  • 「如何给一个模型做性能剖析(找出哪个 kernel 慢)」原书没有覆盖——这是把本章知识用起来的第一步;
  • Triton 这类「用 Python 写 GPU kernel」的新工具原书只在优化库清单里带过一笔16,没展开。

8. 可带走的

  1. GPU = 一万六千个只会乘法的小学生;神经网络全是乘法,所以是主场;
  2. kernel = 被上万线程同时执行的函数;线程→块→网格三层编制;
  3. 1024×1024 矩阵乘 = 派一百万个线程,每人算一个格;
  4. 现代 GPU 内存受限:瓶颈在搬数据,不在算;
  5. 共享内存+分块让每个数从大仓库只读一次;算子融合让中间结果不落仓库;
  6. FlashAttention = 「注意力不落盘」的融合,显存从 O(n²) 到 O(n);
  7. 库优先(cuBLAS/cuDNN);自写 kernel 只为全新算子、显存关键路径、可融合三件事;
  8. 四个坑=一万六千人同时干活的四种乱象:抢格子、挤窗口、走岔路、占座太多;
  9. 性能问题要下潜一层抽象才能看清——少搬一次数,是所有优化共同的名字。

9. 原文地图

主题原书章原文位置
CUDA 是什么Training Neural Networks with CUDA Kernels and Modern Frameworkstext/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:9(搜「Compute Unified Device Architecture」)
CPU 对照 GPU 核数Training Neural Networks with CUDA Kernels and Modern Frameworkstext/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:13(搜「16,000」)
适合 GPU 的运算Training Neural Networks with CUDA Kernels and Modern Frameworkstext/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:15(搜「Matrix multiplications」)
1000×1000 的并行Training Neural Networks with CUDA Kernels and Modern Frameworkstext/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:25(搜「1000x1000」)
kernel 定义与 globalCUDA Kernels Explainedtext/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:31(搜「CUDA kernel is a function」) · text/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:47(搜「global」)
线程-块-网格Thread Hierarchytext/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:35(搜「Thread Hierarchy」)
百万线程例Thread Hierarchytext/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:43(搜「1,048,576」)
内存层级Memory Hierarchytext/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:69(搜「Memory Hierarchy」)
训练循环的 kernel 排班CUDA Kernels in Neural Network Trainingtext/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:87(搜「forward pass」)
矩阵乘 kernel 代码Practical CUDA Kernel Examplestext/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:15(搜「matrix multiplication」)
朴素实现的问题Practical CUDA Kernel Examplestext/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:249(搜「shared memory」)
tiling 演示Practical CUDA Kernel Examplestext/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:319(搜「loaded from global memory only once」)
库优先建议Integration with Deep Learning Frameworkstext/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:571(搜「PyTorch Custom CUDA Extensions」)
四个坑Common Pitfallstext/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:1009(搜「Race Conditions」) · text/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:1027(搜「Bank Conflicts」) · text/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:1029(搜「Divergent Branches」) · text/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:779(搜「Occupancy」)
关键要点清单Common Pitfalls 后的 Key Takeawaystext/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:1035(搜「Key Takeaways」) · text/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:1039(搜「Fusion Is Powerful」)
术语表(章末附录)Appendix: Glossary of Termstext/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:1049(搜「Glossary」)

Footnotes

  1. 出处:「Why GPUs for Deep Learning?」(text/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:13,搜「16,000」)。原文:CPU 8-64 核优化串行,H100 超过 16,000 个 CUDA 核心;GPU 同时计算数千个矩阵元素、每个在不同核心上。

  2. 出处:「CUDA Kernels Explained」第 31 段(text/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:31,搜「CUDA kernel is a function」)。原文:kernel 是在 GPU 上由大量线程并行执行的函数;启动时指定线程数,每个线程处理数据的一小部分。

  3. 出处:「Kernel Syntax」代码(text/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:47,搜「global」)。原文代码:int idx = blockIdx.x * blockDim.x + threadIdx.x; if (idx < n) { data[idx] = data[idx] * 2.0f; }——每个线程翻倍自己那一个元素。

  4. 出处:「Thread Hierarchy」末段(text/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:43,搜「1,048,576」)。原文:1024×1024 矩阵可用 32×32 块、每块 32×32 线程,共 1,048,576 个线程——每个矩阵元素一个;另见 :25(搜「1000x1000」):CPU 毫秒级、GPU 微秒级。

  5. 出处:text/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:35(搜「Thread Hierarchy」)。原文:线程是最小执行单位;线程组成块,块内通过共享内存协作,块可含 256、512 或 1024 线程;块组成网格。

  6. 出处:「Key Takeaways」(text/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:1037,搜「Memory Hierarchy Matters」)。原文:「Modern GPUs are memory-bound. Efficient kernel design focuses on minimizing data movement and maximizing use of fast memory」。 2

  7. 出处:「Practical CUDA Kernel Examples」(tiling 段)(text/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:319,搜「loaded from global memory only once」;:321,搜「Threads cooperate」)。原文:分块后每个元素从全局内存只读一次;线程协作把数据搬进共享内存、在共享数据上计算。

  8. 出处:「Key Takeaways」(text/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:1039,搜「Fusion Is Powerful」)。原文:合并算子减少 kernel 启动开销与显存往返,「often providing bigger gains than raw compute optimization」。

  9. 出处:「Performance Optimizations」(第 07 章文本)(text/42-ch07-7-attentionblock-with-rotary-embedding-gqa-slidi.txt:815,搜「FlashAttention」)。原文:FlashAttention 分块计算注意力、不物化完整注意力矩阵,减少显存 I/O、2-3 倍提速。

  10. 出处:「The Training Loop at Kernel Level」(text/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:87,搜「forward pass」)。原文:前向由矩阵乘 kernel(层激活)、逐元素 kernel(激活函数)、归一化 kernel(均值方差)、损失 kernel 组成;反向从损失梯度 kernel 起、经逐层梯度、权重梯度、梯度汇总;更新阶段是参数更新 kernel(必要时加梯度裁剪 kernel)。

  11. 出处:「Integration with Deep Learning Frameworks」(text/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:571,搜「PyTorch Custom CUDA Extensions」)。原文:PyTorch 支持自定义 CUDA 扩展,常用算子走 cuBLAS/cuDNN 的优化实现。

  12. 出处:「Key Takeaways」(text/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:1041,搜「Libraries Are Highly Optimized」;:1063,搜「Custom Kernels Have Their Place」)。原文:cuBLAS/cuDNN 代表多年优化,能用就用;自写 kernel 适合全新算子、内存关键路径与融合机会。 2

  13. 出处:「Common Pitfalls」(text/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:1009,搜「Race Conditions」;:1043,搜「Bank Conflicts」;:1045,搜「Divergent Branches」;:1047,搜「Occupancy」)。原文:竞态需 __syncthreads() 或原子操作;共享内存 32 bank 全挤同窗会串行;warp 内分支发散会串行执行;寄存器/共享内存占用过高会降低活跃线程数。 2

  14. 出处:「Why GPUs for Deep Learning?」(text/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:15,搜「Matrix multiplications」)。原文:神经网络的重活——大规模矩阵乘、逐元素张量运算、可并行的梯度计算、批处理——天然适配并行。

  15. 出处:「Practical CUDA Kernel Examples」(朴素实现注)(text/95-fm-training-neural-networks-with-cuda-kernels-and-m.txt:253,搜「Real implementations」)。原文:「Real implementations use tiling, shared memory, and register blocking」。

  16. 出处:「Flexibility for Custom Kernels and Optimizations」(第 01 章文本)(text/15-fm-why-cuda-for-building-large-language-models.txt:79,搜「CUTLASS and Triton」)。原文:CUTLASS 与 Triton 简化自定义矩阵乘与注意力 kernel 的编写。