阅读 《大规模并行处理器程序设计》影印版心得 第六章 Performance Consideration

 

6.1 More on Thread Execution

 

warp的概念

 

warp是如何组织的：按x,y,z逐渐增大的方式来线性化多维方式组织的线程，然后从前往后，每32个线程为一个warp

 

The hardware executes an instruction for all threads in the same warp before moving to the next instruction. This is so called SIMD

 

但是有分支的情况下，warp表现就不那么好。就需要两个或多个pass去执行不同的分支。

 

reduction algorithm为例。图6.3显示了常规的并行化方案，最开始是编号为偶数（2的倍数）的线程计算编号为偶数的数组元素和其后续相邻（+1）元素的和，并且把和存储在编号为偶数的数组元素的位置；然后是编号为4倍数的线程计算编号为4的倍数的数组元素和其后续相邻（+2）元素的和，并且把和存储在编号为4的倍数的数组元素的位置；然后是编号为8的倍数的线程计算编号为8的倍数的数组元素和其后续相邻（+4）元素的和，并且把结果存储在编号为8的倍数的数组元素的位置......

 

图6.5显示了一个改进的并行化方案。主要是考虑到warp的结构：threadIdx.x值连续的32个线程组成一个warp。则巧妙设计如下：最开始是编号为0的线程计算编号0的数组元素和编号为256的数组元素的和，并且把和存储在编号为0的数组元素的位置，编号为1的线程计算编号为1的数组元素和编号为257的数组元素的和......；则经过第一轮以后，所有的和都放在了编号小于256的数组元素的位置了。并且，编号为0线程开始的warp（线程0到线程31）执行相同的分支，因此，不会增加多余的pass，节省了时间，提高了性能。

 

但是图6.4的算法并不能完全消除divergence，尤其是第五轮以后还会出现分支执行步调不一致的情况。

 

6.2 Global Memory Bandwidth

 

主要讨论global memory使用时的问题。 目标是最好一次性读出一段连续的内存空间，这样效率最高，性能最好。

 

注意：我们要达到的目的是连续的thread最好能读取连续的global memory空间。图6.6就显示了这样的效果。图6.6A图中，两个连续的thread在每次迭代中，都是读取同样列号的矩阵元素。则这些元素的存储空间不是连续的，对性能有负面影响。而在图6.6B图中，两个连续的thread在每次迭代中，都去读取同样行号的矩阵元素。则这些元素的存储空间是连续的，即同样一个warp（线程号threadIdx.x连续）中的线程读取连续的global memory空间，有助于提升性能。这里理解的一个关键是，warp中的线程的运行是SIMD的，每次迭代时，是先运行第一个线程的读取一个元素的操作，然后是下一个连续线程的读取一个元素的操作。因此，需要把这两个线程读取的空间连续起来，提升性能。

 

当算法是在没有办法让连续号的线程读取联系的存储空间时，可以考虑用tile技术，把global memory中的数据依次拷贝到shared memory中来提升性能。此时，即使相邻线程不读取联系空间，性能也不会太差。因为shared memory带宽比global memory宽多了。

 

 

 

6.3 Dynamic Partitioning of SM Resources

 

SM资源包括寄存器，thread block slot和thread slot，GPU可以动态配置这些资源的数量。但是要注意的是，这些资源都是有数量上限的。不合理的配置资源的数量，可能会导致SM的能力不能被充分利用。

 

对寄存器来说，也可能会发生performance cliff，即一个很小的资源利用的增加，可能会导致程序并行性方面的大的减弱并进而影响程序的性能。一个例子就是增加一个自动变量（寄存器变量），可能会导致一个SM中运行的block数量的减少（因为寄存器变量的数量），并进而导致总线程数量的减少。

 

但是在另外一些情况下，增加自动变量，却可能导致减少global memory访问的开销，并进而提升程序的性能。这里的一个计算很有意思：在一个global memory load和its use之间，如果有4条独立指令，并假定买个指令的执行时间是4个时钟。则4条指令将执行16个时钟。已知一个global memory latency是200个时钟，因此，我们至少需要200/16 = 14个warp来达到zero-overhead scheduling。即当一个线程进入等待global memory latency时，如果有至少14个warp，则一定可以找到一个warp来执行。如果增加自动变量能增加global memory load和its use之间的独立指令条数时，则用更少的warp都可以达到一定可以找到warp来执行的效果。

 

6.4 Data Prefetching

 

Data prefetching，跟6.3中最后谈到的内容有些相似。主要思想是：矩阵相乘算法中，把global memory中存储的Md元素存取到shared memory中，其实需要两步：1）从global memory到registers；2）从registers到shared memory中。但是原来的算法，在这两步之间无independent instructions，无法在内存读写忙时，提供更多的代码来执行。其改进思想如图6.13显示。先进行一个data prefetching（从global memory读到registers），然后进入循环。循环中，先把先前读入到registers中的数据传递到shared memory中，然后同步。然后在把下一步中要用到的数据（next tile）从global memory中读取到registers中，然后进行真正的点积运算（贡献出很多的independent instruction）。这里有两个假定：1）数据读取可以是异步的，即发出读取指令后，程序继续运行。只有程序需要用到要读取的数据时，线程才进入等待状态；

 

 

6.5 Instruction Mix

 

这一节很好理解。就是把循环语句改变为不循环的语句。即unroll the loop。 

 

把for (i = 1; i<=10;i++)  sum+=i;  变为sum = 1+2+3+4+...+10;  这样做，可以提升性能。少了i变量的加减运算，同时提升了对数组元素的访问性能。

 

 

6.6 Thread Granularity

 

主要意思是：很多情况下，把线程进行合并，可以有效利用线程所读取的数据（读一次，利用2次），提升性能。

 

 

6.7 Measured Performance and Summary

 

主要结论是：tiles技术是提升性能的关键。当tiles足够大时，则loop unrolling和data prefetching技术也变得重要起来。

 

这里是主要的一个研究方向：即自动优化技术，在特定硬件上综合各种优化技术达到最大的性能，而且要求这一过程是程序实现，自动进行的。有参考文献可以看。