通用图形处理器架构 #3 SIMT核心

XGundam05 · 发表于 2022-6-11 15:51

在传统的GPU渲染中，GPU需要访问诸如纹理细节贴图等较大的数据，这些数据并不能完全被片上缓存所存储。为了获得更高的性能以及可编程性，我们有必要提高off-chip的缓存带宽（带宽过滤），因此现今的GPU并发执行数千个线程，每一个线程所拥有的存储空间很小， cache可以有效地减少相当数量的片外内存访问。例如在典型的图形负载中，相邻像素的操作具有很高的空间局部性，因此可以被片上缓存命中。

上图展示了一个典型的SIMT核心流水线，这个流水线可以被分为SIMT front-end和SIMD back-end。在这个流水线中存在三个调度循环：一个取指循环、一个指令发射循环和一个寄存器访问循环。取指循环包含了Fetch、I-Cache、Decode和I-Buffer四个阶段，指令发射循环包含了I-Buffer、Score Board、Issue和SIMT-Stack四个部分，最后寄存器访问循环包含了Operand Collector、ALU和Memory三个阶段。
从最简单的系统起步

我们从一个简单的系统开始，将整个系统看作一个GPU附带一个线程调度器。为了提高效率，在英伟达的术语中线程是通过warp进行管理的，因此线程调度器的基本单元是一个warp。在每一个时钟周期线程调度器都会先择一个warp进行调度。在单层循环近似中，warp的程序计数器会通过访问指令缓存来寻找下一条执行的指令。在取到一条指令后，首先要对指令进行解码，操作数寄存器地址会从寄存器堆取回，与此同时会对SIMT线程的mask进行计算。
在执行mask和操作数地址准备好后，执行会以单指令多数据的行为进行，每一个线程执行的功能单元（lane）是由执行mask决定的。在现代CPU设计中，功能单元通常是异构的，这也就意味着每一个功能单元可能只能提供指令功能的一部分。例如英伟达GPU包含了special function unit、load/store unit、浮点计算单元、整数计算单元，在较新的架构中还包含Tensor Core。
所有的功能单元对外暴露的lane的数量与一个warp内的线程数量相同。但实际上为了节省不必要的硬件资源开销，GPU可以通过将执行器件进行流水化或加深流水深度实现在几个周期内执行完一个warp的线程，这可以实现更高的单位面积性能，但代价是能源消耗的增加。
现代GPU的一个重要特征是在程序员的角度看来，所有线程都执行了完全独立的路径，但实际上这只是一个虚拟出来的概念。在硬件实现中，这种运行模型有可能通过传统的分支预测实现，但现在更多是将其与一个被称为预测掩码stack实现的，这被称为SIMT stack。
SIMT stack可以高效的处理每一个线程独立运行时可能遇到的两个问题：首先是嵌套控制流，如果一个分支跳转依赖于其他的分支结果。第二个是整个warp中的线程在一个控制流中跳过了所有的计算。对于复杂控制流来说，SIMT stack可以节约相当数量不必要的开销，在支持分支预测的CPU中通常通过使用多个预测寄存器来处理嵌套的控制流。
do {
t1 = tid*N; // A
t2 = t1 + i;
t3 = data1[t2];
t4 = 0;
if( t3 != t4 ) {
      t5 = data2[t2]; // B
      if( t5 != t4 ) {
         x += 1; // C
      } else {
         y += 2; // D
      }
} else {
      z += 3; // F
}
i++; // G
} while( i < N );
上面是一个嵌套控制流的例子，下面是对应的PTX汇编代码，可以参考。
A:  mul.lo.u32 t1, tid, N;
add.u32 t2, t1, i;
ld.global.u32 t3, [t2];
mov.u32 t4, 0;
setp.eq.u32 p1, t3, t4;
@p1 bra F;
B:  ld.global.u32 t5, [t2];
setp.eq.u32 p2, t5, t4;
@p2 bra D;
C:  add.u32 x, x, 1;
bra E;
D:  add.u32 y, y, 2;
E:  bra G;
F:  add.u32 z, z, 3;
G:  add.u32 i, i, 1;
setp.le.u32 p3, i, N;
@p3 bra A;

上面这张图描述了SIMT stack的工作过程。首先根据上面的代码可以看到整个程序根据分支可以分为ABCDEFG几个程序段。我们假设现在有一个包含4个线程的warp从A block开始运行，我们看到在图a中每一个代码块后都有一个二进制串，例如A/1111就代表当前warp中的四个线程都会运行该代码块。
图c的初始状态我们可以看到栈中保存了指向代码块B、F和G的项，这是因为当我们运行完代码块A后，不同线程的下一条指令根据分支的不同会跳转到不同的代码块，首先我们找到最终这四个线程汇合的代码块G，将其压入栈中，然后依次压入F和B两个代码块。我们可以发现压入栈时我们需要考虑到不同线程在何时会汇合，并保证对应的指令在栈中保持相对顺序，因此在每一项的第一列我们保存了这些线程将会在哪里汇合。
由于我们的栈顶TOS现在指向了B，因此在下一步我们首先执行运行代码块B的三个线程（弹出B），在B中我们也发生了分支，因此根据汇合点我们首先将E压入，保证两个分支执行完后可以在E汇合，然后将代码块C和D压入栈中。代码块入栈的顺序也是一个有趣的问题，一般来说我们会让具有更多活动线程的块先入栈，但是在图c中我们首先压入了F，然后才是B。
每一项的第三列保存了当前代码块需要有哪些线程执行，这里比较简单就不再细说。
最后图b把每一个代码块中执行的线程用图像表示出来，还是很清楚的。
汇合点一般都会选择分支汇合最早的代码块，但是在实际运行中可能会出现例外。
SIMT死锁和非栈的SIMT架构

