2. GPU硬件架构
大家好，我是你们的GPU架构讲师。今天咱们聊聊GPU的“骨架”——硬件架构。说实话，搞GPU优化这么多年，我最大的体会就是：不懂硬件，优化就是瞎蒙。你想想看，连SM长啥样、数据怎么搬的都不知道，怎么写出高性能代码？
这一节，我会从三个核心维度展开：SM（流多处理器）的内部结构、内存层次体系，以及PCIe与NVLink总线。嗯，这三个东西，基本决定了你程序的性能天花板。
一、SM（流多处理器）—— GPU的“心脏”
SM 是什么？说白了，它就是 GPU 里干活的核心单元。一个 GPU 有几十到上百个 SM，每个 SM 又包含一堆计算单元。我习惯把 SM 想象成一个“小工厂”——里面有工人（CUDA Core）、有工头（Warp Scheduler）、有临时仓库（寄存器/共享内存）。
1.1 CUDA Core 与 Tensor Core
每个 SM 里包含几十到上百个 CUDA Core。这些是执行浮点运算的主力。我记得在 Volta 架构刚出来时，NVIDIA 加入了 Tensor Core，专门做矩阵乘加。当时我在做深度学习推理优化，看到 Tensor Core 的吞吐量，直接惊了——比普通 CUDA Core 快了一个数量级。
- CUDA Core：通用计算单元，处理 int32、fp32、fp64。
- Tensor Core：专为矩阵运算设计，支持 fp16、bf16、int8 等混合精度。
- 特殊函数单元（SFU）：处理 sin、cos、sqrt 等复杂指令。
关键点： 每个 SM 的 CUDA Core 数量决定了它的计算密度。但别只看数量——调度效率才是瓶颈。我曾经在项目中遇到一个情况：理论算力 10 TFLOPS，实际只跑出 2 TFLOPS。查了半天，发现是 warp 发散太严重。
1.2 Warp 调度器与寄存器
SM 内部有多个 Warp 调度器。每个调度器负责管理一组 warp（32 个线程）。嗯，这里有个坑：warp 是 GPU 调度的基本单位，不是线程。你写代码时觉得是 1024 个线程在跑，实际上硬件是按 32 个一组来执行的。
寄存器是 SM 里最快的存储，但数量有限。每个 SM 有几十万到上百万个 32 位寄存器。我建议你写 kernel 时，每个线程的寄存器使用量控制在 32 个以内。超过这个数，编译器会 spill 到 local memory，性能直接腰斩。我曾经调试过一个 kernel，寄存器用了 64 个，性能只有优化后的 40%。
个人经验： 用 --ptxas-options=-v 编译选项查看寄存器使用量。如果看到 "spill stores" 或 "spill loads"，赶紧优化。
1.3 共享内存与 L1 缓存
共享内存是 SM 内部的一块 SRAM，由程序员显式管理。它的延迟比全局内存低一个数量级（约 20-30 周期 vs 400-800 周期）。我习惯把共享内存当作“手边的工具台”——频繁访问的数据放上去，别每次都去仓库（全局内存）拿。
L1 缓存是自动管理的，但和共享内存共用同一块物理空间。在 NVIDIA 的架构里，你可以通过 cudaFuncSetAttribute 调整两者的比例。比如，计算密集型任务可以多分给 L1，数据重用多的任务多分给共享内存。
注意： 共享内存容量很小，每个 SM 通常只有 48KB-164KB。别一股脑把大数组放进去。我曾经见过有人把 1MB 的数据塞共享内存——结果编译直接报错。
二、内存层次结构——数据流动的“高速公路”
GPU 的内存层次，说白了就是“速度越快的越贵、容量越小”。理解这个层次，你才能知道数据该放哪。
| 内存类型 | 位置 | 延迟（周期） | 容量 | 管理方式 |
|---|---|---|---|---|
| 寄存器 | SM 内部 | 1 | ~256KB/SM | 编译器自动 |
| 共享内存 / L1 | SM 内部 | ~20-30 | 48-164KB/SM | 显式 / 自动 |
| L2 缓存 | 芯片全局 | ~200 | 几 MB - 几十 MB | 自动 |
| 全局内存（HBM/GDDR） | 显存颗粒 | ~400-800 | 几 GB - 几十 GB | 显式 |
| 常量内存 | 芯片全局 | ~200（有缓存） | 64KB | 显式 |
2.1 全局内存——主战场
全局内存是 GPU 最大的存储空间，也是性能瓶颈的重灾区。它的带宽很高（HBM2e 可达 2 TB/s），但延迟也高。我经常跟学员说：别让线程闲着等数据。解决办法就是“合并访问”——让同一个 warp 的线程访问连续的内存地址。
举个例子：
// 好的合并访问
float val = data[threadIdx.x]; // 连续地址
// 坏的合并访问
float val = data[threadIdx.x * 32]; // 跨步访问，带宽利用率低我曾经优化过一个矩阵转置的 kernel，原始版本带宽利用率只有 15%。通过调整数据布局和 tile 大小，最终达到了 85%。说白了，就是让内存访问模式更“整齐”。
2.2 L2 缓存——自动的“加速器”
L2 缓存是所有 SM 共享的。它的作用就是减少全局内存的访问次数。如果你的数据有局部性（比如反复访问同一块区域），L2 能帮你省不少时间。但注意，L2 不是万能的——如果数据访问模式是随机的，L2 命中率会很低。
我建议： 用 ncu（NVIDIA Nsight Compute）分析 L2 命中率。如果低于 50%，考虑优化数据布局或使用共享内存。
2.3 常量内存与纹理内存
常量内存适合存储只读的、所有线程都访问相同地址的数据（比如模型参数）。它有专门的缓存，延迟低。纹理内存则适合二维/三维空间局部性访问（比如图像处理）。
嗯，这两个用得相对少一些。我个人只在做图像滤波和查找表时用过纹理内存。常量内存倒是经常用——比如把神经网络的权重放进去。
三、PCIe 与 NVLink——GPU 的“血管”
数据怎么从 CPU 搬到 GPU？怎么在多卡之间传输？这就是总线干的事。
3.1 PCIe 总线
PCIe 是 CPU 和 GPU 之间的标准接口。目前主流是 PCIe Gen4 x16，单向带宽约 32 GB/s。Gen5 翻倍到 64 GB/s。但说实话，这个带宽和 GPU 内部动辄 TB/s 的显存带宽比，差了两个数量级。
所以，尽量减少 CPU-GPU 数据传输。我见过有人每个迭代都从 CPU 拷贝数据到 GPU，结果 90% 的时间花在传输上。正确的做法是：一次拷贝，多次计算。
避坑指南： 我曾经在项目中用 cudaMemcpy 做异步传输，但忘了加 cudaStreamSynchronize，结果数据还没传完就开始计算，读到的全是垃圾。嗯，血的教训。
3.2 NVLink 与 NVSwitch
NVLink 是 NVIDIA 的多卡互联技术。它的带宽远高于 PCIe——比如 A100 的 NVLink 3.0，单向带宽 600 GB/s，是 PCIe Gen4 的 18 倍。NVSwitch 则是一个交换机，让多卡之间实现全互联。
如果你做多卡训练或推理，NVLink 是必须的。我建议用 nvidia-smi topo -m 查看 GPU 之间的拓扑结构。如果两张卡之间没有 NVLink，数据传输会走 PCIe，性能会差很多。
举个例子：在 8 卡 A100 上做数据并行训练，如果卡之间走 NVLink，all-reduce 通信时间只占 5%；如果走 PCIe，通信时间可能占到 30% 以上。
核心观点： 总线带宽决定了多卡扩展的效率。NVLink 让多卡像一台 GPU 一样工作，PCIe 则像多台独立的机器。
小结
好了，这一节的内容就到这里。SM 结构决定了计算效率，内存层次决定了数据访问速度，总线决定了数据传输能力。这三者加起来，就是 GPU 性能的“铁三角”。