# CUDA 基础：线程束执行的本质

# 1. 线程束和线程块

线程束是 SM 中基本的执行单元，当一个线程块的网格被启动后，网格中的线程块分布在 SM 中。一旦线程块被调度在一个 SM 上，线程块中的线程会被进一步划分为线程束。一个线程束由 32 个连续的线程组成（目前的 GPU 都是 32 个线程，但不保证未来是 32 个），在一个线程束中，所有的线程按照单指令多线程（SIMT）方式执行；也就是说，所有线程都执行相同的指令，每个线程在私有数据上进行操作。下图展示了线程块的逻辑视图和硬件视图之间的关系：

然而，从硬件的角度来看，所有的线程都被组织成了一维的，线程块可以被配置为一维、二维、三维的。在一个块中，每个线程都有唯一的 ID 。对于一维的线程块，唯一的线程 ID 被存储在 CUDA 的内置变量 threadIdx.x 中，并且，threadIdx.x 中拥有连续值得线程被分组到线程束中。例如，一个有 128 个线程的一维线程块被组织到 4 个线程束里，如下所示：

warp0: thread 0, .........., thread 31
warp1: thread 32, ........., thread 63
warp2: thread 64, ........., thread 95
warp3: thread 96, ........., thread 127

线程块是一个逻辑产物，因为在计算机里，内存总是一维线性存在的，所以执行起来也是一维的访问线程块中的线程，但是我们在写程序的时候却可以以二维三维的方式进行，原因是方便我们写程序，比如处理图像或者三维的数据，三维块就会变得很直接，很方便。

在块中，每个线程有唯一的编号（可能是个三维的编号），threadIdx
网格中，每个线程块也有唯一的编号(可能是个三维的编号)，blockIdx
那么每个线程就有在网格中的唯一编号。

用 $x$ 维度作为最内层的维度， $y$ 维度作为第二个维度， $z$ 维度作为最外层的维度，则二维或三维线程块的逻辑布局可以转化为一维物理布局。例如，对于一个给定的二维线程块，在一个块中每个线程的独特标识符都可以用内置变量 threadIdx 和 blockDim 来计算：

tid = threadIdx.x + threadIdx.y * blockDim.x;

对于一个三维线程块，可以用下面的方式计算：

tid = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.x * blockDim.y;

在 C 语言中，假设三维数组 t 保存了所有的线程，那么 (threadIdx.x, threadIdx.y, threadIdx.z) 就相当于：

t[z][y][x];

一个线程块的线程束的数量可以根据下式确定：

$$ \mathrm{WarpsPerBlock} = \left \lceil \frac{\mathrm{threadsPerBlock}}{\mathrm{warpSize}} \right \rceil $$

因此，硬件总是给一个线程块分配一定数量的线程束。线程束不会在不同的线程块之间分离。如果线程块的大小不是线程束大小的偶数倍，那么在最后的线程束里有些线程就不会活跃。比如说一个在 $x$ 轴中有 40 个线程、在 $y$ 轴中有 2 个线程的二维线程块。从应用程序的角度来看，在一个二维网格中共有 80 个线程。

硬件为这个线程块配置了 3 个线程束，使总共 96 个硬件线程去支持 80 个软件线程。注意，最后半个线程束是不活跃的。即使这些线程未被使用，它们仍然消耗 SM 的资源，如寄存器。

注释
从逻辑角度来看： 线程块是线程的集合，它们可以组织为一维、二维或三维布局。
从硬件角度来看： 线程块是一维线程束的集合。在线程块中线程被组织成一维布局，每 32 个连续线程组成一个线程束。

# 2. 线程束分化

控制流是高级编程语言的基本构造中的一种。GPU 支持传统的、C 风格的、显式的控制流结构，例如，if···then···else、for 和 while。

