CUDA编程：warp

参考： CUDA编程：warp - 知乎 (zhihu.com)

SM中的warp调度器会选择warp执行，

• selected warp(issued warp)
• stalled warp
• eligible warp

被选中的成为selected，没选中但已经具备执行条件的成为eligible，没准备好的成为stalled

1. warp

GPU以一个warp为基本单位管理线程，一个warp的大小通常是32.也就是说，一个线程束对应32个线程。但有时候线程不足以占用一个warp的所有空间，比如下面的情况：

一个简单的场景

由两个线程块组成的线程结构。每个线程块有39个线程。而GPU把线程块分配到SM上，也就是把39个线程分配到一个SM上。因为39不是32的倍数，所以划分为32 + 7的组合。第一个warp可以占满，而第二个warp就只能填7个线程了。这样就造成了资源的浪费。

至于为什么不将32-63放在一个warp中，是因为其中32-39以及40-63这两部分不是同一个线程块的，实际上应该被分配到两个SM，自然更不可能出现在一个warp中了。

所以，应该尽可能让block为32的倍数以避免warp资源的浪费。

2. warp divergence

线程发散：当某些线程需要执行一条指令，而其他线程不需要执行。这种情况就是线程发散。正常环境下，发散的的分支会是的某些线程处于空闲的状态，而其他线程执行分支中的代码。例如：

int val = 0;
if(threadIdx.x % 2 == 0)
    val = threadIdx.x;

通常来说，当某一个wrap中的线程出现发散情况，即warp divergence，会造成一定程度的性能损耗。看下面的粒子

if(tid % 32 < 16)
{
...
}
else
{
...
}

假设有两个warp，warp0和warp1。warp0中的idx为0-31，wrap1中的idx为32-63

出现线程发散的情况

红色部分的线程执行if时，紫色的线程无法执行else，所以会陷入阻塞状态。因为一个warp中的线程必须执行相同的指令，不能一部分执行if中的指令，一部分执行else中的。而不同的warp就没有这个问题，如下图，warp0全部执行if，warp1全部执行else。由此观之，在内核函数中引入if-else可能会引入线程发散的问题。

未出现线程发散

3. Resource partitioning

CUDA性能优化----warp深度解析 - Aamboo - 博客园 (cnblogs.com)

如何理解CUDA的资源划分？

我们已经清楚，内核函数中的抽象的每个block实际上是被划分到每个SM上处理的，SM又以一个wrap为基本单位。

每个SM有一个32位register集合放在register file中，还有固定数量的shared memory，这些资源都被thread瓜分了，由于资源是有限的，所以，如果thread数量比较多，那么每个thread占用资源就比较少，反之如果thread数量较少，每个thread占用资源就较多，这需要根据自己的需求作出一个平衡。

内存图

如果没有足够的资源可供分配，核函数就将启动失败

• warp分类：
  当一个block获得了足够的资源时，就成为active block。block中的warp就称为active warp。active warp又可以被分为下面三类：

1. selected warp

• actively executing

1. 1. stalled warp

• not ready for executing

1. 1. eligible warp

• ready for executing but not currently excuteing

SM中的warp调度器会选择warp执行，被选中的成为selected，没选中但已经具备执行条件的成为eligible，没准备好的成为stalled。
满足准备执行需要以下两个条件：

1. 1. 32 CUDA cores should be free
   2. all arguments of the current instruction for the warp should be ready

例如，Kepler架构GPU任何时刻的active warp数目必须少于或等于64个。selected warp数目必须小于或等于4个。如果一个warp阻塞了，调度器会挑选一个Eligible warp准备去执行。
CUDA编程中应该重视对计算资源的分配：这些资源限制了active warp的数量。因此，我们必须掌握硬件的一些限制，为了最大化GPU利用率，我们必须最大化active warp的数目。
5. Latency hiding
延迟隐藏的概念让我们进一步的理解active warp最大化的性能优势。
首先，什么是latency：

• • 即指令从刚开始被发布到完成执行的时钟周期数。
  • 通常，指令可以分为两种，每种对应一种latency：
    • Arithmetic instruction
    • Memory instruction

顾名思义，Arithmetic instruction latency是一个算术操作的始末间隔。另一个则是指内存传输的始末间隔。

• • GPU各种内存的latency是不同的，通常来说，shared memory的memory latency最低：

从上图可以看出，如果我们需要一个wrap进行算术运算，假设其指令延迟为15时钟周期，如果在这15个时钟周期中不做任何操作，等待算术指令在第16个时钟周期完成后再进行其他指令的发布，那么这15个cycle就是实打实的延迟了，因为我们什么都没干。

幸运的是，CUDA引入了延迟隐藏：

• • 我们就可以将每个指令的latency隐藏于issue其它warp的指令的过程中。
  • 思考如下的例子：
    • w1需要执行一个20clock cycles latency的算术指令。假设SM在t1处选择了w1作为selected warp，因为指令延迟为20，所以在t21处w1执行算术指令：

• • • 在同一个执行context中切换是没有消耗的，因为在整个warp的生命期内，SM处理的每个warp的执行context都是“on-chip”的。所以在t2处，我们改变execution context也是没有关系的，这样的话我们完全可以将其他指令的发布时间点就放在t2处：

• • • 以此类推，如果我们拥有足够的eligible warp，我们完全可以填充满w1的latency：

• • • 此时，我们来分析一下各个warp的指令延迟。假设w2也需要20 clock cycles。如果没有引入latency hiding，那么w2从22开始issue，经过20 cycles到42开始执行。但当我们引入了latency hiding，w2从2就已经issue了，这样其在22便以及可以执行指令。从宏观上看，w1刚在21处执行完，w2遍紧随其后执行完，这就相当于把w2的延迟隐藏在了w1的延迟中！这就是latency hiding
    • 顺带一提，根据上面的例子，我们可以计算一下所需的eligible warp的数量：
      • 假设1SM→128 cores→4 warps parallelly executed
      • 4 x 20 = 80，每个warp的延迟为20，我们需要20个eligible warp来覆盖20个cycle，四个并行执行的warp就对应4个20
      • 所以一个SM需要80个eligible warp

6. Occupancy
Occupancy就是每个SM的active warp占最大warp数目的比例：
Occupancy = Active warps / maximum warps

当一个warp阻塞了，我们希望有其他的eligible warp来occupy内核，使得内核永远保持busy，提高性能。
grid和block的配置准则：

• • • 保证block中thread数目是32的倍数
    • 避免block太小：每个blcok最少128或256个thread
    • 根据kernel需要的资源调整block
    • 保证block的数目远大于SM的数目
    • 多做实验来挖掘出最好的配置

可以通过CUDA Toolkit里面的Occupancy_Calculator来进行计算：