2、GPU的核心思想：SIMT（单指令多线程）架构详解，对比CPU的SIMD，理解并行计算的根本差异。

好，我们直接进入正题。

上一章我们聊了GPU和CPU在宏观上的分工。这一章，我们来扒一扒GPU最核心的“灵魂”——SIMT（单指令多线程）。说白了，不理解SIMT，你就不算真正懂GPU。

2.1 从一条指令说起：CPU怎么干活？

先看看CPU。CPU处理数据，最经典的方式叫SIMD（单指令多数据）。

举个例子。你要给一个数组里每个数加1。在CPU上，你可以写一个循环：

for (int i = 0; i < 8; i++) {
    a[i] = a[i] + 1;
}

CPU会一条一条地执行。但如果你用SIMD指令集（比如AVX），你可以这样：

// 伪代码：一次加载8个整数，一次加1，一次存回
__m256i vec = _mm256_loadu_si256((__m256i*)a);
vec = _mm256_add_epi32(vec, _mm256_set1_epi32(1));
_mm256_storeu_si256((__m256i*)a, vec);

看到了吗？一条指令，同时操作了8个数据。这就是SIMD的核心：数据并行。你有一个“宽”的运算单元，一次塞进去多个数据，它一次性算完。

但这里有个关键点：CPU的SIMD是显式的。程序员必须手动把数据打包，调用特定的指令。编译器有时候能帮你自动向量化，但很多时候不行。我见过不少项目，为了榨干CPU的最后一点性能，工程师花大量时间手写SIMD intrinsics，调试起来非常痛苦。

核心区别一： CPU的SIMD是“数据级并行”，GPU的SIMT是“线程级并行”。

2.2 GPU的SIMT：让线程来“伪装”数据

GPU的思路完全不同。它不搞什么宽向量指令。它搞的是大量轻量级线程。

还是那个加1的例子。在GPU上（比如CUDA），你写的是：

__global__ void add_one(int *a) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    a[idx] = a[idx] + 1;
}

然后你启动N个线程：

add_one<<<1024, 256>>>(a); // 启动 1024*256 = 262144 个线程

每个线程只处理一个数据。看起来是“多指令多数据”对吧？每个线程执行自己的指令？

错。

GPU硬件在执行时，会把32个线程打包成一个Warp（线程束）。这32个线程，在同一个时刻，执行同一条指令。

这就是SIMT的精髓：从程序员角度看，是多个独立的线程；从硬件执行角度看，是单条指令在操作多个数据。

我个人习惯把Warp理解为GPU的“最小执行单元”。CPU的最小执行单元是“线程”，GPU的最小执行单元是“Warp”。一个Warp里的线程，必须步调一致。

2.3 SIMT vs SIMD：根本差异在哪里？

我们来做个对比表格，这样更清楚：

对比维度 CPU SIMD GPU SIMT
并行粒度 数据级（向量元素） 线程级（Warp内的线程）
编程模型 显式向量化（intrinsics/自动向量化） 隐式并行（写普通线程代码）
控制流 无分支（向量内必须一致） 支持分支（但会序列化，有性能代价）
硬件实现 宽ALU（如256位、512位） 大量简单ALU + 共享控制单元
数据访问 需要手动对齐和打包 自动合并访问（coalescing）
典型延迟隐藏 依赖编译器调度 靠线程切换（zero-overhead context switch）

你看，最核心的差异在于控制流处理。

CPU的SIMD，向量里的所有元素必须走同一条路。如果你写：

if (a[i] > 0) {
    b[i] = b[i] * 2;
} else {
    b[i] = b[i] + 1;
}

用SIMD，你得把两个分支都算一遍，然后用mask（掩码）选结果。非常麻烦。

GPU的SIMT呢？一个Warp里的32个线程，如果有的走if，有的走else，硬件会怎么做？

它会先执行所有走if的线程，把走else的线程“屏蔽”掉。然后再反过来执行else的线程。 这叫分支发散（branch divergence）。

我曾经在一个图像处理项目里，因为一个简单的if-else判断，导致性能直接腰斩。排查了半天，才发现是Warp内的线程分支发散太严重。记住：同一个Warp里的线程，尽量走同一个分支。

2.4 为什么GPU能靠“线程切换”隐藏延迟？

CPU遇到缓存未命中，只能干等，几百个周期就浪费了。GPU呢？它不等。

GPU的SM（流多处理器）里，有大量的寄存器文件和线程上下文。当一个Warp因为访存被阻塞时，硬件会立即切换到另一个就绪的Warp。这个切换是零开销的——因为所有上下文都已经在寄存器里了。

你想想看，一个SM里可能有64个Warp（2048个线程）。只要其中有一部分Warp在等待数据，另一部分就能立刻顶上。这就是GPU的“延迟隐藏”魔法。

相比之下，CPU的线程切换代价很高，要保存和恢复一堆寄存器、栈指针、程序计数器。所以CPU通常只跑少量线程，靠大缓存和分支预测来减少延迟。

核心区别二： CPU用“缓存+分支预测”来减少延迟，GPU用“大量线程+零开销切换”来隐藏延迟。

2.5 实际编程中的启示

理解了SIMT，你写GPU代码时就会注意几点：

  • 尽量让Warp内的线程访问连续的内存地址。这样硬件可以把多个访问合并成一次大的内存事务。这就是“合并访问（coalesced access）”。
  • 避免Warp内的分支发散。如果实在避免不了，尽量让分支条件基于Warp ID或线程ID的某些位，这样同一个Warp的线程会走同一条路。
  • 不要害怕启动大量线程。GPU就是为大量线程设计的。启动的线程越多，硬件越有机会通过切换来隐藏延迟。

我建议你在写kernel时，先想想这个Warp会怎么执行。如果32个线程各干各的，那性能大概率好不了。反过来，如果32个线程步调一致，那GPU就能跑得飞起。

2.6 小结

嗯，这一章我们讲清楚了SIMT的核心思想。它和CPU的SIMD，本质上是两种不同的并行哲学：

  • SIMD：把数据压扁，塞进宽ALU里一次性处理。
  • SIMT：把任务拆成大量线程，让硬件用单指令驱动它们并行执行。

下一章，我们会深入GPU的内存层次。你会发现，SIMT的执行模型和GPU的内存设计是紧密耦合的。理解了SIMT，你才能理解为什么GPU有那么多奇奇怪怪的内存类型——全局内存、共享内存、本地内存、纹理内存……

好，今天就到这里。记住：Warp是GPU的灵魂，分支发散是性能的杀手。

← 上一章 📖 返回目录 下一章 →