在上一期,我们提到的一个问题:
在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进行同步的时候,应当怎么做呢?
这个问题下期分解。