CPU 拥有复杂的硬件以执行分支预测，也就是在每个条件检查中预测应用程序的控制流会使用哪个分支。如果预测正确，CPU 中的分支只需付出很小的性能代价。如果预测不正确，CPU 可能会停止运行很多个周期，因为指令流水线被清空了。我们不必完全理解为什么 CPU 擅长处理复杂的控制流。这个解释只是作为对比的背景。当我们的程序包含大量的分支判断时，从程序角度来说，程序的逻辑是很复杂的，因为一个分支就会有两条路可以走，如果有 10 个分支，那么一共有 1024 条路走，CPU 采用流水线化作业，如果每次等到分支执行完再执行下面的指令会造成很大的延迟，所以现在处理器都采用分支预测技术，而 CPU 的这项技术相对于 GPU 来说高级了不止一点点，而这也是 GPU 与 CPU 的不同，设计初衷就是为了解决不同的问题。CPU 适合逻辑复杂计算量不大的程序，比如操作系统，控制系统，GPU 适合大量计算简单逻辑的任务，所以被用来算数。

GPU 是相对简单的设备，它没有复杂的分支预测机制。一个线程束中的所有线程在同周期中必须执行相同的指令，如果一个线程执行一条指令，那么线程束中的所有线程都必须执行该指令。如果在同一线程束中的线程使用不同的路径通过同一个应用程序，这可能会产生问题。例如，思考下面的语句:

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

假设在一个线程束中有 16 个线程执行这段代码，cond 为 true，但对于其他 16 个来说 cond 为 false 。一半的线程束需要执行 if 语句块中的指令，而另一半需要执行 else 语句块中的指令。在同一线程束中的线程执行不同的指令，被称为线程束分化。我们已经知道，在一个线程束中所有线程在每个周期中必须执行相同的指令，所以线程束分化似乎会产生一个悖论。

如果一个线程束中的线程产生分化，线程束将连续执行每一个分支路径，而禁用不执行这一路径的线程。线程束分化会导致性能明显地下降。在前面的例子中可以看到，线程中并行线程的数量减少了一半: 只有 16 个线程同时活跃地执行，而其他 16 个被禁用了。条件分支越多，并行性削弱越严重。此外，线程束分化只发生在同一个线程束中。在不同的线程束中，不同的条件值不会引起线程束分化。

执行过程如下：

因为线程束分化导致的性能下降就应该用线程束的方法解决，根本思路是避免同一个线程束内的线程分化，而让我们能控制线程束内线程行为的原因是线程块中线程分配到线程束是有规律的而不是随机的。这就使得我们根据线程编号来设计分支是可以的，补充说明下，当一个线程束中所有的线程都执行 if 或者，都执行 else 时，不存在性能下降；只有当线程束内有分歧产生分支的时候，性能才会急剧下降。

线程束内的线程是可以被我们控制的，那么我们就把都执行 if 的线程塞到一个线程束中，或者让一个线程束中的线程都执行 if ，另外线程都执行 else 的这种方式可以将效率提高很多。

举以下一个低效的核函数为例：

__global__ void mathKernel1(float *c) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    float a, b;
    a = b = 0.0f;
    if (tid % 2 == 0) {
        a = 100.0f;
    } else {
        b = 200.0f;
    }
    c[tid] = a + b;
}

这种情况下，线程束内的线程会产生分化，因为线程束内的线程会有一半执行 if ，另一半执行 else ，这样就会导致性能下降。我们可以通过下面的方式来优化：

__global__ void mathKernel2(float *c) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    float a, b;
    a = b = 0.0f;
    if ((tid / warpSize) % 2 == 0) {
        a = 100.0f;
    } else {
        b = 200.0f;
    }
    c[tid] = a + b;
}

假设只配置一个大小为 64 的一维线程块，那么只有 2 个线程束，第一个线程束内的线程编号 tid 从 0 到 31， tid / warpSize 都等于 0，那么都执行 if 语句；第二个线程束内的线程编号 tid 从 32 到 63， tid / warpSize 都等于 1，那么都执行 else 语句。这样就避免了线程束内的线程分化，效率较高。

在 CUDA 中，对线程束分化的评价指标为分支效率 (branch efficiency)，它是一个 0 到 100 之间的百分比，表示线程束中的线程在同一周期中执行的分支指令的百分比。分支效率越高，性能越好。分支效率的计算公式如下：

$$ \mathrm{branch\ efficiency} = \frac{\mathrm{branches - divergent\ branches}}{\mathrm{branches}} $$

