第2章计算核心架构

2. 计算核心架构：流处理器(SP)与线程束(Warp)调度、SIMT执行模型详解

好，咱们今天聊点硬核的。GPU 里头最核心的东西，就是它的计算核心。你想想看，一个 GPU 动辄几千个核心，这些核心到底是怎么工作的？它们怎么协同？怎么调度？嗯，这就是我们这一章要讲清楚的事。

我个人习惯把 GPU 的计算核心想象成一个「超级工厂」。流处理器（SP）就是工厂里的工人，线程束（Warp）就是工人的班组，而 SIMT 执行模型，就是工厂的管理制度。咱们一个一个拆开看。

2.1 流处理器（SP）：最基础的算力单元

流处理器，英文叫 Streaming Processor，简称 SP。在 NVIDIA 的术语里，它也叫 CUDA Core。说白了，这就是 GPU 里干活的最小单元。

一个 SP 能做什么？它能做整数运算、浮点运算、逻辑运算、地址计算等等。但注意，它不能独立执行指令。它必须依附于更大的调度单元。

核心要点：SP 是执行单元，不是控制单元。它只管算，不管「怎么算」和「什么时候算」。

我在项目中遇到过一个问题：有同事把 SP 的数量直接等同于性能，觉得 SP 越多性能一定越好。其实不然。你想想看，如果调度跟不上，SP 再多也是闲置的。就像工厂里工人再多，如果物料送不过来，也只能干瞪眼。

每个 SP 内部通常包含：

ALU（算术逻辑单元）：做加减乘除、逻辑判断
FPU（浮点运算单元）：处理浮点数计算
寄存器文件：存放临时数据
本地数据缓存：减少访存延迟

一个典型的 GPU 芯片里，SP 会被组织成更大的组。比如 NVIDIA 的架构里，32 个 SP 组成一个 SM（Streaming Multiprocessor）。这个 SM 才是真正的调度单位。

2.2 线程束（Warp）：GPU 的调度原子

好，现在问题来了。GPU 有几千个 SP，如果每个 SP 都独立调度，那调度器的复杂度会爆炸。怎么办？

GPU 的解决方案是：把线程分组。这个组就叫线程束（Warp）。在 NVIDIA 的架构里，一个 Warp 包含 32 个线程。这 32 个线程共享同一个指令，但处理不同的数据。

我的经验：理解 Warp 是理解 GPU 性能的关键。我曾经花了两周时间调一个 kernel，死活跑不满利用率。后来发现是 Warp 调度出了问题——分支发散导致大量线程闲置。嗯，从那以后我再也不敢忽视 Warp 的调度细节了。

Warp 调度的核心机制：

零开销上下文切换：Warp 之间的切换不需要保存/恢复寄存器，因为每个 Warp 有自己的寄存器空间
细粒度多线程：当一个 Warp 在等待内存访问时，调度器立即切换到另一个 Warp
锁步执行：同一个 Warp 内的所有线程执行同一条指令，只是数据不同

你可能会问：为什么是 32 个线程？不是 16 个，也不是 64 个？

这跟硬件设计有关。32 这个数字，刚好能填满一个 SM 里的计算单元，同时不会让调度器太复杂。我记得在早期的 GPU 架构里，Warp 大小是 16 或者 24，后来才统一到 32。这是经过大量权衡的结果。

2.3 SIMT 执行模型：单指令多线程

SIMT，全称 Single Instruction Multiple Threads。这个名字起得很有意思。它跟 SIMD（单指令多数据）很像，但又有本质区别。

SIMD 是：一条指令，操作多个数据。比如 AVX 指令，一次处理 8 个 float。

SIMT 是：一条指令，多个线程各自执行，每个线程有自己的寄存器、自己的地址、自己的状态。

说白了，SIMT 是 SIMD 的「高级版」。它保留了 SIMD 的高效性，又给了程序员更大的灵活性。

关键区别：在 SIMD 里，程序员要手动管理数据布局。在 SIMT 里，程序员写的是标量代码，硬件自动帮你并行化。这就是为什么 CUDA 编程比 SSE/AVX 编程简单得多的原因。

SIMT 的执行流程是这样的：

程序员写一个 kernel 函数，看起来是串行的
硬件创建 N 个线程，每个线程执行同样的 kernel
每 32 个线程组成一个 Warp，共享指令流
Warp 调度器每次选择一个 Warp，发射一条指令
Warp 内的 32 个线程同时执行这条指令（但数据不同）

