3、GPU架构基础:GPU核心架构(流处理器、线程束、内存层次)、SIMT执行模型、NVIDIA与AMD GPU架构对比
好,咱们今天聊聊GPU的底子。说实话,很多做异构计算的工程师,一开始都容易把GPU当成“很多核的CPU”。这个理解其实挺危险的。我刚开始接触GPU编程那会儿,也犯过这个错——拿CPU的思维去写GPU代码,结果性能惨不忍睹。
GPU和CPU的设计哲学完全不同。CPU是“单兵作战能力强”,GPU是“人海战术”。你想想看,一个CPU核心能处理复杂的逻辑分支、乱序执行、大缓存预取,而一个GPU核心呢?它只做一件事:算。简单粗暴,但效率极高。
3.1 流处理器(Streaming Multiprocessor, SM)
NVIDIA把它的计算单元叫做SM(Streaming Multiprocessor),AMD叫CU(Compute Unit)。名字不同,但本质差不多——都是GPU内部的基本计算模块。
一个SM里面包含什么?我直接说关键:
- CUDA核心(NVIDIA的叫法):就是最基础的算术逻辑单元。一个SM里通常有几十到上百个。
- 共享内存:这块内存很小,但速度极快。我在项目中经常用它来做线程间的数据交换,比全局内存快两个数量级。
- 寄存器文件:每个线程都有自己的寄存器,但总量有限。我曾经踩过一个坑——寄存器用太多,导致每个SM能同时跑的线程数大幅下降,性能直接腰斩。
- 线程束调度器:这个后面细说,它是SM的大脑。
核心要点:SM是物理概念,线程块(Thread Block)是逻辑概念。一个线程块会被调度到某个SM上执行。SM里的资源(寄存器、共享内存)是所有线程块共享的,所以资源分配要精打细算。
3.2 线程束(Warp)与SIMT执行模型
这是GPU最精髓的地方。CPU是SIMD(单指令多数据),GPU是SIMT(单指令多线程)。有什么区别?
SIMD是显式的——程序员得自己写向量化代码。SIMT是隐式的——程序员写的是线程,硬件自动把32个线程打包成一个Warp(NVIDIA的叫法,AMD叫Wavefront,通常是64个线程)。
我举个例子你就明白了:
// 假设我们处理一个数组
__global__ void add(int *a, int *b, int *c) {
int idx = threadIdx.x + blockIdx.x * blockDim.x;
c[idx] = a[idx] + b[idx];
}
这段代码看起来是每个线程独立执行。但实际上,硬件会把连续的32个线程(比如线程0到31)组成一个Warp,它们在同一时刻执行同一条指令。这就是SIMT的核心——一个Warp里的线程,执行相同的指令,但操作不同的数据。
注意:如果Warp里的线程走了不同的分支(比如if-else),那就会发生“线程束分化”。我曾经在一个项目里,因为一个简单的if判断,导致Warp里的线程一半走if,一半走else,性能直接掉了50%。解决办法?尽量让同一个Warp里的线程走相同的分支。
为什么会这样?因为SM里的线程束调度器一次只能给一个Warp发一条指令。如果线程走的分支不同,那硬件只能先执行if分支,再执行else分支,相当于串行化了。
3.3 内存层次:从快到慢
GPU的内存层次,说白了就是“越快的越贵,越贵的越小”。我画个简化的层次:
| 内存类型 | 位置 | 速度 | 容量 | 作用域 |
|---|---|---|---|---|
| 寄存器 | SM内部 | ~1周期 | 几十KB/SM | 单个线程 |
| 共享内存 | SM内部 | ~5周期 | 几十~上百KB/SM | 线程块内 |
| L1缓存 | SM内部 | ~10周期 | 几十~上百KB/SM | SM内 |
| L2缓存 | 芯片共享 | ~100周期 | 几MB | 全局 |
| 全局内存 | 显存 | ~400周期 | 几GB~几十GB | 全局 |
嗯,这里要注意:共享内存和L1缓存通常是共享同一块物理SRAM的。你可以配置它们的比例。我个人的习惯是,如果线程块内需要频繁交换数据,就多分点给共享内存;如果数据局部性比较好,就多分点给L1。
避坑指南:我曾经在做一个矩阵转置的优化时,发现全局内存访问模式不对——线程访问的地址是跨步的,导致每次内存访问都浪费带宽。后来改成共享内存做中转,性能提升了3倍。记住:合并访问是GPU内存优化的第一原则。
3.4 NVIDIA vs AMD:架构对比
说实话,这两家的架构思路挺不一样的。我简单对比一下:
| 特性 | NVIDIA(以Ampere为例) | AMD(以RDNA2为例) |
|---|---|---|
| 基本计算单元 | SM(含128个CUDA核心) | CU(含64个流处理器) |
| 线程束大小 | 32线程/Warp | 64线程/Wavefront |
| 共享内存 | 可配置,最大164KB/SM | 固定大小,通常128KB/CU |
| 编程模型 | CUDA(专有) | ROCm/HIP(开源) |
| 缓存架构 | 统一L1/共享内存 | 分离的L1和共享内存 |
我个人觉得,NVIDIA的架构更灵活一些——共享内存和L1可以动态调整,适合不同场景。AMD的架构更固定,但好处是延迟更可预测。
还有一个关键区别:线程束大小。NVIDIA是32,AMD是64。这意味着什么?如果你写一个简单的向量加法,NVIDIA的Warp里32个线程同时干活,AMD的Wavefront里64个线程同时干活。但Wavefront越大,分支分化的代价也越大——64个线程里只要有一个走不同分支,整个Wavefront就得串行化。
我记得有一次帮客户迁移代码,从NVIDIA的CUDA转到AMD的ROCm。代码逻辑完全一样,但性能就是上不去。后来发现是Wavefront大小不同导致的——原来针对32线程优化的分支逻辑,到了64线程上,分支分化的概率翻了一倍。解决办法?重新组织数据,让同一个Wavefront里的线程尽量走相同路径。
总结一下:GPU架构的核心就三件事——流处理器怎么组织、线程怎么调度、内存怎么访问。理解了这三件事,你就能写出高效的GPU代码。别被那些花里胡哨的术语吓到,说白了就是:让每个线程都忙起来,让每次内存访问都值回票价。
下一章咱们聊聊CPU和GPU怎么协同工作——这才是异构计算真正有意思的地方。