第2章：GPU架构详解：NVIDIA GPU架构演进、流多处理器（SM）、CUDA核心、内存层次结构、Warp与线程束

好，咱们进入正题。这一章，我打算把NVIDIA GPU的底裤扒开给你看。你可能会觉得，搞异构计算嘛，会写CUDA代码不就行了？嗯，我以前也这么想。直到有一次，我写了个核函数，性能死活上不去，折腾了两天，最后发现是线程束发散的问题。从那以后，我养成了一个习惯：写代码之前，先想想GPU里面到底是怎么跑的。

2.1 NVIDIA GPU架构演进：从Fermi到Hopper

先聊聊架构的演变。说白了，NVIDIA这些年就是在干一件事：让更多的计算单元，更快地拿到数据。

我个人把架构演进分成几个关键节点：

• Fermi (2010)：真正的转折点。它第一次引入了完整的L1/L2缓存层次，还有真正的ECC支持。我记得当时做科学计算，没有ECC简直要命，跑个三天三夜，结果一个比特翻转，全白干。
• Kepler (2012)：这代引入了动态并行，就是GPU自己可以启动新的核函数，不用CPU管。嗯，想法很好，但我在项目中试过，调度开销其实不小，小任务别用。
• Maxwell (2014)：能效比大幅提升。它把SM（流多处理器）拆成了更小的处理块，叫SMM。说白了，就是让调度更灵活，减少资源浪费。
• Pascal (2016)：统一内存地址空间来了！CPU和GPU终于可以共享指针了。我当年为了这个特性，重构了整整一个代码库，但值了。
• Volta (2017)：Tensor Core首次登场。这玩意儿专门为矩阵乘法设计的，AI训练直接起飞。做传统HPC的同事可能觉得跟自己没关系，但后来我发现，很多线性代数运算也能用它加速。
• Turing (2018)：加入了RT Core做光线追踪。游戏卡用户狂喜，但做通用计算的，其实更关注它改进的整数和浮点并发执行能力。
• Ampere (2020)：第三代Tensor Core，支持稀疏矩阵计算。还有，L2缓存翻倍了。我在做某个流体模拟项目时，L2缓存大小直接决定了能不能把整个网格数据塞进去，省了不知道多少全局内存访问。
• Hopper (2022)：引入了DPX指令集，专门加速动态规划算法。还有Transformer Engine，说白了就是为AI大模型量身定做的。

2.2 流多处理器（SM）：GPU的心脏

SM是什么？你可以把它想象成一个迷你CPU，但比CPU简单粗暴得多。一个GPU里，有几十到上百个SM。每个SM，才是真正干活的地方。

SM内部包含：

• CUDA核心（执行算术运算）
• 共享内存（L1缓存）
• 寄存器文件
• Warp调度器
• 加载/存储单元
• 特殊函数单元（比如算sin、cos、sqrt）

你想想看，一个SM要同时管理成百上千个线程。它怎么做到的？靠的就是硬件线程调度。CPU切换线程要保存上下文，开销很大。GPU呢？它每个线程的上下文都在寄存器里，切换就是换个指针的事，几乎零开销。

2.3 CUDA核心：没那么神秘

很多人以为CUDA核心就是GPU里的一个小处理器，跟CPU核心差不多。其实不是。CUDA核心本质上是一个高度简化的浮点/整数运算单元。它没有复杂的控制逻辑，没有分支预测，没有乱序执行。说白了，就是傻大黑粗地算。

一个SM里有多少个CUDA核心？不同架构不一样：

架构	每个SM的CUDA核心数
Fermi	32
Kepler	192
Maxwell	128
Pascal	64
Volta	64
Turing	64
Ampere	64
Hopper	128

注意看，Kepler有192个，但Pascal只有64个。是不是退步了？不是。因为Kepler的CUDA核心其实很弱，频率低，功能少。Pascal的单个核心强多了。所以，别只看数量，要看实际吞吐量。

2.4 内存层次结构：速度与容量的博弈

GPU的内存层次，跟CPU很像，但更极端。从上到下，速度越来越慢，容量越来越大：

1. 寄存器：最快，每个线程私有。一个SM里寄存器总数有限，比如每个SM有65536个32位寄存器。你每个线程用的寄存器越多，SM上能同时跑的线程就越少。这是个经典的trade-off。
2. 共享内存 / L1缓存：同一个线程块内的线程可以访问。速度很快，但容量小，通常几十KB到一百多KB。我在项目中经常用它来做数据重用，比如矩阵分块乘法，把小块数据先搬到共享内存里，再计算，能快好几倍。
3. L2缓存：所有SM共享。容量从几MB到几十MB。它缓存全局内存的访问。如果你的数据访问有局部性，L2能帮你省很多带宽。
4. 全局内存：就是显存。容量最大（几个GB到几十GB），但延迟最高（几百个时钟周期）。所有线程都能访问，但要做好合并访问，不然带宽利用率惨不忍睹。
5. 常量内存：只读，有缓存。适合存储所有线程都需要的常量数据，比如物理参数、查找表。
6. 纹理内存：也是只读，有专门的缓存和硬件插值单元。做图像处理时特别好用。

2.5 Warp与线程束：真正的执行单元

这是最容易被忽视，但最重要的概念。你写代码时，感觉是成千上万个线程在并行执行。但实际上，GPU是以32个线程为一组来执行的。这一组，就叫一个Warp（线程束）。

一个Warp里的32个线程，执行同一条指令。如果它们走不同的分支（比如if-else），那就会发生线程束发散。什么意思？就是一部分线程走if，另一部分走else，它们不能同时执行，只能串行。性能直接打折扣。

举个例子：

__global__ void divergent_kernel(float* data, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n / 2) {
        data[idx] = data[idx] * 2.0f;  // 一半线程走这里
    } else {
        data[idx] = data[idx] * 0.5f;  // 另一半走这里
    }
}

这个代码，如果n是1024，线程块大小是256，那每个线程块里有128个线程走if，128个走else。一个Warp有32个线程，如果这32个线程刚好一半走if一半走else，那这个Warp就要串行执行两条路径。性能损失50%。

怎么避免？尽量让同一个Warp里的线程走同一个分支。比如，让数据按某种规律排列，或者用位运算代替分支。

另外，Warp的调度也是门学问。当一个Warp在等待内存访问结果时，SM会立刻切换到另一个Warp来执行。这就是所谓的零开销上下文切换。所以，要隐藏内存延迟，你需要足够多的Warp。一般来说，每个SM上至少要有4-8个活跃的Warp，才能把内存延迟隐藏得比较好。

好了，这一章的内容就这些。GPU架构其实不复杂，但细节很多。你只要记住：SM是核心，Warp是执行单元，内存层次是瓶颈。下一章，我们会真正开始写CUDA代码，把这些理论用起来。