以上线程束分化例子的完整代码如下：

PKUcoldkeyboard

cuda-demo

Waiting for api.github.com...

unkown

Waiting...

编译命令为：（强制 CUDA 编译器不利用分支预测去优化内核，使用 Tesla T4 GPU）

nvcc -g -G -arch=sm_75 -o simpleDivergence simpleDivergence.cu

运行结果为：

代码中的 Warmup 部分是提前启动一次 GPU，因为第一次启动 GPU 时会比第二次速度慢一些，具体原因未知，可以去查一下 CUDA 的相关技术文档了解内容。我们可以通过 Nvidia Nsight Compute 来查看分支效率（旧版的 nvprof 被弃用了，metrics 参数对应的修改可以参考 CUDA 编程性能分析工具 nvprof/ncu –metrics 参数含义，而且运行 ncu 时候必须使用 root 权限），结果如下所示：

[58735] simpleDivergence@127.0.0.1
  warmingup(float *) (1, 1, 1)x(64, 1, 1), Context 1, Stream 7, Device 0, CC 7.5
    Section: Command line profiler metrics
    ---------------------------------------------------- ----------- ------------
    Metric Name                                          Metric Unit Metric Value
    ---------------------------------------------------- ----------- ------------
    smsp_sass_average_branch_targets_threads_uniform.pct                  100.00%
    ---------------------------------------------------- ----------- ------------

  mathKernel1(float *) (1, 1, 1)x(64, 1, 1), Context 1, Stream 7, Device 0, CC 7.5
    Section: Command line profiler metrics
    ---------------------------------------------------- ----------- ------------
    Metric Name                                          Metric Unit Metric Value
    ---------------------------------------------------- ----------- ------------
    smsp_sass_average_branch_targets_threads_uniform.pct                   83.33%
    ---------------------------------------------------- ----------- ------------

  mathKernel2(float *) (1, 1, 1)x(64, 1, 1), Context 1, Stream 7, Device 0, CC 7.5
    Section: Command line profiler metrics
    ---------------------------------------------------- ----------- ------------
    Metric Name                                          Metric Unit Metric Value
    ---------------------------------------------------- ----------- ------------
    smsp_sass_average_branch_targets_threads_uniform.pct                  100.00%
    ---------------------------------------------------- ----------- ------------

  mathKernel3(float *) (1, 1, 1)x(64, 1, 1), Context 1, Stream 7, Device 0, CC 7.5
    Section: Command line profiler metrics
    ---------------------------------------------------- ----------- ------------
    Metric Name                                          Metric Unit Metric Value
    ---------------------------------------------------- ----------- ------------
    smsp_sass_average_branch_targets_threads_uniform.pct                   71.43%
    ---------------------------------------------------- ----------- ------------

  mathKernel4(float *) (1, 1, 1)x(64, 1, 1), Context 1, Stream 7, Device 0, CC 7.5
    Section: Command line profiler metrics
    ---------------------------------------------------- ----------- ------------
    Metric Name                                          Metric Unit Metric Value
    ---------------------------------------------------- ----------- ------------
    smsp_sass_average_branch_targets_threads_uniform.pct                  100.00%
    ---------------------------------------------------- ----------- ------------

CUDA 的 nvcc 编译器仍然是在 mathKernel1 和 mathKernel3 上执行有限的优化，以保证分支效率在 50% 以上。注意，mathKernel2 不报告分支分化的唯一原因是它的分支粒度是线程束大小的倍数。此外，把 mathKernel1 中的 if…else 语句分离为 mathKernel3 的多个 if 语句，可以使分支分化的数量翻倍。

# 3. 资源分配

前面提到，每个 SM 上执行的基本单位是线程束，也就是说，单指令通过指令调度器广播给某线程束的全部线程，这些线程同一时刻执行同一命令，当然也有分支情况，也有很多线程束没执行，那么这些没执行的线程束情况又如何呢？可以将这些没执行的线程束分为两类：一类是已经激活的，也就是说这类线程束其实已经在 SM 上准备就绪了，只是没轮到它执行，这时候它的状态为阻塞，另一类是可能分配到 SM 了，但是还没上片，这类就称之为未激活线程束。而每个 SM 上有多少个线程束处于激活状态，取决于以下资源：

