🫣线程束和线程束分化|Warp

source from 《Professional CUDA C Programming》By Max Grossman&Ty McKercher.I made some changes according to my own understanding.

线程束是SM中基本的执行单元。当一个线程块的网格被启动后，网格中的线程块分 布在SM中。一旦线程块被调度到一个SM上，线程块中的线程会被进一步划分为线程束。 一个线程束由32个连续的线程组成，在一个线程束中，所有的线程按照单指令多线程 （SIMT）方式执行；也就是说，所有线程都执行相同的指令，每个线程在私有数据上进 行操作。

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

$$\text{一个线程块中线程束的数量}=\text{向正无穷取整}\left(\frac{\text{一个线程块中线程的数量 }}{\text{线程束大小}}\right)$$

因此，硬件总是给一个线程块分配一定数量的线程束。线程束不会在不同的线程块之 间分离。如果线程块的大小不是线程束大小的倍数，那么在最后的线程束里有些线程就 不会活跃。

下图一个在x轴中有40个线程、在y轴中有2个线程的二维线程块。从应用 程序的角度来看，在一个二维网格中共有80个线程。 硬件为这个线程块配置了3个线程束，使总共96个硬件线程去支持80个软件线程。注 意，最后半个线程束是不活跃的。即使这些线程未被使用，它们仍然消耗SM的资源，如寄存器。

线程块：逻辑角度与硬件角度

• 从逻辑角度来看，线程块是线程的集合，它们可以被组织为一维、二维或三维布局。

• 从硬件角度来看，线程块是一维线程束的集合。在线程块中线程被组织成一维布局， 每32个连续线程组成一个线程束。

线程束分化

CPU拥有复杂的硬件以执行分支预测，也就是在每个条件检查中预测应用程序的控制 流会使用哪个分支。如果预测正确，CPU中的分支只需付出很小的性能代价。如果预测不 正确，CPU可能会停止运行很多个周期，因为指令流水线被清空了。我们不必完全理解为 什么CPU擅长处理复杂的控制流。这个解释只是作为对比的背景。

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

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

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

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

注意，线程束分化只发生在同一个线程束中。在不同的线程束中，不同的条件值不会 引起线程束分化。

为了获得最佳的性能，应该避免在同一线程束中有不同的执行路径。例如，假设有两个分支，下面展示了简单的算术内核示例。我们可以用一个偶数和奇 数线程方法来模拟一个简单的数据分区，目的是导致线程束分化。该条件（tid%2==0）使 偶数编号的线程执行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;
  }

如果使用线程束方法（而不是线程方法）来交叉存取数据，可以避免线程束分化，并 且设备的利用率可达到100%。条件（tid/warpSize）%2==0使分支粒度是线程束大小的倍 数；偶数编号的线程执行if子句，奇数编号的线程执行else子句。这个核函数产生相同的 输出，但是顺序不同。

 __global__ void mathKernel2(void) {
   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;
 }

分支效率

分支效率被定义为未分化的分支与全部分支之比，可以使用以下公式来计算：

$$\text{分支效率}=100\times\left(\frac{\text{分支数 - 分化分支数 }}{\text{分支数}}\right)$$

如果我们使用nvprof查看mathKernel1核函数的分支效率，会发现他的分支效率居然是100%。这个奇怪的现象是 CUDA编译器优化导致的结果，它将短的、有条件的代码段的断定指令取代了分支指令 （导致分化的实际控制流指令）。 在分支预测中，根据条件，把每个线程中的一个断定变量设置为1或0。这两种条件流 路径被完全执行，但只有断定为1的指令被执行。断定为0的指令不被执行，但相应的线程 也不会停止。只有在条件语 句的指令数小于某个阈值时，编译器才用断定指令替换分支指令。因此，一段很长的代码 路径肯定会导致线程束分化。

此外，大家也可以做做试验将mathKernel1中的if...else语句修改为多个if语句，可 以使分化分支的数量翻倍。

占用率

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

$$\text{占用率}=\frac{\text{活跃线程束数量}}{\text{最大线程束数量}}$$

小结

• 当一个分化的线程采取不同的代码路径时，会产生线程束分化

• 不同的if-then-else分支会连续执行

• 尝试调整分支粒度以适应线程束大小的倍数，避免线程束分化

• 不同的分化可以执行不同的代码且无须以牺牲性能为代价