GPU硬件架构：从SM到内存层级，一次讲透
大家好，我是你们的老朋友。今天咱们聊聊GPU的硬件架构。说实话，搞CUDA编程这么多年，我最大的感触就是：不懂硬件，写出来的代码就像蒙着眼睛开车。你想想看，连车子的引擎怎么转都不知道，怎么能开得快？
这一章，我会带你从流式多处理器（SM）开始，一路走到内存层级。嗯，内容有点多，但都是干货。
1. 流式多处理器（SM）——GPU的心脏
SM，全称Streaming Multiprocessor，是GPU最基本的计算单元。你可以把它想象成一个小型处理器集群。每个SM内部包含多个CUDA核心、共享内存、寄存器文件、warp调度器等组件。
我个人习惯把SM比作一个工厂车间：
- CUDA核心 = 车间里的工人
- 共享内存 = 车间里的工作台
- 寄存器 = 工人手里的工具
- Warp调度器 = 车间主任
我在项目中遇到过一个问题：某个kernel在GTX 1080上跑得飞快，换到V100上反而慢了。后来一查，原来是SM数量不同导致的线程块分配策略问题。你看，不了解SM的架构，连性能问题都定位不了。
关键点：每个SM独立执行线程块（Thread Block），多个SM并行工作。GPU的性能，很大程度上取决于SM的数量和每个SM的资源容量。
2. CUDA核心——没那么简单
很多人以为CUDA核心就是CPU核心的缩小版。其实不然。CUDA核心是轻量级的计算单元，专门为并行计算设计。
每个CUDA核心包含：
- 整数运算单元（INT ALU）
- 浮点运算单元（FP ALU）
- 一些控制逻辑
注意，CUDA核心没有复杂的分支预测、乱序执行这些CPU才有的东西。为什么？因为GPU的设计哲学是「用更多的核心，做更简单的事」。
我曾经犯过一个错误：在kernel里写了大量if-else分支，结果性能惨不忍睹。后来才明白，warp内的分支分歧会让CUDA核心串行执行，白白浪费了并行能力。
避坑指南：尽量减少warp内的分支。如果实在避免不了，让同一warp内的线程走相同路径。
3. Warp调度——GPU的隐形指挥官
Warp是GPU执行的基本单位，通常包含32个线程。为什么是32？这是NVIDIA多年经验总结出来的最优值。
Warp调度器的工作方式很有意思：
- 它维护一个就绪warp池
- 每个时钟周期，选择其中一个就绪warp发射指令
- 如果某个warp在等待内存访问，就切换到另一个warp
说白了，warp调度就是GPU隐藏延迟的绝招。当某个warp在等数据时，调度器立刻切换到另一个warp，让计算单元始终忙碌。
我记得有一次优化一个矩阵乘法kernel，发现occupancy（占用率）很高，但性能就是上不去。后来用nvprof一看，warp stall（停顿）占了大量时间。原来是共享内存bank conflict导致的。调整数据布局后，性能直接翻倍。
注意：warp调度是硬件自动完成的，但你可以通过优化代码来「帮助」调度器。比如：减少寄存器使用量、避免共享内存bank conflict、合理使用__syncthreads()。
4. 内存层级——从快到慢，从贵到便宜
GPU的内存层级，说白了就是速度与容量的权衡。越靠近计算单元的内存越快，但容量越小；越远的内存越慢，但容量越大。
| 内存类型 | 位置 | 速度 | 容量 | 作用域 |
|---|---|---|---|---|
| 寄存器 | SM内部 | ~1周期 | 每个线程最多255个 | 线程私有 |
| 共享内存 | SM内部 | ~5周期 | 每个SM 48-164KB | 线程块内共享 |
| 本地内存 | 显存 | ~100周期 | 取决于显存 | 线程私有 |
| 全局内存 | 显存 | ~400周期 | GB级别 | 所有线程 |
| 常量内存 | 显存（带缓存） | ~100周期（缓存命中时更快） | 64KB | 所有线程（只读） |
| 纹理内存 | 显存（带缓存） | ~100周期（缓存命中时更快） | 取决于显存 | 所有线程（只读） |
4.1 寄存器——最快的，也是最稀缺的
寄存器是每个线程私有的，速度最快。但每个SM的寄存器总数有限（比如V100有65536个）。如果你每个线程用太多寄存器，SM能同时容纳的线程数就会减少，occupancy就低了。
我有个习惯：写kernel时先用--maxrregcount限制寄存器数量，看看性能变化。有时候少用几个寄存器，occupancy上去了，性能反而更好。
4.2 共享内存——程序员手中的利器
共享内存是可编程的缓存，由同一个线程块内的所有线程共享。它的速度比全局内存快两个数量级。
共享内存最常见的用途是数据复用。比如矩阵乘法中，把子矩阵加载到共享内存，然后反复使用，避免重复访问全局内存。
// 共享内存示例：矩阵分块
__shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
__shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
// 每个线程加载一个元素到共享内存
As[ty][tx] = A[row * W + tx];
Bs[ty][tx] = B[ty * W + col];
__syncthreads();
// 从共享内存读取，进行计算
for (int k = 0; k < BLOCK_SIZE; ++k) {
sum += As[ty][k] * Bs[k][tx];
}小技巧：使用共享内存时，注意bank conflict。如果多个线程访问同一个bank的不同地址，就会发生冲突，导致串行化。解决办法是调整数据布局，比如加padding。
4.3 全局内存——慢，但容量大
全局内存是GPU的主存，所有线程都能访问。但它的延迟高达400-800个时钟周期。所以，尽量减少全局内存访问次数。
访问全局内存时，要遵循合并访问（coalesced access）原则：同一warp内的线程，访问连续的内存地址。这样硬件可以把多次访问合并成一次，效率高得多。
我曾经接手过一个项目，代码里到处都是随机访问全局内存。优化后改成连续访问，性能提升了5倍。嗯，有时候性能问题就是这么简单粗暴。
4.4 本地内存——寄存器的「溢出区」
当寄存器不够用时，编译器会把一些变量「溢出」到本地内存。本地内存虽然物理上在显存里，但每个线程只能访问自己的那份。
注意：本地内存的访问速度很慢，因为它本质上就是全局内存。如果发现kernel性能不佳，可以用--ptxas-options=-v查看本地内存使用量。如果太多，说明寄存器压力过大，需要优化。
4.5 常量内存——只读，但带缓存
常量内存只有64KB，但它的缓存机制很特别：同一warp内的线程访问同一地址时，只需要一次内存访问。所以，常量内存特别适合存储查找表、系数等只读数据。
4.6 纹理内存——为空间局部性而生
纹理内存最初是为图形学设计的，但用在通用计算中也有奇效。它支持硬件插值和边界处理，而且对二维/三维访问模式有优化。
我个人很少用纹理内存，除非是图像处理或者需要插值的场景。大多数情况下，共享内存+全局内存的组合就够用了。
5. PCIe与NVLink——GPU之间的高速公路
多GPU系统里，数据传输是绕不开的话题。PCIe和NVLink就是GPU之间、GPU与CPU之间的通信通道。
| 特性 | PCIe 4.0 x16 | NVLink 2.0 | NVLink 3.0 |
|---|---|---|---|
| 单向带宽 | ~16 GB/s | ~50 GB/s | ~100 GB/s |
| 延迟 | ~1μs | ~200ns | ~200ns |
| 拓扑 | 树形 | 混合立方体 | 全互联 |
NVLink的优势很明显：带宽高、延迟低。如果你在多GPU系统上做分布式训练或大规模计算，NVLink能显著减少通信开销。
我记得有一次在DGX-1上做实验，8块V100通过NVLink全互联。数据在GPU之间传输，速度比PCIe快了近5倍。嗯，贵有贵的道理。
注意：使用多GPU时，尽量让数据在GPU之间直接传输（通过NVLink），避免经过CPU中转。使用cudaMemcpyPeer或统一内存（Unified Memory）可以简化编程。
6. 知识体系总览
下面这张图，是我自己画的GPU硬件架构知识体系。你可以把它当作本章的「地图」：
这张图展示了GPU硬件架构的核心组件和它们之间的关系。从GPU整体，到SM，再到CUDA核心、Warp调度器和内存层级，最后到各种内存类型和互连技术。你可以把它当作学习路线图。
写在最后
GPU硬件架构的知识，说白了就是理解资源在哪里、速度有多快、怎么用最划算。SM是计算的核心，warp调度是隐藏延迟的关键，内存层级决定了数据访问的效率，而PCIe/NVLink则影响着多GPU系统的通信。
我刚开始学CUDA时，总觉得这些硬件细节太枯燥。直到后来做项目，一次次被性能问题折磨，才明白不懂硬件，优化无从谈起。希望这一章能帮你打下坚实的基础。
下一章，我们会深入CUDA编程模型，看看怎么用代码把这些硬件资源用起来。嗯，到时候见。