这里有个坑，我必须提醒你：分支发散。

避坑指南：我曾经写过一个 kernel，里面有个 if-else 分支。当时没在意，结果性能惨不忍睹。后来一分析，发现同一个 Warp 里的线程，有的走 if 分支，有的走 else 分支。硬件没办法同时执行两个分支，只能先执行 if 分支（此时走 else 的线程被屏蔽），再执行 else 分支（此时走 if 的线程被屏蔽）。这样一来，Warp 的执行效率直接砍半。

所以，写 GPU 代码时，要尽量避免 Warp 内的分支发散。如果实在避免不了，也要让同一个 Warp 内的线程走相同的分支。

2.4 调度器的具体实现

咱们来看看调度器到底是怎么工作的。以 NVIDIA 的架构为例，每个 SM 里有多个 Warp 调度器。每个调度器负责管理一组 Warp。

调度器的核心数据结构：

组件	作用	我的备注
Warp 状态表	记录每个 Warp 的状态（就绪、等待、执行中）	这个表是硬件实现的，查表速度极快
指令缓存	存放 Warp 要执行的指令	指令缓存 miss 会导致 Warp 停顿
寄存器文件	每个 Warp 独享的寄存器空间	寄存器分配不当会限制 Warp 数量
Scoreboard	跟踪指令间的依赖关系	类似 CPU 的记分牌，但更简单

调度策略其实很简单：每个周期，调度器从就绪的 Warp 中选一个，发射一条指令。选哪个？通常是轮询（Round-Robin），或者优先选等待时间最长的。

你想想看，如果 SM 里有 64 个 Warp，每个 Warp 执行一条指令需要 4 个周期，那理论上调度器可以隐藏 256 个周期的访存延迟。这就是 GPU 能容忍高延迟的原因——靠的是大量的并行 Warp，而不是靠缓存。

2.5 实战：如何利用 Warp 调度优化性能

光说不练假把式。咱们来看一个实际的优化案例。

假设我们要做一个向量加法：C = A + B。最简单的写法：

__global__ void vec_add(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];
    }
}

这个代码能跑，但性能不是最优的。为什么？因为每个线程只算一个元素，访存模式是连续的，但计算密度太低。

优化思路：让每个线程算多个元素，提高计算密度。

__global__ void vec_add_opt(float *A, float *B, float *C, int N) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    // 每个线程处理 4 个元素
    int stride = blockDim.x * gridDim.x;
    for (int i = idx; i < N / 4; i += stride) {
        float4 a = ((float4*)A)[i];
        float4 b = ((float4*)B)[i];
        ((float4*)C)[i] = make_float4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w);
    }
}

这个优化有什么好处？

减少了线程总数，降低了调度开销
使用了 float4 向量化加载，提高了内存带宽利用率
每个线程做更多计算，提高了 ALU 利用率

小技巧：我在实际项目中，经常用这种「线程合并」的技巧。把多个小任务合并到一个线程里，往往能带来 20%-50% 的性能提升。但要注意，不要合并太多，否则会导致寄存器压力过大，反而降低并行度。

2.6 总结与思考

好，咱们这一章的内容就到这里。来回顾一下核心要点：

SP 是执行单元，只管算，不管调度
Warp 是调度原子，32 个线程一组，共享指令流
SIMT 是执行模型，让程序员写标量代码，硬件自动并行
分支发散 是性能杀手，要尽量避免
Warp 调度 靠大量并行来隐藏延迟，而不是靠缓存

最后，我想留一个问题给你思考：如果让你设计下一代 GPU 的调度器，你会把 Warp 大小改成多少？为什么？

这个问题没有标准答案。但思考的过程，能帮你更深入地理解 GPU 的架构设计哲学。

下一章，咱们聊聊内存层次结构。嗯，那又是一个大坑，到时候我把我踩过的坑都告诉你。

📚 课程目录

2. 计算核心架构：流处理器(SP)与线程束(Warp)调度、SIMT执行模型详解

2.1 流处理器（SP）：最基础的算力单元

2.2 线程束（Warp）：GPU 的调度原子

2.3 SIMT 执行模型：单指令多线程

2.4 调度器的具体实现

2.5 实战：如何利用 Warp 调度优化性能

2.6 总结与思考