参考: 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又可以被分为下面三类:
- selected warp
- actively executing
-
- stalled warp
- not ready for executing
- ready for executing but not currently excuteing
SM中的warp调度器会选择warp执行,被选中的成为selected,没选中但已经具备执行条件的成为eligible,没准备好的成为stalled。
满足准备执行需要以下两个条件:
-
- 32 CUDA cores should be free
- 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来进行计算: