GPU架构与编程模型：从SM到Warp的微观世界
大家好，我是你们的老朋友。今天我们来聊聊GPU架构和CUDA编程模型。说实话，这部分内容我讲了不下几十遍，但每次备课还是会发现新的感悟。GPU这东西，你越深入越觉得它精妙。
NVIDIA CUDA核心架构：不止是“很多核心”
很多人以为GPU就是一堆CPU核心的简单堆叠。其实不然。我刚开始接触CUDA时也这么想，结果写出来的程序性能惨不忍睹。后来才明白，GPU的架构设计有它独特的哲学。
NVIDIA的GPU架构，从Fermi到Ampere再到Hopper，核心思想一直没变：用大量轻量级线程掩盖延迟。你想想看，CPU是“快而少”，GPU是“多而慢”。但“慢”不代表差，关键是让这些“慢”的线程同时干活，总吞吐量就上去了。
核心概念：CUDA核心 ≠ CPU核心。CUDA核心本质上是一个浮点运算单元（FPU），它没有完整的控制逻辑和缓存。真正负责调度的是SM（Streaming Multiprocessor）。
我在项目中遇到过一位同事，他非要把CUDA核心数当成CPU核心数来用，结果每个线程都做大量分支判断，性能直接腰斩。嗯，这就是没理解架构的代价。
SM与Warp调度：真正的“指挥官”
SM（流式多处理器）才是GPU的“大脑”。每个SM包含多个CUDA核心、共享内存、寄存器文件和调度器。而Warp，是SM调度的基本单位——32个线程为一组。
为什么会是32？这不是拍脑袋定的。我记得NVIDIA的工程师说过，这是经过大量仿真验证后的最优值。32个线程同时执行同一条指令，这就是SIMT（单指令多线程）的精髓。
| 概念 | 说明 | 我的经验 |
|---|---|---|
| SM | 硬件执行单元，包含所有计算资源 | SM数量决定了并行度上限 |
| Warp | 32个线程的调度组 | Warp内线程必须执行相同指令 |
| Warp调度器 | 每个SM有多个，零开销切换 | 这是掩盖延迟的关键 |
Warp调度有个特点：零开销上下文切换。CPU切换线程要保存寄存器、刷新缓存，开销很大。但GPU的Warp切换？说白了就是换个指针。为什么能做到？因为每个Warp的寄存器都是物理独立的，不需要保存恢复。
避坑指南：我曾经写过一个kernel，每个线程都做大量if-else分支。结果Warp内的32个线程走了不同分支，性能直接降到1/32。记住：Warp内分支越少越好，最好没有。
内存层次结构：数据搬运的艺术
GPU的内存层次，说白了就是“距离计算单元越近越快，但越小”。我习惯把它分成四层：
- Global Memory（全局内存）：显存，容量最大（几GB到几十GB），但延迟最高（几百个周期）。所有线程都能访问。
- Shared Memory（共享内存）：每个SM独有，容量小（几十KB到一百多KB），延迟低（几个周期）。同一个Block内的线程共享。
- Local Memory（本地内存）：每个线程私有，但物理上存在显存里。寄存器不够用时自动使用，性能很差。
- Register（寄存器）：最快，每个线程私有。但数量有限（每个SM几千到几万个）。
你想想看，如果每个线程都去Global Memory读数据，那延迟根本掩盖不住。我做过一个矩阵乘法的优化，从直接读Global Memory改成用Shared Memory做分块，性能提升了8倍。这就是内存层次结构的威力。
注意：Local Memory是个“隐形杀手”。我曾经调试一个程序，发现性能远低于预期。查了半天，原来是每个线程用了太多局部变量，寄存器放不下，自动溢出到Local Memory。解决办法？减少每个线程的变量数，或者调整编译器选项。
CUDA编程模型：Kernel、Grid、Block、Thread
CUDA的编程模型，说白了就是“怎么组织你的线程”。我刚开始学的时候，总觉得Grid、Block、Thread这些概念很抽象。后来画了一张图，一下子就明白了。
这张图我画了很多遍。你看，一个Grid包含多个Block，每个Block包含多个Thread。Kernel函数就是在Thread上执行的。我习惯这样理解：
- Kernel：你写的那个函数，用
__global__修饰。它是GPU执行的入口。 - Grid：一次Kernel调用产生的所有线程的集合。你可以理解成“任务的总量”。
- Block：Grid的划分单位。同一个Block内的线程可以协作（通过Shared Memory和同步）。
- Thread：最小的执行单元。每个Thread执行Kernel函数的一份拷贝。
为什么要有Block？说白了是为了让线程能协作。我做过一个图像处理的项目，每个像素点需要参考周围8个像素的值。如果每个线程独立去读Global Memory，那效率太低了。我把图像分成16x16的Block，每个Block内的线程先把数据加载到Shared Memory，然后互相协作计算。性能提升了5倍。
// 一个典型的CUDA Kernel
__global__ void vectorAdd(float *A, float *B, float *C, int N) {
// 计算当前线程的全局索引
int idx = blockIdx.x * blockDim.x + threadIdx.x;
// 边界检查
if (idx < N) {
C[idx] = A[idx] + B[idx];
}
}这段代码看起来简单，但里面有几个关键点：
blockIdx.x：当前Block在Grid中的索引blockDim.x：每个Block包含的线程数threadIdx.x：当前Thread在Block中的索引
我刚开始写CUDA时，经常搞混这几个变量。后来我总结了一个口诀：“Block定位大范围，Thread定位小范围”。blockIdx告诉你“你在哪个组”，threadIdx告诉你“你在组里的哪个位置”。
我的习惯：设计Block大小时，我通常选16x16或32x8。为什么？因为Warp是32个线程一组，Block的维度最好是Warp大小的整数倍。这样能避免Warp内出现“空洞”，浪费计算资源。
总结一下
GPU架构和CUDA编程模型，说白了就是两件事：理解硬件怎么干活，学会怎么组织线程。SM和Warp是硬件的核心，Grid/Block/Thread是软件的抽象。把这两层对应起来，你就能写出高效的CUDA程序。
我记得第一次用CUDA加速一个科学计算程序，从CPU的2小时缩短到GPU的3分钟。那种成就感，嗯，至今难忘。但背后的坑也不少，比如内存访问模式、Warp发散、寄存器溢出...这些我们后面会慢慢聊。
好了，今天就到这里。记住：理解架构，才能驾驭性能。