CUDA性能优化----warp深度解析

本文转自：http://blog.163.com/wujiaxing009@126/blog/static/71988399201701224540201/

一、引言

CUDA性能优化----sp, sm, thread, block, grid, warp概念中提到：逻辑上，CUDA中全部thread是并行的，可是，从硬件的角度来讲，实际上并非全部的thread可以在同一时刻执行，接下来咱们将深刻学习和了解有关warp的一些本质。

 

二、Warps and Thread Blocks

warp是SM的基本执行单元。一个warp包含32个并行thread，这32个thread执行于SIMT模式。也就是说全部thread执行同一条指令，而且每一个thread会使用各自的data执行该指令。

block能够是1D、2D或者3D的，可是，从硬件角度看，全部的thread都被组织成一维的，每一个thread都有个惟一的ID。每一个block的warp数量能够由下面的公式计算得到：

  一个warp中的线程必然在同一个block中，若是block所含线程数目不是warp大小的整数倍，那么多出的那些thread所在的warp中，会剩余一些 inactive的thread，也就是说，即便凑不够warp整数倍的thread，硬件也会为warp凑足，只不过那些thread是inactive状态， 须要注意的是，即便这部分thread是inactive的，也会消耗SM资源，这点是编程时应避免的。

 

三、Warp Divergence(warp分歧)

控制流语句广泛存在于各类编程语言中，GPU支持传统的、C-style的显式控制流结构，例如if…else,for,while等等。

CPU有复杂的硬件设计能够很好的作分支预测，即预测应用程序会走哪一个path分支。若是预测正确，那么CPU只会有很小的消耗。和CPU对比来讲，GPU就没那么复杂的分支预测了。

这样问题就来了，由于全部同一个warp中的thread必须执行相同的指令，那么若是这些线程在遇到控制流语句时，若是进入不一样的分支，那么同一时刻除了正在执行的分支外，其他分支都被阻塞了，十分影响性能。这类问题就是warp divergence。

注意，warp divergence问题只会发生在同一个warp中。下图展现了warp divergence问题：

 

为了得到最好的性能，就须要避免同一个warp存在不一样的执行路径。避免该问题的方法不少，好比这样一个情形，假设有两个分支，分支的决定条件是thread的惟一ID的奇偶性，kernel函数以下(simpleWarpDivergence.cu)：

__global__ void mathKernel1(float *c) 

{
int tid = blockIdx.x * blockDim.x + threadIdx.x;
float a, b;
a = b = 0.0f;
if (tid % 2 == 0)
a = 100.0f;
else
b = 200.0f;
c[tid] = a + b;
}

一种方法是，将条件改成以warp大小为步调，而后取奇偶，代码以下：

__global__ void mathKernel2(void) 

{
int tid = blockIdx.x * blockDim.x + threadIdx.x;
float a, b;
a = b = 0.0f;
if ((tid / warpSize) % 2 == 0)
a = 100.0f;
else
b = 200.0f;
c[tid] = a + b;
}

经过测试发现两个kernel函数性能相近，到这里你应该在奇怪为何两者表现相同呢，其实是由于当咱们的代码很简单，能够被预测时，CUDA的编译器会自动帮助优化咱们的代码。（ 稍微提一下GPU分支预测，这里一个被称为预测变量的东西会被设置成1或者0，全部分支都会执行，可是只有预测变量值为1时，该分支才会获得执行。当条件状态少于某一个阈值时，编译器会将一个分支指令替换为预测指令。）所以，如今回到自动优化问题，一段较长的代码就可能会致使warp divergence问题了。

可使用下面的命令强制编译器不作优化：

$ nvcc -g -G -arch=sm_20 simpleWarpDivergence.cu -o simpleWarpDivergence

四、Resource Partitioning(资源划分)

一个warp的context包括如下三部分：

1. Program counter
2. Register
3. Shared memory

再次重申，在同一个执行context中切换是没有消耗的，由于在整个warp的生命期内，SM处理的每一个warp的执行context都是“on-chip”的。

每一个SM有一个32位register集合放在register file中，还有固定数量的shared memory，这些资源都被thread瓜分了，因为资源是有限的，因此，若是thread数量比较多，那么每一个thread占用资源就比较少，反之若是thread数量较少，每一个thread占用资源就较多，这须要根据本身的需求做出一个平衡。

资源限制了驻留在SM中blcok的数量，不一样的GPU，register和shared memory的数量也不一样，就像Fermi和Kepler架构的差异。若是没有足够的资源，kernel的启动就会失败。下图是计算能力为2.x和3.x的device参数对比：

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

1. Selected warp
2. Stalled warp
3. Eligible warp

SM中 warp调度器每一个cycle会挑选active warp送去执行，一个被选中的warp称为 Selected warp，没被选中，可是已经作好准备被执行的称为 Eligible warp，没准备好要被执行的称为 Stalled warp。warp适合执行须要知足下面两个条件：

1. 32个CUDA core有空
2. 全部当前指令的参数都准备就绪