程序计数器
寄存器
共享内存

线程束一旦被激活来到片上，那么它就不会再离开 SM 直到执行结束。

每个 SM 都有 32 位的寄存器组，每个架构寄存器的数量不一样，其存储于寄存器文件中，为每个线程进行分配，同时，固定数量的共享内存，在线程块之间分配。

一个 SM 上被分配多少个线程块和线程束取决于 SM 中可用的寄存器和共享内存，以及内核需要的寄存器和共享内存大小。当 kernel 占用的资源较少，那么更多的线程处于活跃状态，相反则线程越少。

寄存器资源的分配

共享内存的分配

上面讲的主要是线程束，如果从逻辑上来看线程块的话，可用资源的分配也会影响常驻线程块的数量。特别是当 SM 内的资源没办法处理一个完整块，那么程序将无法启动。

以下是资源列表：

当寄存器和共享内存分配给了线程块，这个线程块处于活跃状态，所包含的线程束称为活跃线程束。活跃的线程束又分为三类：

选定的线程束
阻塞的线程束
符合条件的线程束

当 SM 要执行某个线程束的时候，执行的这个线程束叫做选定的线程束，准备要执行的叫符合条件的线程束，如果线程束不符合条件还没准备好就是阻塞的线程束。满足下面的要求，线程束才算是符合条件的：

32 个 CUDA 核心可以用于执行
执行所需要的资源全部就位

Kepler 活跃的线程束数量从开始到结束不得大于 64，可以等于。任何周期选定的线程束小于等于 4 。由于计算资源是在线程束之间分配的，且线程束的整个生命周期都在片上，所以线程束的上下文切换是非常快速的。下一节将说明如何通过大量的活跃的线程束切换来隐藏延迟。

# 4. 延迟隐藏

SM 依赖线程级并行，以最大化功能单元的利用率，因此，利用率与常驻线程束的数量直接相关。在指令发出和完成之间的时钟周期被定义为指令延迟。当每个时钟周期中所有的线程调度器都有一个符合条件的线程束时，可以达到计算资源的完全利用。这就可以保证，通过在其他常驻线程束中发布其他指令，可以隐藏每个指令的延迟。

与在 CPU 上用 C 语言编程相比，延迟隐藏在 CUDA 编程中尤为重要。CPU 核心是为同时最小化延迟一个或两个线程而设计的，而 GPU 则是为处理大量并发和轻量级线程以最大化吞吐量而设计的。GPU 的指令延迟被其他线程束的计算隐藏。

考虑到指令延迟，指令可以被分为两种基本类型：

算术指令
内存指令

算术指令延迟是一个算术操作从开始到它产生输出之间的时间。内存指令延迟是指发送出的加载或存储操作和数据到达目的地之间的时间。对于每种情况，相应的延迟大约为：

算术操作为 10～20 个周期
全局内存访问为 400～800 个周期

下图是阻塞线程束到可选线程束的过程逻辑图：

其中线程束 0 （Warp 0）阻塞两段时间后恢复可选模式，但是在这段等待时间中，SM 没有闲置。那么至少需要多少线程，线程束来保证最小化延迟呢？可以根据利特尔法则（Little’s Law）提供一个合理的近似值。它起源于队列理论中的一个定理，也可以用于 GPU 中：

$$ \mathrm{所需线程束}=\mathrm{延迟}\times \mathrm{吞吐量} $$

注释
注意带宽和吞吐量的区别，带宽一般指的是理论峰值，最大每个时钟周期能执行多少个指令，吞吐量是指实际操作过程中每分钟处理多少个指令。简单来说，带宽通常是指理论峰值，而吞吐量是指已达到的值。

这个可以想象成一个瀑布，像这样，绿箭头是线程束，只要线程束足够多，吞吐量是不会降低的：

假设在 kernel 里一条指令的平均延迟是 5 个周期。为了保持在每个周期内执行 6 个线程束的吞吐量，则至少需要 30 个未完成的线程束。

