第2章：GPU硬件架构：流处理器（SM/CU）详解、显存层次结构（HBM/GDDR）、PCIe与NVLink互联
大家好，我是你们的老朋友。今天我们来聊聊GPU的硬件架构。说实话，很多搞GPU编程的朋友，写了几年代码，却对芯片内部长什么样一知半解。我个人觉得，不理解硬件，你写的代码就像在黑暗中摸索——能跑，但跑不快。
这一章，我们聚焦三个核心模块：流处理器（SM/CU）、显存层次结构，以及互联总线（PCIe/NVLink）。这三样东西，说白了就是GPU的“心脏”、“粮仓”和“血管”。
2.1 流处理器：SM与CU
先聊聊SM（Streaming Multiprocessor）。NVIDIA叫它SM，AMD叫它CU（Compute Unit）。名字不同，本质一样——都是GPU里干活的最小独立单元。
一个SM里面有什么？我给大家拆开看看：
- CUDA Core / 流处理器单元：真正干算术的。一个SM里通常有64~128个。
- Tensor Core：专门干矩阵乘法的。AI训练全靠它。
- Warp Scheduler：负责调度线程。一次调度32个线程（一个warp）。
- Shared Memory：SM内部的高速缓存。程序员可以手动控制。
- Register File：寄存器堆。每个线程私有，容量有限。
核心概念：SM是硬件概念，Thread Block是软件概念。一个Thread Block会被调度到一个SM上执行。SM里的资源（寄存器、共享内存）是有限的，决定了你能开多少个线程。
我记得有一次做性能优化，发现kernel occupancy（占用率）只有30%。查了半天，原来是每个线程用了太多寄存器，导致SM能同时容纳的线程块数量骤减。后来我强制限制了寄存器使用量（__launch_bounds__），占用率直接拉到80%，性能翻倍。嗯，这就是理解SM资源的意义。
2.2 显存层次结构
显存，就是GPU的“粮仓”。但粮仓也有远近之分。离计算单元越近，速度越快，容量越小。
典型的显存层次如下：
| 层级 | 典型容量 | 带宽 | 延迟 | 作用 |
|---|---|---|---|---|
| HBM / GDDR（全局显存） | 16~80 GB | ~2 TB/s（HBM2e） | ~200~400 cycles | 存储所有数据 |
| L2 Cache | ~40 MB | ~4 TB/s | ~100 cycles | 缓存全局显存数据 |
| L1 Cache / Shared Memory | ~128 KB per SM | ~10 TB/s | ~30 cycles | 线程块内共享数据 |
| Register | ~256 KB per SM | 极高 | 1 cycle | 线程私有变量 |
你想想看，从HBM读一个数要400个时钟周期，而从寄存器读只要1个周期。这差距有多大？所以优化的核心思路就是：让数据尽量靠近计算单元。
我的经验：写CUDA代码时，我习惯先用Shared Memory做数据复用。比如矩阵分块，把子块先搬到Shared Memory里，再让线程去算。这样能减少90%以上的全局显存访问。我曾经在一个图像处理项目里，靠这招把处理时间从12ms降到了3ms。
HBM和GDDR有什么区别？简单说：
- HBM（High Bandwidth Memory）：堆叠式设计，带宽极高，但容量相对小。常用于高端计算卡（A100、H100）。
- GDDR（Graphics DDR）：传统设计，容量大，带宽也不错。常用于游戏卡（RTX 4090用GDDR6X）。
说白了，HBM是跑车，GDDR是卡车。跑车快但装得少，卡车装得多但慢一点。
2.3 PCIe与NVLink互联
GPU不是孤岛。它需要和CPU通信，也需要和其他GPU通信。这就靠互联总线。
PCIe（Peripheral Component Interconnect Express）
这是最通用的互联方式。CPU和GPU之间，通过PCIe总线交换数据。目前主流是PCIe 4.0 x16，单向带宽约16 GB/s。PCIe 5.0翻倍到32 GB/s。
但问题来了：GPU内部显存带宽是2 TB/s，PCIe只有16 GB/s。差了100多倍！所以PCIe是典型的瓶颈。我经常跟团队说：尽量减少CPU和GPU之间的数据搬运。能一次传完就别分十次。
避坑指南：我曾经遇到一个项目，每次迭代都把数据从CPU拷到GPU，算完再拷回去。结果PCIe带宽占满，GPU利用率不到20%。后来改成异步传输 + 流水线（pipeline），把计算和传输重叠起来，利用率才拉到80%以上。记住：PCIe传输是同步的噩梦。
NVLink
NVLink是NVIDIA自家的高速互联技术。它绕过了PCIe，直接在GPU之间建立高速通道。NVLink 3.0的单向带宽可达50 GB/s，是PCIe 4.0的三倍多。
NVLink有什么用？多卡训练！比如你用8张A100做分布式训练，NVLink让卡间通信几乎无延迟。我见过一个场景：用NVLink互联的8卡系统，训练吞吐量比纯PCIe互联高出40%。
为什么？因为NVLink支持GPU Direct RDMA，数据可以直接从一个GPU的显存传到另一个GPU的显存，完全不用经过CPU和系统内存。这就像两个邻居直接开了一扇门，不用再绕到小区大门了。
总结一下：
- 单卡优化：关注SM资源、显存层次、Shared Memory使用。
- 多卡优化：关注NVLink带宽、GPU Direct、异步传输。
- CPU-GPU通信：关注PCIe带宽，尽量减少传输次数。
好了，这一章就到这里。理解这些硬件细节，你写代码的时候心里就有谱了。下一章我们聊聊线程模型和warp调度——那才是真正让SM跑起来的核心机制。
公众号:蓝海资料掘金营，微信deep3321