GPU计算基础：从架构演进到协同工作
大家好，我是你们的老朋友。今天咱们聊聊GPU计算最核心的东西——架构、编程模型，还有CPU和GPU怎么配合干活。这些东西看着基础，但我在实际项目中见过太多人栽跟头了。说白了，不理解底层，你写出来的核函数性能就是上不去。
一、GPU架构演进：从图形卡到通用计算利器
GPU最早是干嘛的？给游戏渲染画面的。但后来大家发现，这玩意儿做并行计算太猛了。我2012年刚入行那会儿，用的还是Fermi架构，当时觉得已经够强了。现在回头看，真是小巫见大巫。
核心架构演进路线：
- Fermi (2010)：首次引入真正的L1/L2缓存层次，支持ECC。我当年用GTX 580跑科学计算，ECC帮了大忙。
- Kepler (2012)：引入SMX单元，每个SMX有192个CUDA核心。功耗控制做得不错。
- Maxwell (2014)：进一步优化能效比，L2缓存翻倍。我印象最深的是它的共享内存带宽提升明显。
- Pascal (2016)：统一内存模型来了！CPU和GPU终于能共享指针了。嗯，这里要注意，统一内存不是万能的，性能开销你得心里有数。
- Volta (2017)：Tensor Core横空出世，AI训练直接起飞。独立的线程调度栈，让warp调度更灵活。
- Turing (2018)：加入RT Core，实时光线追踪。同时Tensor Core升级到第二代。
- Ampere (2020)：第三代Tensor Core，支持稀疏计算。L2缓存飙到40MB，我跑大模型时明显感觉缓存命中率上去了。
- Hopper (2022)：Transformer Engine，专门为AI设计。DPX指令集，动态规划加速。
你想想看，从Fermi到Hopper，CUDA核心数量翻了上百倍，内存带宽从不到200GB/s飙到3TB/s以上。为什么会这样？因为应用需求在推着硬件走。我做过一个流体力学模拟项目，用Kepler跑要3天，换到Ampere上只要6小时——这就是架构演进的力量。
二、CUDA编程模型：你写的代码到底怎么跑？
CUDA编程模型，说白了就是三个抽象层次：线程(Thread)、线程块(Block)、网格(Grid)。我刚开始学的时候，总觉得这玩意儿绕。后来想明白了，这就是GPU的"组织架构"。
我的个人习惯：写核函数前，先画个线程布局图。把数据怎么分、线程怎么排，在纸上画清楚。这招帮我避免了很多逻辑错误。
来看个最简单的例子：
// 向量加法：C = A + B
__global__ void vecAdd(float *A, float *B, float *C, int N) {
int idx = threadIdx.x + blockIdx.x * blockDim.x;
if (idx < N) {
C[idx] = A[idx] + B[idx];
}
}
// 调用方式
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(A, B, C, N);这里有几个关键点：
- threadIdx：线程在块内的索引。我见过新手直接拿它当全局索引用，结果数据对不上。
- blockIdx：块在网格内的索引。配合blockDim，就能算出全局ID。
- blockDim：每个块的线程数。这个值不是随便设的，要考虑SM的资源限制。
我曾经踩过的坑：有一次做矩阵乘法，我设了1024个线程每块，结果跑起来特别慢。后来一查，每个线程用的寄存器太多，SM能同时驻留的线程块数量被严重限制了。记住：线程数不是越多越好，要平衡资源占用。
CUDA的执行模型是SIMT（单指令多线程）。什么意思？就是32个线程组成一个warp，它们执行同一条指令，但处理不同的数据。如果一个warp里的线程走了不同的分支（比如if-else），那性能就惨了——这叫warp divergence。
三、GPU与CPU协同工作原理：谁该干什么活？
CPU和GPU，一个像项目经理，一个像施工队。CPU擅长复杂逻辑控制、串行任务；GPU擅长简单重复、大规模并行计算。我见过有人非要用GPU做字符串解析，结果性能还不如单核CPU——这就是没搞清楚分工。
典型的协同工作流程：
- CPU准备数据：从磁盘读数据、做预处理、分配显存。
- CPU将数据拷贝到GPU：通过PCIe总线。这一步往往是瓶颈，我建议能用异步传输就别用同步的。
- CPU启动GPU核函数：CPU下发指令，GPU开始干活。CPU可以去做别的事（异步执行）。
- GPU执行计算：成千上万个线程同时跑。CPU可以轮询或者等中断。
- CPU将结果拷回：计算完成，数据回传。
这里有个概念叫流(Stream)。你可以把流想象成一条流水线。多个流可以并行执行，实现计算和传输的重叠。我做过一个视频处理项目，用3个流同时处理不同帧，吞吐量直接翻倍。
// 使用流实现异步传输和计算
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
// 在stream1中：传输数据A，计算A
cudaMemcpyAsync(d_A, h_A, size, cudaMemcpyHostToDevice, stream1);
kernelA<<<grid, block, 0, stream1>>>(d_A);
// 在stream2中：传输数据B，计算B
cudaMemcpyAsync(d_B, h_B, size, cudaMemcpyHostToDevice, stream2);
kernelB<<<grid, block, 0, stream2>>>(d_B);
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);避坑指南：我曾经在项目里同时开了8个流，结果发现PCIe带宽被争抢，每个流都在等。后来我改成2个流，配合流水线操作，效率反而更高。记住：流的数量不是越多越好，要根据硬件资源来定。
还有一个重要概念——内存层次。GPU有全局显存、共享内存、寄存器、局部内存、常量内存、纹理内存。每个的访问速度和容量都不一样。我建议你把最频繁访问的数据放共享内存，能快一个数量级。
| 内存类型 | 位置 | 访问延迟 | 作用域 | 生命周期 |
|---|---|---|---|---|
| 寄存器 | SM内部 | ~1 cycle | 单个线程 | 核函数执行期间 |
| 共享内存 | SM内部 | ~20 cycles | 线程块内 | 核函数执行期间 |
| 全局显存 | 显存颗粒 | ~400 cycles | 所有线程 | 程序运行期间 |
| 常量内存 | 显存颗粒 | ~400 cycles (有缓存) | 所有线程 | 程序运行期间 |
| 纹理内存 | 显存颗粒 | ~400 cycles (有缓存) | 所有线程 | 程序运行期间 |
嗯，这里要注意：共享内存虽然快，但容量有限（通常几十KB）。我见过有人把整个数据集塞共享内存，结果编译都过不了。合理做法是分块处理——把数据切成小块，一块一块搬到共享内存里算。
核心要点总结：
- GPU架构从Fermi到Hopper，核心数、带宽、专用单元都在进化
- CUDA编程模型：线程→线程块→网格，三层抽象要理解透
- CPU负责控制流和串行任务，GPU负责大规模并行计算
- 内存层次决定性能上限，共享内存是你的好朋友
- 流和异步操作能隐藏传输延迟，但别贪多
最后说一句：这些基础概念，你花一周时间搞透，后面写核函数能少走三个月弯路。我当年就是太着急上手写代码，结果调优调到头秃。基础不牢，地动山摇啊。