第三章 GPU的硬件组成：流式多处理器（SM）、计算单元（CUDA Core/Tensor Core）、内存层次结构
好，咱们今天聊点硬核的。GPU到底长什么样？
很多软件工程师写CUDA代码，写了好几年，却从没想过硬件是怎么执行这些指令的。我个人觉得，这就像开车不懂发动机原理——能开，但永远开不到极致。
这一章，我带你把GPU的“内脏”翻出来看看。
3.1 流式多处理器（SM）——GPU的“核心团队”
先问一个问题：GPU里有几百甚至几千个核心，它们是怎么组织起来的？
答案是——SM（Streaming Multiprocessor）。你可以把SM理解成一个“小分队”。每个SM内部有自己的计算单元、缓存、调度器。GPU就是由多个SM组成的。
举个例子：NVIDIA的A100有108个SM，每个SM内部又有64个FP32 CUDA Core。所以总核心数就是108×64=6912个。
SM的核心职责：
- 线程调度：SM负责把线程块（Thread Block）分配给内部的CUDA Core。我习惯把SM比作一个项目经理，它决定哪个线程块什么时候跑、跑在哪个核心上。
- 指令发射：每个SM内部有多个Warp Scheduler（经线调度器），每个调度器每周期可以发射一条指令给一组线程（32个线程为一组，称为一个Warp）。
- 资源管理：寄存器、共享内存这些资源，都由SM统一分配。
重要概念：Occupancy（占用率）
SM能同时容纳的Warp数量是有限的。如果每个线程块用太多寄存器，SM能同时跑的线程块就少，占用率就低。我曾经在一个项目里，为了追求单线程性能，给每个线程分配了128个寄存器，结果占用率掉到25%，整体性能反而下降了40%。
嗯，这就是典型的“捡了芝麻丢了西瓜”。
3.2 计算单元：CUDA Core 与 Tensor Core
SM内部有两种“打工人”：CUDA Core和Tensor Core。它们分工不同，各司其职。
3.2.1 CUDA Core——通用计算的主力
CUDA Core是GPU里最基础的计算单元。每个CUDA Core可以执行整数、浮点数的加减乘除等基本运算。
你写的每个CUDA线程，最终都会映射到一个CUDA Core上执行。但注意：一个CUDA Core同一时间只能执行一个线程的一条指令。所以，线程越多，越能“喂饱”这些核心。
CUDA Core的规格（以A100为例）：
| 类型 | 数量/SM | 总数量 | 主要用途 |
|---|---|---|---|
| FP32 Core | 64 | 6912 | 单精度浮点运算 |
| INT32 Core | 64 | 6912 | 整数运算、地址计算 |
| FP64 Core | 4 | 432 | 双精度浮点运算（科学计算） |
你可能会问：为什么FP64核心这么少？说白了，游戏和AI训练不需要那么高的精度。双精度是给科学计算准备的，所以数量少，但每个FP64 Core性能很强。
3.2.2 Tensor Core——AI时代的“特种兵”
Tensor Core是NVIDIA从Volta架构开始引入的专用计算单元。它专门为矩阵乘法而生。
为什么需要它？
深度学习里，全连接层、卷积层，本质上都是矩阵乘法。如果用CUDA Core做，一次只能算一个乘加（FMA）。但Tensor Core一次可以算一个4×4矩阵的乘法，也就是64个乘加操作。
Tensor Core的工作方式：
- 输入：两个4×4的矩阵（A和B）
- 输出：一个4×4的矩阵（C）
- 一次操作完成：D = A × B + C
我记得第一次用Tensor Core写矩阵乘法时，性能直接翻了4倍。但要注意：Tensor Core对数据精度有要求，通常需要FP16或BF16输入，FP32累加。
我的建议：
如果你的计算是矩阵乘法密集型（比如AI推理、训练），尽量用cuBLAS或cuDNN库，它们会自动调用Tensor Core。手写CUDA代码调用Tensor Core比较麻烦，需要用到WMMA（Warp Matrix Multiply-Accumulate）API。
3.3 内存层次结构——速度与容量的博弈
GPU的内存系统，说白了就是一个“金字塔”。越往上越快，但容量越小；越往下越慢，但容量越大。
你想想看，如果所有数据都放在最慢的全局内存里，那计算单元大部分时间都在等数据。所以，理解内存层次结构，是优化性能的关键。
3.3.1 寄存器（Register）——最快，但最稀缺
寄存器是每个线程私有的，速度最快（一个周期内读写）。每个SM有固定数量的寄存器，比如A100每个SM有65536个32位寄存器。
关键点：
- 每个线程能用的寄存器数量是有限的。如果每个线程用太多寄存器，SM能同时跑的线程数就减少。
- 编译器会自动分配寄存器，但你也可以用
__launch_bounds__来限制每个线程的最大寄存器数。
我曾经踩过的坑：
有一次我写了一个复杂的kernel，里面用了很多局部变量。编译后发现每个线程用了96个寄存器，导致SM只能同时跑8个Warp（正常可以跑64个）。性能惨不忍睹。
后来我把一些中间结果放到共享内存里，寄存器用量降到32个，占用率从12.5%提升到50%，性能翻了3倍。
3.3.2 共享内存（Shared Memory）——线程块内的“高速缓存”
共享内存是每个SM内部的，同一个线程块内的所有线程都可以访问。速度比全局内存快很多（大约20-30倍延迟），但容量有限（A100每个SM最多164KB）。
共享内存的典型用途：
- 数据复用：多个线程需要读取同一份数据时，先加载到共享内存，避免重复访问全局内存。
- 线程间通信：同一个线程块内的线程可以通过共享内存交换数据，比用全局内存快得多。
- 合并访问：把非合并的全局内存访问，先整理到共享内存，再合并读取。
举个例子：矩阵乘法中，每个线程需要读取A的一行和B的一列。如果直接读全局内存，每个线程都要访问很多次。但如果我们把A的一行和B的一列先加载到共享内存，所有线程就可以共享这些数据，大大减少全局内存访问。
3.3.3 全局内存（Global Memory）——最大，但最慢
全局内存就是GPU的显存。容量大（几十GB），但延迟高（几百个周期）。所有线程都可以访问，但速度最慢。
优化全局内存访问的关键：
- 合并访问（Coalesced Access）：相邻的线程访问相邻的内存地址。这是最重要的优化手段。
- 对齐访问：内存地址最好是128字节对齐。
- 减少访问次数：能用共享内存或寄存器，就别用全局内存。
一个简单的性能对比：
| 内存类型 | 延迟（周期） | 带宽（GB/s） | 容量 |
|---|---|---|---|
| 寄存器 | 1 | 极高 | 每个SM 256KB |
| 共享内存 | ~20 | ~10000 | 每个SM 164KB |
| L1缓存 | ~30 | ~5000 | 每个SM 192KB |
| L2缓存 | ~200 | ~2000 | 40MB |
| 全局内存 | ~400 | ~2000 | 40GB+ |
3.4 总结：如何利用硬件特性写出高性能代码？
好了，硬件结构讲完了。那对我们写代码有什么实际指导？
- 控制寄存器用量：每个线程不要用太多局部变量。如果发现占用率低，先检查寄存器使用量。
- 善用共享内存：数据复用场景，优先考虑共享内存。但注意不要超过SM的容量限制。
- 全局内存访问要合并：这是最基础也最重要的优化。写代码时，确保线程索引和内存地址是线性对应的。
- 考虑使用Tensor Core：如果你的计算是矩阵乘法，别自己手写，用cuBLAS。
我个人习惯是：先写一个能跑的版本，然后用NVIDIA的Nsight Compute工具分析性能瓶颈。看看是计算受限还是内存受限，再针对性地优化。
记住一句话：GPU的硬件设计，就是为了让你用并行度来掩盖延迟。所以，尽量让SM“忙起来”，别让任何一个核心闲着。
下一章，我们会深入讨论线程模型和Warp调度。到时候你会发现，硬件和软件的结合，远比想象中有趣。