Warp

Warp 是 NVIDIA CUDA 中的一个术语，代表一组同时执行相同指令的线程。一个 Warp 包含一定数量的线程（NVIDIA GPU 上通常是 32 个线程）。Warp 与 AMD GPU 中的 Wavefront 是一个概念。

需要强调的是：Warp 中的线程以一个单一的指令流同时执行相同的指令序列。这种方式，即单指令流多数据流（SIMD）并行性，是建立在所有线程可以在任何时候执行相同操作的假设之上的。

线程发散

分支指令会非常影响 GPU 的性能。

依赖于数据的分支指令（例如if-else语句）可能导致线程之间的执行路径分岔。当一个 Warp 中的线程执行分支指令且分支判断的结果在不同的线程之间有所不同时，GPU 必须处理每个分支路径。这就是所谓的线程发散（Thread Diverge）。

例如，假设一个 Warp 是 4 个线程，Kernel 程序如下：

int i = 0;
int thread_id = (int)(get_global_id(0)); // 获取当前线程的 ID
if (thread_id < 2) { // 当前线程的 ID 如果小于 2，则给 i 变量加一
    i += 1;
}
i -= 1;

第一个 Warp 中的四个线程 ID 分别为：0、1、2和3。2、3 号线程逻辑上不用执行加一操作，但由于 0、1 号线程需要执行且一个 Warp 中的线程永远执行的是相同的指令，因此 2、3 号线程也需要陪跑。这种陪跑的现象可以看成执行过程中的“气泡”，降低了硬件的利用率。

硬件实现

硬件如何处理线程发散。

每个 Warp 设置一个 PC，代表 Warp 中的所有线程当前要执行的指令的 PC。

Warp 中的每个线程再设置一个 NEXT_PC，代表当前线程接下来要跳转到的 PC。在没有分支指令的情况下， NEXT_PC 等于 PC + 1。

假设一个 Warp 包含 4 个线程。

当 4 个线程的 NEXT_PC 不一样时，代表发生了线程发散。线程发散发生之后，选择 NEXT_PC 最小的线程来执行，其他线程陪跑。陪跑线程的 NEXT_PC 在陪跑过程中不变。

退出分支执行时，同样的，选择 NEXT_PC 最小的线程来执行。

举个例子，

__kernel void
test(__global int* OUT) {
  int tid = (int)(get_global_id(0)); // Inst 0
  if (tid < 2) {                     // Inst 1
    OUT[tid] += 10;                  // Inst 2
  }
  OUT[tid] += 100;                   // Inst 3
}

执行过程如下，

图中，

每一行的最前面的 PC 为 Warp 正在执行的指令 PC，紧接着的带方框的 4 个 PC 为四个线程的 NEXT_PC
Warp PC 为 Inst 1 时，发生了分支发散（Diverge）
灰色方框代表此时线程处于陪跑状态
EOS（End of Shader）代表线程结束。

线程合并

我们称退出分支执行并更新线程的执行状态的操作为线程合并（Thread Reconverge）。线程发散只有在线程的实际执行过程中才能发现，因此需要硬件自行判断，但是线程合并一般不让硬件判断，因为编译器可以提前知道。编译器在需要做线程合并的指令做好标记，硬件识别出该标记后，做线程合并操作，即选择 NEXT_PC 最小的线程来执行。

举例：嵌套分支

嵌套分支的处理逻辑同上，为了让读者更好的理解，给出一个更加复杂的情形的执行过程，加深理解。

__kernel void
test(__global int* OUT) {
/*Inst  0*/ int tid = (int)(get_global_id(0));
/*Inst  1*/ if (tid < 2) {
/*Inst  2*/   OUT[tid] += 10;
/*Inst  3*/   if (tid < 1) {
/*Inst  4*/     OUT[tid] -= 10;
/*Inst  -*/   } else {
/*Inst  5*/     OUT[tid] -= 5;
/*Inst  -*/   }
/*Inst  -*/ } else {
/*Inst  6*/   OUT[tid] += 20;
/*Inst  7*/   if (tid < 3) {
/*Inst  8*/     OUT[tid] -= 10;
/*Inst  -*/   } else {
/*Inst  9*/     OUT[tid] -= 5;
/*Inst  -*/   }
/*Inst  -*/ }
/*Inst 10*/ OUT[tid] += 100;
}

降低线程发散带来的性能降低的方法

线程发散会带来 GPU 执行性能的降低。减少分支发散的方法一般有：

减少一个 Warp 包含的线程的数量。例如，假设程序逻辑为：当线程 ID 小于 2 时执行 if 语句内容，且一个 Warp 是 4 个线程。那么此时会发生线程分散。如果设置一个 Warp 只包含 2 个线程，那么就不会发生线程分散了。该方法降低了线程发散的概率。
- 英伟达一个 Warp 包含 32 个线程
- ARM 一个 Warp 包含 4/8/16 个线程
将一个 Warp 中的线程分多次执行。例如，还是上面所述的程序，假设一个 Warp 包含 4 个线程，第一个时钟执行前两个线程，第二个时钟执行后两个线程。这也会降低线程分散的概率。
- 细心的读者会发现，这样做硬件资源减半，程序执行时间也加倍了。解决办法是：在保持硬件资源不变的情况下，该方案可以一次调度两个 Warp 执行。
- 英伟达的 Fermi 架构就是这样做的
分支合并
参看：fetch.php (cmu.edu)