对于算术运算来说，其所需的并行可以表示成隐藏算术延迟所需要的操作数量。下面的表格出了 Fermi 和 Kepler 设备所需的操作数量。示例中的算术运算是一个 32 位的浮点数乘加运算 (a + b $\times$ c)，表示在每个 SM 中每个时钟周期内的操作数量。吞吐量因不同的算术指令而不同。

吞吐量由 SM 中每个周期内的操作数量确定，而执行一条指令的一个线程束对应 32 个操作。因此，为保持计算资源的充分利用，对于 Fermi GPU 而言，每个 SM 中所需的线程束数量通过计算为 $640 \div 32 = 20 $ 个线程束。因此，算术运算所需的并行可以用操作的数量或线程束的数量来表示。这个简单的单位转换表明，有两种方法可以提高并行：

指令级并行（ILP）：一个线程中有很多独立的指令
线程级并行（TLP）：很多并发地符合条件的线程

同样，与指令周期隐藏延迟类似，内存隐藏延迟是靠内存读取的并发操作来完成的，需要注意的是，指令隐藏的关键目的是使用全部的计算资源，而内存读取的延迟隐藏是为了使用全部的内存带宽，内存延迟的时候，计算资源正在被别的线程束使用，所以我们不考虑内存读取延迟的时候计算资源在做了什么，我们的根本目的是把计算资源，内存读取的带宽资源全部使用满，这样就能达到理论的最大效率。同样下表根据利特尔法则给出了需要多少线程束来最小化内存读取延迟，不过这里有个单位换算过程，机器的性能指标内存读取速度给出的是 GB/s 的单位，而我们需要的是每个时钟周期读取字节数，所以要用这个速度除以频率，例如 Tesla C2070 的内存带宽是 144 GB/s，转化成时钟周期： $\frac{144\mathrm{GB/s}}{1.566 \mathrm{GHz}}=92\mathrm{B/t}$，这样就能得到单位时间周期的内存带宽了。即下表的数据：

需要说明的是这个速度不是单个 SM 的而是整个 GPU 设备的。Fermi 需要并行的读取 74KB 的数据才能让 GPU 带宽满载，如果每个线程读取 4 个字节，我们大约需要 18500 个线程，大约 579 个线程束才能达到这个峰值。

所以，延迟的隐藏取决于活动的线程束的数量，数量越多，隐藏得越好，但是线程束的数量又受到上面的说的资源影响。所以这里就需要寻找最优的执行配置来达到最优的延迟隐藏。

那么我们怎么样确定一个线程束的下界呢，使得当高于这个数字时 SM 的延迟能充分的隐藏，其实这个公式很简单，也很好理解，就是 SM 的计算核心数乘以单条指令的延迟，比如 32 个单精度浮点计算器，每次计算延迟 20 个时钟周期，那么我需要最少 $32 \times 20 =640$ 个线程使设备处于忙碌状态。然而，这只是一个下边界。

# 5. 占用率

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

$$ \mathrm{Occupancy} = \frac{\mathrm{Active\ Warps}}{\mathrm{Max\ Warps}} $$

通过以下代码可以查询设备的最大线程束数量：

int dev = 0;
cudaDeviceProp deviceProp;
CHECK(cudaGetDeviceProperties(&deviceProp, dev));

log_info("Device %d: %s", dev, deviceProp.name);
log_info("Number of SMs: %d", deviceProp.multiProcessorCount);
log_info("Total amount of constant memory: %4.2f KB", deviceProp.totalConstMem / 1024.0);
log_info("Total amount of shared memory per block: %4.2f KB",
          deviceProp.sharedMemPerBlock / 1024.0);
log_info("Total number of registers available per block: %d", deviceProp.regsPerBlock);
log_info("Warp size: %d", deviceProp.warpSize);
log_info("Maximum number of threads per block: %d", deviceProp.maxThreadsPerBlock);
log_info("Maximum number of threads per multiprocessor: %d",
          deviceProp.maxThreadsPerMultiProcessor);
