阅读 《大规模并行处理器程序设计》影印版心得 第六章 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的数组元素的和......；则经过第一轮以后，所有的和都放在了编号小于25