CUDA-Performance considerations

Programming Massively Parallel Processors

Posted by BY J on July 24, 2019

性能考量

在不同的应用程序中,由于不同因素,经常会遇到性能瓶颈。我们需要考虑清楚是算法策略还是硬件因素的影响,基于这个认识,本章将给出算法模式的类型对高性能表现的影响。

全局内存带宽

影响CUDA Kernel表现的主要因素之一是对全局内存的数据访问。本节将介绍Memory coalescing技术,该技术经常与tiling技术一起应用以高效利用全局内存带宽。 CUDA设备中全局内存的实现是由DRAMs技术实现的。数据位存储在叫做DRAM cell的小电容器中。 现代DRAM用并行技术提高数据访问率,通常叫做内存访问吞吐量。 通常电容器越大,速度越快。但是现在每单元电容器的体积越来越小,所以速度并没有提升。 每次访问DRAM一个位置,实际上访问了小范围的连续位置。DRAM上传感器都是并行工作的,每一个连续位置内的传感器都存储了数据的一位。一旦访问到一个,其他的都可以高速传输到处理器中。这种方式称为DRAM bursts。如果我们可以充分利用这种特性,DRAM则可以很快的提供数据,这比随意访问任意位置高效的多。 认识到DRAM的burst特性后,当前GPU 允许编程者将线程访问组织成适当形式以实现高效数据访问。这利用了warp线程中执行相同指令的特点。例如当线程0访问N位置,线程1访问N+1,等,则他们会被固化成单一指令进行访问。

5.3 Warps and SIMD HardWare

当前的CUDA设备会捆绑线程执行程序,此种策略会因为代码的类型造成表现差异。 第三章,有介绍每一个线程块被分成了warps.Warps的执行是通过SIMD硬件执行的。这种技术可以减少硬件制作成本,低电力消耗以及实现固化内存访问。然而在实现过程中,warp的大小很容易改变。当前,所有的CUDA设备warp配置几乎都设置为每个warp包含32个线程。 display ###说明: 将线程组织成warp的动机如上图所示:每一个处理器有一个控制单元可以取并解码指令。相同的控制信号进入多个处理单元,每一个都执行了一个warp的线程中的一个。因为所有的处理单元被相同的指令控制,则执行结果的不同是因为register file中存储的不同数据。这在处理器设计中就叫做Single-Instruction-Multiple-Data(SIMD). 现代处理器中的控制单元相当复杂。多个处理单元共享一个控制单元可以减小硬件制作成本以及电力消耗。这种设计在未来处理器中已成为一种趋势。 对于一个block的大小不是32的倍数,最后的warp将被其他线程填充以凑够32个线程。 对于二维数组,以行展开,组成warp.三维则以x,y组成二维数组,按z方向展开。 在if-else的例子中,相同warp中的线程会执行不同的执行路径,我们定义这种情况为线程的diverge。Thread Divergence 会严重影响程序的执行时间,因为它们是串行执行。 如果一个循环条件是建立在线程的索引序号上,那么就会发生线程的divergence. 一个广泛发生线程divergence的原因是在映射线程到数据,处理边界条件时。不过,对表现的影响随着数据的大小发生变化。通常,向量越大,随表现的影响越小。(32的倍数,长度为100的向量,1/4个warp会D,1000,则1/32。) 在一些重要的并行算法中,Control divergence会随着参与计算的线程的数量的增加而增加。下面将以reduction algorithm为例进行介绍 reduction算法从数组值中取出一个值。该值可以是和,最大值胡总和最小值。所有这些类型的reduction算法的计算结构都相同。首先,可以通过串行计算实现,每个元素均被访问一次,算法的执行时间与元素数量成正比。算法复杂度为O(N)。 因此,可以想到通过并行计算以减少执行时间。可以类比世界杯足球比赛。可以想象1024个队伍,通过10轮就可以得出最后的胜利者,前提是需要有足够的球场。 下图展现了实现并行求和reduction的核函数。原始数组存储在全局内存中。每一个线程块减少了数组的一部分,通过将该部分的元素加载进共享内存并在这些元素上执行并行reduction。代码将输入数组X中的元素从全局内存中加载进共享内存中。reduction是原地进行的,这就意味着共享内存的一些元素会被部分求和值替代。核函数的for循环中,每一个迭代都会实现一轮reduction。

__shared__ float partialSum[SIZE];
partialSum[threadIdx.x] = X[blockIdx.x * blockDim.x + threadIdx.x];
unsigned int t = threadIdx.x;
for(unsigned int stride = 1; stride < blockDim.x; stride *= 2)
{
__syncthreads();
if(t % (2 * stride) == 0)
    partialSum[t] += partialSum[t + stride];
}

__syncthreads()确保了for-loop循环中之前迭代的所有部分和值已经被产生了并且之前的任意一个线程被允许开始当前迭代。此种方式下,所有进入第二次迭代的线程将利用第一次迭代产生的值。在第一轮迭代后,偶数元素将被第一轮的求和值替代掉。 第二轮迭代后,元素下标索引的4的倍数的元素将被求和值替换掉。最后一轮,整个数组的总和将在partialSum[0]得到。 for循环中,stride首先初始化为1.在第一次迭代时,if条件将选取偶数线程执行两个相邻元素的加法。下图展示了核函数的执行。线程与数组元素以水平方向展示。由线程执行的迭代从上到下在竖直方向展示。每一行则展示了迭代后数组的元素内容。 5-1 Divergence分析: 在第一轮迭代中偶数线程将执行加法操作。在接下来的迭代中,更少的线程将执行if之后的加法操作,但是所有的线程在每一轮迭代中仍将执行判断操作。此divergence可以通过对算法的轻微改变实现性能提升。

__shared__ float partial[SIZE]
partialSum[threadIdx.x] = X[blockIdx.x + blockDim.x * blockIdx.x];
unsigned int t = threadIdx.x;
for(unsigned int stride = blockDim.x / 2; stride >=1; stride = stride>>1)
{
__syncthreads();
if(t < stride)
    partialSum[t] += partialSum[t + stride]; // Line 7
}

修改版本将stride初始化为数组长度的一半,表现差异性分析表现: 两者为什么会有不同呢?答案就在执行Line 7的线程与不执行Line 7的线程,他们的位置差异。 在第一轮迭代中,所有线程索引小于数组长度一半的threadIdx.x执行了Line7. 对于一个512个元素的数组, 第一轮迭代中,Thread 0到255执行了加法,256到511没有。在第一轮迭代之后,两个元素的加和存储在了元素0到255中。由于warp中32个线程,也就是warp 0到warp 7执行了加法,而8到15则略过了。由于每个warp中的所有线程都执行了相同指令,所以并没有divergence发生。 读者应该注意到,此图中并没有完全消除divergence。在5th个循环开始,执行Line7的线程将小于32.但是这也将发生divergence的循环由10降到了5. 5-2 两图中虽然改变很小,但是造成的算法表现差异却很大。这需要编程者很好的理解SIMD硬件中线程的执行,从而作出相应的调整。