log_info("Maximum number of warps per multiprocessor: %d",
          deviceProp.maxThreadsPerMultiProcessor / 32);

return 0;

输出结果为：

可以看到 RTX4090 最大 64 个线程束每个 SM。

内核使用寄存器的数量会影响 SM 内线程束的数量，nvcc 的编译选项也有手动控制寄存器的使用。也可以通过调整线程块内线程的多少来提高占用率，当然要合理不能太极端：

小的线程块：每个线程块中线程太少，会在所有资源没用完就达到了线程束的最大要求
大的线程块：每个线程块中太多线程，会导致每个 SM 中每个线程可用的硬件资源较少。

一个确定网格和线程块大小的基本准则如下：

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

尽管在每种情况下会遇到不同的硬件限制，但它们都会导致计算资源未被充分利用，阻碍隐藏指令和内存延迟的并行的建立。占用率唯一注重的是在每个 SM 中并发线程或线程束的数量。然而，充分的占用率不是性能优化的唯一目标。内核一旦达到一定级别的占用率，进一步增加占用率可能不会改进性能。为了提高性能，可以调整很多其他因素。

# 6. 同步

栅栏同步是一个原语，它在许多并行编程语言中都很常见。在 CUDA 中，同步可以在两个级别执行：

线程块内同步
系统级别

块级别的就是同一个块内的线程会同时停止在某个设定的位置，用

__syncthreads();

这个函数完成，这个函数只能同步同一个块内的线程，不能同步不同块内的线程，想要同步不同块内的线程，就只能让核函数执行完成，控制程序交换主机，这种方式来同步所有线程。当__syncthreads 被调用时，在同一个线程块中每个线程都必须等待直至该线程块中所有其他线程都已经达到这个同步点。线程产生的所有全局内存和共享内存访问，将会在栅栏后对线程块中所有其他的线程可见。该函数可以协调同一个块中线程之间的通信，但它强制线程束空闲，从而可能对性能产生负面影响。

在不同的块之间没有线程同步。块间同步，唯一安全的方法是在每个内核执行结束端使用全局同步点；也就是说，在全局同步之后，终止当前的核函数，开始执行新的核函数。

不同块中的线程不允许相互同步，因此 GPU 可以以任意顺序执行块。这使得 CUDA 程序在大规模并行 GPU 上是可扩展的。

# 7. 可扩展性

对于任何并行应用程序而言，可扩展性是一个理想的特性。可扩展性意味着为并行应用程序提供了额外的硬件资源，相对于增加的资源，并行应用程序会产生加速。例如，若一个 CUDA 程序在两个 SM 中是可扩展的，则与在一个 SM 中运行相比，在两个 SM 中运行会使运行时间减半。一个可扩展的并行程序可以高效地使用所有的计算资源以提高性能。可扩展性意味着增加的计算核心可以提高性能。串行代码本身是不可扩展的，因为在成千上万的内核上运行一个串行单线程应用程序，对性能是没有影响的。并行代码有可扩展的潜能，但真正的可扩展性取决于算法设计和硬件特性。

能够在可变数量的计算核心上执行相同的应用程序代码的能力被称为透明可扩展性。一个透明的可扩展平台拓宽了现有应用程序的应用范围，并减少了开发人员的负担，因为它们可以避免新的或不同的硬件产生的变化。可扩展性比效率更重要。一个可扩展但效率很低的系统可以通过简单添加硬件核心来处理更大的工作负载。一个效率很高但不可扩展的系统可能很快会达到可实现性能的上限。

CUDA 内核启动时，线程块分布在多个 SM 中。网格中的线程块以并行或连续或任意的顺序被执行。这种独立性使得 CUDA 程序在任意数量的计算核心间可以扩展。

下图展示了 CUDA 架构可扩展性的一个例子。左侧的 GPU 有两个 SM，可以同时执行两个块；右侧的 GPU 有 4 个 SM ，可以同时执行 4 个块。不修改任何代码，一个应用程序可以在不同的 GPU 配置上运行，并且所需的执行时间根据可用的资源而改变。

# 参考资料

[1] CUDA C 编程权威指南，机械工业出版社，（美）程润伟（John Cheng）等著