3、GPU计算基础：并行计算概念、SIMT模型、线程层次结构、内存层次结构

好，咱们进入正题。这一章是GPU编程的基石。说白了，你后面写的所有核函数，都跑不出今天讲的这几个概念。我当年刚接触CUDA时，就是没搞懂线程层次和内存模型，结果写出来的程序比CPU还慢。嗯，那滋味可不好受。

3.1 并行计算：从串行到并行的思维转变

我们平时写C代码，是一条指令一条指令地执行。这叫串行。GPU不一样，它一次能跑成千上万个线程。这就是并行。

并行计算的核心思想很简单：把一个大任务拆成很多小任务，同时处理。你想想看，一个数组有100万个元素，你要给每个元素加1。CPU得循环100万次。GPU呢？开100万个线程，每个线程只处理一个元素。一秒搞定。

但这里有个坑。不是所有任务都能并行。有些任务有数据依赖——比如计算斐波那契数列，后一项依赖前一项。这种任务强行并行，反而会出错。我在项目中遇到过一位同事，把有依赖的循环直接改成并行，结果数据全乱了。排查了两天才找到原因。

3.2 SIMT模型：单指令多线程

GPU的并行模型叫SIMT——Single Instruction, Multiple Threads。翻译过来就是「单指令，多线程」。

什么意思呢？就是一组线程，在同一时刻，执行同一条指令。但每个线程处理的数据可以不同。这和CPU的SIMD（单指令多数据）有点像，但更灵活。

举个例子。你写了一句代码：c = a + b。在GPU上，32个线程会同时执行这条加法指令。但每个线程拿到的a和b，可能是数组里不同位置的元素。

为什么会这样设计？因为GPU的硬件架构就是这样。它把32个线程绑成一个组，叫Warp。一个Warp里的线程，步调必须一致。如果某个线程走了一条不同的分支（比如if-else），那其他线程就得等着。这就是所谓的「线程束发散」。

3.3 线程层次结构：Grid / Block / Warp

GPU的线程组织方式，分三个层级。我习惯从大到小来讲。

层级	名称	说明
最大	Grid（网格）	一个核函数启动时，对应一个Grid。包含多个Block。
中间	Block（线程块）	一组线程的集合。Block内的线程可以共享内存、同步。
最小	Thread（线程）	实际执行计算的最小单元。

再往下，还有硬件层面的Warp。一个Warp包含32个线程。Block里的线程，会被分成若干个Warp来调度。

你写代码时，需要指定Grid和Block的维度。比如：

// 定义Grid和Block的维度
dim3 gridDim(4, 4);      // 4x4 = 16个Block
dim3 blockDim(32, 32);   // 32x32 = 1024个线程/Block

// 启动核函数
myKernel<<<gridDim, blockDim>>>(d_data);

这段代码启动了16个Block，每个Block有1024个线程。总共16384个线程。每个线程可以通过blockIdx和threadIdx来知道自己是谁。

3.4 内存层次结构：速度与容量的博弈

GPU的内存，和CPU一样，也有层次结构。越靠近计算单元的内存，速度越快，但容量越小。我画个简化的层次图给你看：

内存类型	位置	速度	容量	作用域
寄存器	每个线程私有	最快	极小（几十个）	单个线程
共享内存	每个Block私有	快	小（几十KB）	Block内所有线程
全局内存	所有线程共享	慢	大（几GB）	所有线程 + Host
常量内存	只读	快（有缓存）	小（64KB）	所有线程
纹理内存	只读	快（有缓存）	大	所有线程

这里我要重点说一下共享内存。它是个宝贝。Block内的线程可以通过共享内存交换数据，速度比全局内存快两个数量级。但要注意，共享内存是手动管理的——你得自己把数据从全局内存搬到共享内存里。

我举个例子。矩阵乘法中，每个输出元素需要读取一整行和一整列的数据。如果每次都从全局内存读，带宽根本不够。正确的做法是：把子矩阵加载到共享内存里，然后让Block内的线程协作计算。

// 共享内存声明（在核函数内部）
__shared__ float shared_data[BLOCK_SIZE][BLOCK_SIZE];

// 从全局内存加载到共享内存
shared_data[threadIdx.y][threadIdx.x] = global_data[global_idx];
__syncthreads();  // 同步！确保所有线程都加载完了

// 然后从共享内存读取，进行计算
float val = shared_data[threadIdx.y][threadIdx.x];

3.5 全局内存的访问模式：合并访问

全局内存虽然慢，但如果你访问方式正确，也能榨出不错的带宽。关键就是合并访问。

什么意思呢？当Warp里的32个线程，访问连续的内存地址时，硬件会把这32次访问合并成一次大的内存事务。效率极高。反之，如果每个线程访问的地址是随机分散的，那就得拆成很多次小事务，带宽利用率直线下降。

我建议你记住这个原则：让相邻的线程访问相邻的地址。比如：

// 好的访问模式：合并访问
int idx = blockIdx.x * blockDim.x + threadIdx.x;
float val = data[idx];  // 线程0访问data[0]，线程1访问data[1]...

// 坏的访问模式：非合并访问
int idx = blockIdx.x * blockDim.x + threadIdx.x;
float val = data[threadIdx.x * gridDim.x + blockIdx.x];  // 地址跳跃

嗯，这个优化点，在后面的章节中会反复出现。你现在先有个印象。

3.6 小结

这一章的内容，是GPU编程的「内功心法」。并行计算是思想，SIMT是硬件模型，线程层次是组织方式，内存层次是资源约束。四者缺一不可。

我个人觉得，初学者最容易犯的错误是：只关注算法逻辑，忽略了内存访问模式。结果写出来的核函数，计算效率只有理论峰值的十分之一。所以，我建议你每写一个核函数，都问自己三个问题：

1. 我的线程组织方式合理吗？Block大小是不是32的倍数？
2. 我的内存访问是合并的吗？相邻线程访问相邻地址了吗？
3. 我有没有用到共享内存？能不能把重复访问的数据缓存起来？

把这三个问题想清楚，你的GPU程序基本就成功了一半。下一章，我们会动手写第一个CUDA程序，到时候这些概念都会用上。