GPU的并行结构#

• SM（流式多处理器） 是GPU的基本计算单元，每个SM可同时执行数百个线程
• 一个GPU通常包含多个SM，因此整体可同时执行数千个线程
• 内核启动后，线程块被分配到可用的SM上执行
• 线程块一旦绑定到某个SM，其线程只能在该SM上运行
• 多个线程块可以共享同一个SM，具体取决于该SM的剩余资源

线程束（Warp）#

• CUDA采用 SIMT（单指令多线程） 架构
• 每 32个线程 组成一个线程束（warp）
• 同一线程束内的所有线程同步执行相同指令
• 每个线程拥有独立的：
  • 指令地址计数器
  • 寄存器状态
  • 执行路径（允许分支发散）

SIMT vs SIMD#

为什么32是关键数字#

• 32是SM以SIMD方式同时处理的工作粒度
• 线程数量按32对齐时，硬件资源利用率最高
• 非32倍数会造成线程束内部分线程空闲，降低效率

SIMT的三个核心特征（SIMD不具备）#

1. 每个线程有独立的指令地址计数器
2. 每个线程有独立的寄存器状态
3. 每个线程可以有独立的执行路径

为什么需要 Warp？#

Block是软件层面的组织单位，Warp是硬件层面的执行单位。两者不在同一个维度。

1. 寄存器资源限制#

每个线程需要独立的寄存器状态。假设每个线程用32个寄存器：

Warp太大 → 同时活跃的Warp数量锐减 → 无法通过Warp切换隐藏内存延迟

GPU隐藏延迟的核心机制就是快速切换Warp：当一个Warp等待内存数据时，SM立刻换另一个Warp执行。如果SM上只有2个Warp，一个等内存时另一个也在等，SM就空转了。

2. 分支发散（Divergence）的代价#

SIMT架构下，同Warp内线程走不同分支时必须串行执行：

1
if (threadIdx.x % 2 == 0) {
2
    A();  // 偶数线程执行
3
} else {
4
    B();  // 奇数线程执行
5
}

• Warp大小32：一半线程等，一半执行，浪费50%执行周期
• Warp大小1024：极端情况下1023个线程等1个线程，浪费99.9%执行周期

Warp越大，分支发散的惩罚越重，并行效率越低。

3. 同步开销#

Warp内所有线程共享指令流，但Warp之间需要显式同步。Warp越大：

• 同步屏障涉及线程越多，等待最慢线程的时间越长
• 共享内存竞争概率增加

4. 硬件复杂度和时钟频率#

• 32路并行：控制逻辑相对简单，布线短，时钟频率可以做高
• 1024路并行：需要巨大的指令广播网络和结果收集网络，面积、功耗、延迟都会爆炸，时钟频率被迫降低

GPU追求高吞吐量，靠的是大量小单元高频运行，而不是少量大单元低频运行。

5. 内存访问合并（Coalescing）#

Warp内线程的内存访问模式被硬件优化为合并成一次事务。32个线程的地址通常容易对齐和合并；1024个线程的地址范围太大，合并效率反而下降，内存带宽利用率降低。

Fermi#

Kepler#