硬件基础：GPU架构概览（SM、Warp、Memory Hierarchy）、CPU与GPU协同
各位同学，今天我们聊聊GPU的硬件架构。说实话，很多做AI训练的同学，写了半年CUDA，可能都没搞清楚GPU内部到底怎么工作的。我个人觉得，不理解硬件，写出来的算子性能就是碰运气。
先问大家一个问题：为什么GPU能比CPU快那么多？
答案其实很简单——GPU是「人多力量大」的典型代表。CPU是几个精英干将，GPU是成千上万的普通工人。但要让这些工人高效协作，你得懂他们的组织架构。
SM：GPU的基本计算单元
SM，全称Streaming Multiprocessor，流式多处理器。你可以把它想象成一个「小工厂」。
每个SM内部有自己的计算核心、寄存器、共享内存、调度器。一个GPU有多少个SM，基本决定了它的算力上限。
我记得第一次看NVIDIA的架构白皮书时，有个数据让我印象很深：A100有108个SM，每个SM有64个FP32核心。算下来总共6912个核心。你想想看，一个CPU才几个核？
但这里有个坑——并不是SM越多性能就线性增长。为什么？因为还有内存带宽、调度效率这些瓶颈。我曾在项目里遇到过，把算子从V100迁移到A100，SM数量翻倍了，但性能只提升了40%。后来一查，是内存访问模式没优化好。
Warp：GPU的调度单位
Warp这个词，直译是「 warp 扭曲」，但在GPU里，它指的是32个线程组成的一个执行单元。
为什么是32？这是NVIDIA的设计选择。32个线程在SM里以SIMT（单指令多线程）方式执行。说白了，就是32个人听同一个指令，但每个人处理不同的数据。
这里有个非常重要的概念——Warp Divergence（线程束分化）。
假如你的代码里有if-else分支，32个线程里有些走if，有些走else。那SM会怎么做？它会先执行if路径，把走else的线程「屏蔽」掉，然后再反过来执行else路径。这样一来，性能直接砍半。
所以，写CUDA代码时，尽量让同一个warp里的线程走相同的执行路径。这是最基本的优化意识。
Memory Hierarchy：GPU的存储体系
GPU的内存层次，说白了就是「越靠近计算单元，越快但越小」。我们从上往下看：
| 存储层级 | 大小 | 带宽 | 延迟 | 作用域 |
|---|---|---|---|---|
| 寄存器 | 每个SM约64K~256K | 最高 | ~1 cycle | 单个线程 |
| 共享内存 | 每个SM 48KB~164KB | 很高 | ~20 cycles | 单个Block |
| L1缓存 | 每个SM 128KB~192KB | 高 | ~30 cycles | 单个SM |
| L2缓存 | 整个GPU 40MB~80MB | 中等 | ~200 cycles | 整个GPU |
| 全局内存 | 整个GPU 40GB~80GB | 低 | ~400-800 cycles | 所有线程 |
看到这个延迟差距了吗？寄存器1个周期，全局内存要400个周期。你想想看，如果每个线程都频繁访问全局内存，那GPU大部分时间都在「等数据」，算力再高也白搭。
我个人的习惯是：能放寄存器的放寄存器，能放共享内存的放共享内存，实在不行再考虑全局内存。这个优先级顺序，是优化性能的第一原则。
CPU与GPU协同：异构计算的精髓
CPU和GPU怎么配合？说白了就是「CPU管调度，GPU管计算」。
典型的流程是这样的：
- CPU把数据从主存拷贝到GPU显存
- CPU启动GPU kernel（核函数）
- GPU执行计算
- CPU等待GPU完成（或者做其他事）
- GPU把结果拷贝回主存
这里有个关键点——数据传输是瓶颈。PCIe 4.0的带宽是32GB/s，而GPU显存带宽是2TB/s。差了60多倍。所以，尽量减少CPU和GPU之间的数据传输。
我记得有个项目，需要做大量小batch的推理。每次推理都从CPU拷贝数据到GPU，然后启动一个很小的kernel。结果发现，数据传输时间占了总时间的80%。后来我改成「批量处理」，把多个小batch拼成一个大batch，一次性传输，性能提升了5倍。
另外，CPU和GPU可以异步执行。CPU启动kernel后，不用傻等，可以去做别的事。比如准备下一批数据。这就是所谓的「流水线」技术。
// 伪代码示例：CPU与GPU异步流水线
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
for (int i = 0; i < num_batches; i++) {
// 在stream1上传输数据
cudaMemcpyAsync(d_input[i], h_input[i], size,
cudaMemcpyHostToDevice, stream1);
// 在stream2上执行计算
myKernel<<<grid, block, 0, stream2>>>(d_input[i], d_output[i]);
// 在stream1上回传结果
cudaMemcpyAsync(h_output[i], d_output[i], size,
cudaMemcpyDeviceToHost, stream1);
}
cudaDeviceSynchronize();你看，这样数据传输和计算就重叠了。GPU在算第i批的时候，CPU已经在准备第i+1批的数据了。嗯，这才是高效利用硬件的方式。
知识体系总览
下面这张图，是我自己画的GPU架构概览。你可以把它当作本章的「思维导图」：
这张图里，你可以看到CPU通过PCIe连接到GPU，GPU内部有多个SM，每个SM有自己的寄存器、共享内存、L1缓存和Warp调度器。所有SM共享L2缓存和全局内存。
嗯，这就是GPU架构的核心。理解了这些，你写CUDA代码时就会知道：
- 为什么共享内存比全局内存快20倍
- 为什么warp divergence会拖慢性能
- 为什么数据传输是瓶颈
- 为什么需要让GPU「吃饱」
这些不是理论，是我在实际项目中踩过的坑。希望你能少走弯路。