第三章：GPU核心概念速览——从CPU到GPU的思维转变

好，我们正式开始接触GPU的核心概念。

说实话，很多做CPU出身的朋友，第一次看GPU架构都会觉得别扭。我也是这样。当年我在做并行计算加速项目时，拿着CPU的多线程思维去写CUDA代码，结果性能惨不忍睹。后来我才明白——CPU和GPU的思维方式，从根本上就是两码事。

3.1 CPU vs GPU：两种哲学

先看一张对比表，这是我个人习惯的总结方式：

对比维度	CPU	GPU
设计目标	低延迟，单线程快	高吞吐，大量线程并行
核心数量	几个到几十个	几百到几千个
缓存策略	大缓存，复杂分支预测	小缓存，简单控制逻辑
适合任务	串行、分支密集	数据并行、计算密集
线程切换代价	高（上下文切换）	极低（硬件调度）

你想想看，CPU的每个核心都像个全能选手——能跑操作系统、能处理中断、能做复杂分支预测。但GPU不一样，它把大量晶体管都堆在了计算单元上。说白了，GPU就是用更多的简单核心，去碾压那些可以并行的计算任务。

3.2 SIMT模型初探

GPU最核心的执行模型，叫SIMT（Single Instruction, Multiple Threads）。

嗯，这个名字有点绕。我换个说法：一条指令，驱动一堆线程。

举个例子。假设你要给一万个像素做颜色变换：

// CPU做法：循环遍历
for (int i = 0; i < 10000; i++) {
    pixels[i] = transform(pixels[i]);
}

// GPU做法：每个线程处理一个像素
__global__ void transform_kernel(Pixel* pixels) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    pixels[idx] = transform(pixels[idx]);
}

CPU是一个接一个地处理。GPU呢？它让所有线程同时执行transform这个函数。但注意——所有线程执行的是同一段代码，只是处理的数据不同。这就是SIMT的精髓。

3.3 线程与线程束（Warp）

好，接下来是GPU里最重要的概念之一——线程束（Warp）。

你可能会想：GPU有几千个线程，难道CPU一个一个去调度？当然不是。GPU把线程分组管理，这个组就叫线程束。在NVIDIA的架构里，一个线程束固定包含32个线程。

为什么是32？这跟硬件设计有关。我记得当年看Fermi架构的白皮书时，里面提到32这个数字是经过大量权衡的——太小了硬件利用率不够，太大了分支发散代价太高。

线程束的执行规则很简单：

• 同一线程束内的线程，执行同一条指令
• 如果线程出现分支（if-else），部分线程会被"屏蔽"
• 线程束是硬件调度的基本单位

举个例子：

__global__ void kernel(float* data, int n) {
    int idx = threadIdx.x;
    if (idx < n) {
        data[idx] = data[idx] * 2.0f;
    } else {
        data[idx] = 0.0f;
    }
}

假设线程束里有32个线程，但只有前20个满足idx < n。那么这32个线程会先执行if分支——但只有前20个真正干活，后12个被"屏蔽"。然后反过来，后12个执行else分支，前20个被屏蔽。

3.4 从线程到线程束的映射

你可能会问：我写的代码明明是操作单个线程，硬件怎么知道哪些线程属于同一个线程束？

答案是：按线程ID连续分组。在CUDA里，threadIdx.x从0开始，0-31是第一个线程束，32-63是第二个，以此类推。

所以，当你写这样的代码时：

int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid % 2 == 0) {
    // 偶数线程
} else {
    // 奇数线程
}

嗯，这就有问题了。因为在一个线程束里，0和1是相邻的线程，但它们走了不同的分支。这就是典型的线程束内分支发散。

3.5 思维转变总结

从CPU到GPU，你需要记住三件事：

1. 别想着优化单线程性能——GPU的优势在于同时跑成千上万个线程
2. 小心分支发散——同一线程束内的线程，尽量走同一执行路径
3. 理解线程束是基本调度单位——你的代码最终是以32个线程为一组执行的

我个人觉得，最难转变的就是第二条。很多从CPU转过来的开发者，习惯了用if-else做精细控制，到了GPU上反而成了性能杀手。嗯，这需要一点时间适应。

下一章，我们会深入GPU的内存层次结构。你会发现，内存访问模式对性能的影响，比计算本身还要大。到时候我会分享一个我踩过的坑——因为没注意内存合并访问，导致带宽利用率不到10%。