大模型与AI底层技术揭秘（36）分裂没有出路

在上一期，我们提到的一个问题：

在GPU程序中，如果有if-else这样的分支，在一个warp中，不同的CUDA Core走到了不同的分支，此时scheduler应当如何发射后续的指令呢？

其实这个问题并不复杂。

如果GPU执行这样的代码：

if (condition) {
  //do sth.
} else {
  //do other sth.
}

假设一个warp中有32个线程，其中16个满足判断条件，另外16个不满足。那么，满足判断条件的16个线程会执行if分支中的语句，而另外16个线程则处于挂起状态。待这16个执行线程把if分支中的语句执行完毕后，它们会进入挂起状态，而另外16个执行线程则开始执行else分支中的语句，执行完毕后，这32个线程才可以继续并行执行后续的语句。

我们假定一个最坏的情况：某一个warp中，有32个线程各自执行各自的语句，那么实际上这32个线程只能串行执行。

因此，在GPU程序设计中，我们要记住：

统一才是正道，分裂没有出路！

为什么GPU的并行执行，在发生程序分支时只能串行执行呢？

这和GPU的架构有着紧密的关系。

在方老师多次为大家介绍的“图灵机”中，其核心运行机制是，读取（由二进制编码表示的）指令，并送到ALU进行执行。这一运行机制一直传承下来，无论是x86为代表的CISC，还是ARM为代表的RISC，在这一点上没有根本的差异。这一机制叫做指令发射。

当然，“图灵机”的执行过程是完全串行的。读取指令，对指令译码并发射到ALU，ALU执行指令，对存储器访问和写回存储器这五个步骤，必须逐个执行。

为了提升指令执行的效率，CPU设计还引入了一些新的机制。例如，流水线机制：

流水线机制可以让CPU中取指、译码发射、执行、内存访问和内存写回这五部分电路一直处于工作状态，理论上能提升5倍的指令吞吐量。

超线程则是另一种提升指令吞吐量的手段。所谓超线程就是在一个CPU核心中，搞多套取指单元、译码发射单元、访存单元和寄存器列（register file），并被操作系统视为多个vCPU。虽然每个CPU核心只有一个ALU，但由于ALU只有在执行指令时才使用，因此，超线程产生的多个vCPU实际上性能与真实的物理核心相差无几。支持超线程的CPU核心拥有多个取指单元和译码发射单元，可以让多条指令（看上去是）并行执行，（其实是时分复用ALU），因此，这种执行方式也可以被称为多指令多线程（MIMT）。

而在GPU中并非如此。

GPU的每个SM（Streaming Processor）中往往拥有上千个CUDA核，每个CUDA核都拥有自己的寄存器列。在Turing以后的架构中还有可能拥有Tensor核。但是，每个SM的取指单元和译码发射单元是小于这个数量的。

在安培（Ampere）架构的NVIDIA GPU中，每个SM有64套取指单元和译码发射单元，可以同时执行最多64个Warp；图灵（Turing）架构的NVIDIA GPU中，每个SM可以同时执行最多4个Warp；帕斯卡（Pascal）架构的NVIDIA GPU中，每个SM可以同时执行最多2个Warp。

SM中分配给每个Warp的只有一套取指单元和译码发射单元，也就是说，Warp中所有的线程要共用一个指令指针（Instruction Pointer）。因此，我们需要尽量减少同一个Warp中的分支指令。

如果实在要使用if-else这样的分支，我们可以让不同的分支运行到不同的wap，比如这样写：

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

第一个线程束内的线程编号tid从0到31，tid/warpSize都等于0，那么就都执行if分支。第二个线程束内的线程编号tid从32到63，tid/warpSize都等于1，执行else分支。这样就可以保证warp内没有分支，效率极高。

上期遗留的第二个问题是：

不同的执行次序会导致不同的CUDA Core的执行时间发生差异，那么，需要warp中，各个CUDA进行同步的时候，应当怎么做呢？

这个问题下期分解。