CUDA 执行模型

CUDA 执行模型

CUDA 执行模型能够提供有助于在指令吞吐量和内存访问方面编写高效代码的见解。

1 GPU 架构概述

GPU 架构是围绕一个流式多处理器(SM)的拓展阵列搭建的，通过复制这种结构来实现 GPU 的硬件并行。

如图所示，其中包含了 Fermi SM 关键组件：

• CUDA Cores
• Shared Memory / L1 Cache
• Register File
• 加载/存储单元(Load / Store Units)
• 特殊功能单元(Special Function Units)
• 线程束调度器(Warp Scheduler)

SM

GPU 中的每一个 SM 都能支持数百个线程并发执行，每个 GPU 通常有多个 SM。当启动一个核函数时，它的 block 会被分布在可用的 SM 上执行。一个 block 一旦被调度到一个 SM 上之后，其中的 thread 只会在那个 SM 上并发执行。多个 block 可以被分配到同一个 SM 上，这是根据 SM 资源的可用性进行调度的。同一 thread 中的指令利用指令级并行性进行流水线化。

因为 SM 有限，虽然从编程模型层面看所有线程都是并行执行的，但是从更底层的角度看，所有 block 也是分批次的在物理层面的机器上执行，block 里不同的 thread 可能进度都不一样，但是同一个 warp 内的 thread 拥有相同的进度。
并行可能会就会引起竞争，多个 thread 以未定义的顺序访问同一个数据，就导致了不可预测的行为，CUDA 只提供了一种块内同步的方式，块之间没办法同步。同一个 SM 上可以有不止一个常驻的线程束，有些在执行，有些在等待，他们之间状态的转换是不需要开销的。

warp

CUDA 采用单指令多线程(SIMT)架构来管理和执行线程，每 32 个线程一组，称为线程束(warp)。warp 中的所有 thread 同时执行相同的指令，每个 thread 都有自己的地址计数器和寄存器状态，利用自身的数据执行当前的指令。每个 SM 都将分配给它的 block 划分到包含 32 个 thread 的 warp 中，然后再可用的硬件资源上调度执行。

SIMT

SIMT 架构与 SIMD 架构相似，两者都是将相同的指令广播给多个执行单元来实现并行。一个关键的区别是 SIMD 要求同一个向量中的所有元素要在一个统一的同步组中一起执行，而 SIMT 允许属于同一 warp 的多个 thread 独立执行。尽管一个 warp 中的所有 thread 在相同的程序地址上同时开始执行，但是单独的 thread 仍有可能有不同的行为。SIMT 确保可以编写独立的线程级并行代码、标量线程以及用于协调线程的数据并行代码。

2 warp 的本质

2.1 线程束和线程块

$$一个线程块中线程束数量=\lceil \frac{一个线程块中线程的数量}{线程束大小}\rceil$$

2.2 线程束分化

if (cond)
{
    ...
}
else
{
    ...
}

假设这段代码是核函数的一部分，那么当一个 warp 的 32 个 thread 执行这段代码的时候，如果其中 16 个执行 if 中的代码段，而另外 16 个执行 else 中的代码块，同一个 warp 中的 thread，执行不同的指令，这叫做 warp 的分化。
我们知道在每个指令周期，warp 中的所有 thread 执行相同的指令，但是 warp 又是分化的，所以这似乎是相悖的，但是事实上这两个可以不矛盾。解决矛盾的办法就是每个 thread 都执行所有的 if 和 else 部分，当一部分 cond 成立的时候，执行 if 块内的代码，有一部分线程 cond 不成立，那么这些 thread 等待，等到都执行完成 if 之后再执行下一个指令，线程束分化会产生严重的性能下降。条件分支越多，并行性削弱越严重。

注意 warp 分化研究的是一个 warp 中的 thread，不同 warp 中的分支互不影响。

2.3 占用率

在每个 CUDA 核心里指令是顺序执行的，当一个 warp 阻塞时，SM 切换执行其他符合条件的 warp。理想情况下，我们想要有足够的 warp 占用设备的核心。占用率是每个 SM 中活跃的 warp 占最大 warp 数量的比值。
$$占用率=\frac{活跃线程束数量}{最大线程束数量}$$

grid 和 block 大小准则：

• 保持每个块中线程数量是线程束大小(32)的倍数
• 避免 block 太小：每个 block 至少要有 128 或 256 个 thread
• 根据内核资源的需求调整 block 大小
• block 的数量要远远多于 SM 的数量，从而在设备中可以显示有足够的并行
• 通过实验得到最佳执行配置和资源使用情况

2.4 同步

CUDA同步这里只讲两种：

• 块级：在设备执行过程中等待一个 block 中多有 thread 达到同一点
• 系统级：等待主机和设备完成所有的工作

使用下面的代码实现块级同步

__syncthreads();

3 避免分支分化

3.1 并行归约问题

3.2 并行归约中的分化

3.3 改善并行归约的分化

3.4 交错配对的归约