例如，Kepler架构GPU任什么时候刻的active warp数目必须少于或等于64个。selected warp数目必须小于或等于4个（由于scheduler有4个？不肯定，至于4个是否是太少则不用担忧，kernel启动前，会有一个 warmup操做，可使用cudaFree()来实现）。若是一个warp阻塞了，调度器会挑选一个Eligible warp准备去执行。

CUDA编程中应该重视对计算资源的分配：这些资源限制了active warp的数量。所以，咱们必须掌握硬件的一些限制，为了最大化GPU利用率，咱们必须最大化active warp的数目。

 

五、Latency Hiding(延迟隐藏)

指令从开始到结束消耗的clock cycle称为指令的latency。当每一个cycle都有eligible warp被调度时，计算资源就会获得充分利用，基于此，咱们就能够将每一个指令的latency隐藏于issue其它warp的指令的过程当中。

和CPU编程相比， latency hiding对GPU很是重要。CPU cores被设计成能够最小化一到两个thread的latency，可是GPU的thread数目可不是一个两个那么简单。

当涉及到指令latency时，指令能够被区分为下面两种：

1. Arithmetic instruction
2. Memory instruction

顾名思义，Arithmetic  instruction latency是一个算术操做的始末间隔。另外一个则是指load或store的始末间隔。

两者的latency大约为：

1. 10-20 cycle for arithmetic operations
2. 400-800 cycles for global memory accesses

下图是一个简单的执行流程，当warp0阻塞时，执行其余的warp，当warp变为eligible时从新执行。

  你可能想要知道怎样评估active warps 的数量来hide latency。Little’s Law能够提供一个合理的估计：

对于Arithmetic operations来讲，并行性能够表达为用来hide  Arithmetic latency的操做的数目。下表显示了Fermi和Kepler架构的相关数据，这里是以(a + b * c)做为操做的例子。不一样的算术指令，throughput（吞吐）也是不一样的。

这里的throughput定义为每一个SM每一个cycle的操做数目。因为每一个warp执行同一种指令，所以每一个warp对应32个操做。因此，对于Fermi来讲，每一个SM须要640/32=20个warp来保持计算资源的充分利用。这也就意味着，arithmetic operations的并行性能够表达为操做的数目或者warp的数目。两者的关系也对应了两种方式来增长并行性：

1. Instruction-level Parallelism（ILP）：同一个thread中更多的独立指令
2. Thread-level Parallelism （TLP）：更多并发的eligible threads

对于Memory operations，并行性能够表达为每一个cycle的byte数目。

由于memory throughput老是以GB/Sec为单位，咱们须要先做相应的转化。能够经过下面的指令来查看device的memory frequency：

$ nvidia-smi -a -q -d CLOCK | fgrep -A 3 "Max Clocks" | fgrep "Memory"

以Fermi架构为例，其memory frequency多是1.566GHz，Kepler的是1.6GHz。那么转化过程为：

 

乘上这个92能够获得上图中的74，这里的数字是针对整个device的，而不是每一个SM。

有了这些数据，咱们能够作一些计算了，以Fermi架构为例，假设每一个thread的任务是将一个float(4 bytes)类型的数据从global memory移至SM用来计算，你应该须要大约18500个thread，也就是579个warp来隐藏全部的memory latency。

Fermi有16个SM，因此每一个SM须要579/16=36个warp来隐藏memory latency。

 

六、Occupancy(占用率)

当一个warp阻塞了，SM会执行另外一个eligible warp。理想状况是，每时每刻到保证cores被占用。Occupancy就是每一个SM的active warp占最大warp数目的比例：

咱们可使用cuda库函数的方法来获取warp最大数目：

cudaError_t cudaGetDeviceProperties(struct cudaDeviceProp *prop, int device);

而后用 maxThreadsPerMultiProcessor来获取具体数值。

grid和block的配置准则：

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

Occupancy专一于每一个SM中能够并行的thread或者warp的数目。无论怎样，Occupancy不是惟一的性能指标，当Occupancy达到某个值时，再作优化就可能再也不有效果了，还有许多其它的指标须要调节。

 

七、Synchronize(同步)

同步是并行编程中的一个广泛问题。在CUDA中，有两种方式实现同步：

1. System-level：等待全部host和device的工做完成
2. Block-level：等待device中block的全部thread执行到某个点

由于CUDA API和host代码是异步的，cudaDeviceSynchronize能够用来停下CPU等待CUDA中的操做完成：

cudaError_t cudaDeviceSynchronize(void);

由于block中的thread执行顺序不定，CUDA提供了一个函数来同步block中的thread。

__device__ void __syncthreads(void);

当该函数被调用时，block中的每一个thread都会等待全部其余thread执行到某个点来实现同步。

 

八、结束语

CUDA性能优化是一个多方面、复杂的问题，深刻了解warp的概念和特性是CUDA性能优化的一个关键和开始。

 

参考： CUDA性能优化----sp, sm, thread, block, grid, warp概念