基于栈的SIMT实现有可能会导致一种被称为“SIMT死锁”的问题，具体将在下面进行描述。为了解决这个问题，英伟达在Volta架构上提出了一种新的线程管理方法，被称为“Independent Thread Scheduling独立线程调度”。

上图描述了一种可能出现SIMT死锁的情形，这里面出现了一个atomicCAS函数，其实这就是很常见的T&S锁。我们会首先比较互斥变量mutex是否为0，如果是0那么我们就将它置为1，否则什么也不做，等待下一次循环继续比较。这个操作对于每一个线程都是原子化的，即一次只有一个线程可以访问mutex并执行赋值操作，那么每次就有可能产生一个线程成功将mutex写入1，进入临界区。这个时候就出现了分支，一个线程需要退出循环而其他线程需要停留在循环中。回忆我们之前的stack结构，此时栈顶指向了B，我们的汇合点在C，也就是说我们必须等待所有线程完成B后才能弹出C。但是此时我们栈顶的那一个线程只有执行了C才可以释放锁，我们想执行C必须执行所有B，这样就产生了SIMT死锁。
为了解决这个问题，英伟达对每一个warp添加了一个“汇合屏障 convergence barriers”来替代原来的stack结构。

这张图展示了convergence barriers所需的信息，这些信息将会被存放在warp scheduler的硬件寄存器当中。Barrier Participation Mask用来追踪当前warp参与了那些汇合屏障，每一个warp可能有多个汇合屏障。我们还需要知道有哪些线程到达了某一个给定的屏障，这样我们就可以及时释放屏障让线程可以继续运行，Barrier State就是用来追踪有哪些线程到达了给定的屏障。Thread State用来追踪在当前warp内所有线程的状态，是可以执行、阻塞在某个屏障还是失效。失效状态的存在是为了解决可能存在的SIMT死锁问题。如果我们没有失效状态，当一个线程失效后，其他到达屏障的线程也无法跨越屏障，进而被阻塞。由于嵌套分支的存在，对一个warp我们无法知道嵌套分支的深度可能有多少，因此在一个warp中我们需要任意数量的convergence barriers。
为了使用合并屏障，我们引入一个特殊的指令“ADD”。在运行ADD后，所有活动的线程会被加入当前屏障，如果发生了分支，系统会自动选择对应的子集并更新Thread Active使得子集中的线程可以继续执行。与stack based的系统相比，使用汇合屏障的系统可以在不同分支间任意切换，这使得我们避免了潜在的死锁问题。
当一个分支到达屏障时，我们需要使用“WAIT”指令在Barrier Stat寄存器中更新对应分支线程的状态，将它们置为block。一旦所有的分支都到达当前屏障并且都执行了对应的WAIT指令，系统就可以将这个屏障中所有的线程重新置为active。为了实现不同分支之间的切换，英伟达使用了一个名为YIELD的指令来实现非直接分支的功能，在此不继续深入讨论。

上图为stack based的warp分支执行。

上图为使用汇合屏障的warp分支执行。
Warp Scheduling

在一个SM处理器中我们可能有很多个Warp，我们也需要通过Warp调度来隐藏可能的cache miss。那么当我们应该以什么顺序来选择下一个执行的warp呢？一个最简单的想法是按照编号顺序轮流执行，轮流执行可以给每一个warp分配近乎相同的执行时间。如果每一个核心执行warp的数量乘以指令发射的时间超过了访存延迟，那么所有的核心将会一直处于繁忙，因此我们可以使用这种方式来增加每一个核心的吞吐量。
但是这其中也有需要我们权衡的地方，如果想要使不同的warp在每一个周期都可以发射自己的指令，那么对应的我们需要给每一个warp配备独立的寄存器。因此，增加每个核的warp数会增加用于寄存器的芯片面积，对于一个固定的芯片面积，增加每个核心的warp会减少每个芯片的总核心数。
在实际应用中，取决于运行程序的不同，我们对warp的调度应该采取不同的策略。例如如果我们进行图形任务，不同的warp间访问的纹理数据具有时空局部性，因此我们可以使用轮流执行来平衡每一个warp执行的时间，同时也可以充分利用cache hit带来的性能提升。同样，对于DRAM临近区域的访问也会具有更高的效率，因此对于这种情形，我们更愿意采用轮询的执行方式。但是如果每一个warp访问的数据不具有时空局部性，也就是说访问的位置差别较大，此时我们应该倾向于重复执行同一个warp来获取最大的cache hit，避免warp切换导致的访存开销。

		自动登录	找回密码
密码			立即注册

通用图形处理器架构 #3 SIMT核心

本帖子中包含更多资源