第2章 GPU核心概念

GPU核心概念：SIMT执行模型、Warp/Wavefront、线程块与网格

好，我们直接进入正题。今天聊的这几个概念，是理解GPU架构的基石。说白了，你搞懂了它们，再看CUDA或者ROCm的编程模型，就会有种「原来如此」的感觉。

我个人习惯把SIMT、Warp、线程块和网格，看作是一套「从硬件到软件」的抽象层级。咱们一层层剥开来看。

很多人刚接触GPU时，会把它跟CPU的SIMD（单指令多数据）混为一谈。其实差别很大。

SIMT，全称是Single Instruction, Multiple Threads。翻译过来就是「单指令，多线程」。你想想看，CPU的SIMD，比如AVX指令，一次操作8个float数据，但本质上还是一个线程在执行。而GPU的SIMT，是让多个独立的线程，在同一时刻执行同一条指令。

嗯，这里要注意：每个线程都有自己的寄存器、程序计数器。它们只是指令相同，但数据可以完全不同。我在项目中遇到过不少新手，以为Warp里的线程必须走一样的分支，其实不是的——只是效率会打折。

核心要点：SIMT模型下，每个线程是独立的执行实体。硬件通过Warp来管理这些线程的并行执行。

为什么会这样设计？因为GPU要处理海量的数据并行任务。如果每个线程都独立取指、译码，那芯片面积和功耗根本扛不住。所以硬件把多个线程打包成一个组，共享指令流。这就是Warp的由来。

在NVIDIA的术语里叫Warp，在AMD那边叫Wavefront。名字不同，本质一样——都是硬件调度和执行的最小单位。

以NVIDIA为例，一个Warp包含32个线程。这32个线程在同一个周期里执行同一条指令。如果遇到分支，比如if-else，那部分线程会被「屏蔽」掉，等另一条路径走完再切换回来。这就是所谓的「分支发散」问题。

我曾经在一个图像处理项目里，因为没注意分支发散，导致性能直接腰斩。后来我把条件判断改成了用掩码计算，Warp里的线程就全跑满了。避坑指南：尽量让Warp内的线程走相同的控制流。

个人经验：写kernel时，尽量让线程ID与数据访问模式对齐。比如用threadIdx.x作为数组索引，这样Warp内的线程访问连续地址，能最大化利用显存带宽。

好，现在从硬件视角切换到软件视角。你在写CUDA程序时，会定义两个层级：Grid和Block。

硬件怎么映射呢？一个Block会被分配到同一个SM（Streaming Multiprocessor）上。SM里的Warp调度器，再把Block里的线程拆成多个Warp来执行。

我记得有一次调试一个矩阵乘法的kernel，发现性能上不去。后来一查，是Block大小设得太小，导致SM里的Warp数量不足，无法隐藏访存延迟。调整Block大小后，吞吐量直接翻倍。

注意：Block内的线程可以通过共享内存（Shared Memory）通信，但Block之间是独立的，无法直接同步。跨Block的同步只能通过kernel启动边界来实现。

你想想看，为什么要有Grid和Block两层？直接一个大的线程集合不行吗？原因在于硬件资源有限。一个SM能同时容纳的线程数、寄存器数、共享内存大小都是固定的。通过Block这个中间层，我们可以灵活控制每个SM上的负载。

下面我用一张SVG图，把这几层的关系串起来。你一看就明白了。

这张图从顶向下展示了：Grid包含多个Block，每个Block包含多个Warp，每个Warp包含32个线程。硬件调度以Warp为单位，但编程时我们以Block和Grid来组织逻辑。

最后，我总结几个实际项目中容易踩的坑：

Block大小不是越大越好：每个SM能同时容纳的Warp数量有限。Block太大，会导致SM上活跃的Block数量减少，影响延迟隐藏。
Warp发散要警惕：如果Warp内32个线程走了不同的分支，性能会下降。我建议用__ballot_sync或__shfl_sync这类指令来优化。
网格维度要匹配数据：Grid和Block的维度（1D/2D/3D）最好与你的数据布局一致。比如处理2D图像，就用dim3(blockDim.x, blockDim.y)。

一个小技巧：调试时可以用cudaOccupancyMaxPotentialBlockSize这个API，它会帮你算出当前硬件上最优的Block大小。省得自己瞎猜。

好了，SIMT、Warp、Block、Grid这几个概念，今天就聊到这儿。它们就像GPU世界的「语法规则」，搞懂了，后面学内存模型、调度器、缓存架构都会轻松很多。

公众号:蓝海资料掘金营，微信deep3321