目录
Parallel Reduction 并行归约:线程和线程的调度和分割
Warp 分割
Memory Coalescing访存合并
Bank冲突:在shared memory
SM资源动态分割:硬件
数据预读
指令混合
循环展开
有效的数据并行算法+针对GPU架构特性的优化 = 最优性能
Parallel Reduction并行规约
回顾Parallel Reduction(sum)
规约:把一组很大的数据通过某种综合性的运算获得一组很少的数据,数据在不断的减少。
Google mapreduce 模型也是用到了数据的并行规约
将数据两两求和,第一组得到4个,第二组得到2个,最后得到结果
类似于淘汰赛n个元素进行log(n)个回合
第一次线程访问邻近的数据,第二次需要跳着访问,第三次线程跳着四个数据去访问。
假定在shared memory里面做并行规约
__shared__ float partialSum[];
// ... load into shared memory,累加的基本元素放在shared memory里面,提升访存性能
unsigned int t = threadIdx.x;//假定线程ID,
for (unsigned int stride = 1; stride < blockDim.x; stride *= 2)//假定有8个线程,第一次3次循环,步长stride=1、2、4
{
__syncthreads();//同步,保证每一步做完之后在做下一步操作.
if(t % (2 * stride) == 0)//t线程标号*步长stride=1、2、4;
partialSum[t] += partialSum[t + stride];//加上对应步长的位置,在同一块shared memory里面进行累加;当步长增加时,多余的线程,做相同操作(一个warp),不操作寄存器
}
八个数据累加
第一轮、线程1、3、5、7闲置(0开始);n个元素实际只需要n/2个线程
第二轮、线程2、6闲置。
第三轮、线程4闲置。
总之每一轮需要的线程数减半。
改进一下:
排列:不取临近的数据,只改变步长,步长变成4、2、1
__shared__ float partialSum[];
// ... load into shared memory
unsigned int t = threadIdx.x;
for (unsigned int stride = blockDim.x/2; stride < blockDim.x; stride /= 2)
{
__syncthreads();
if(t < stride)
partialSum[t] += partialSum[t + stride];
}
八个线程启动
第一轮、线程4、5、6、7闲置(0开始);n个元素实际只需要n/2个线程
第二轮、线程2、3闲置。
第三轮、线程1闲置。
总之每一轮需要的线程数减半。线程所处位置不同。
之前线程交叉减半,之后是前一半和后一半
if(t % (2 * stride) == 0)//stride=1、2、4; partialSum[t] += partialSum[t + stride]; | if(t < stride) partialSum[t] += partialSum[t + stride]; |
stride=1、2、4 | stride=4、2、1 |
将提前闲置的资源进行释放,第一中因为warp在,因此无法将占用资源释放出来,第二种相率更高一点。
Warp分割:块内线程如何划分warp
通晓warp分割有助于:减少分支发散、让warp今早完工。释放占用资源
- Block被划分为以连续的32为单位的线程组叫做warp。(织布机里的线束0-31,32-63)
- Warp是最基本的调度单元。以warp为单元发射线程指令。
- Warp一直执行相同指令(SIMT),同步执行
- 每一个线程只能执行自己的代码路径。若出现分支发散,divergent,大部分都在等待,把所有分支都完成,完成工作时间延长,warp间可以做不同的事情。
- Fermi SM有2个warp调度器(Tesla has 1)
- 设备切换没有时间代价,GPU上下文已经存在实际空间里面,只是需要将开关拨到实际单元
- 许多warps在一起可以隐藏访存延时
分割原则:threadIdx连续增加的一组
一维的block
threadIdx.x 0~512(G80/GT200)
第n个warp
起始线程ID:32n
结尾线程ID:32(n + 1) - 1
如果块大小不是32的倍数,最后一个warp将被填充
Warp 0 | Warp 1 | Warp 2 | Warp 3 |
0...31 | 32...63 | 64...95 | 96...127 |
二维Block,以行作为主元linearized order
增长threadIdx意味着
增长threadIdx.x
始于行threadIdx.y == 0,1,2...
三维Block,以行作为主元linearized order
始于threadIdx.z == 0,1,2...
分割为二维block
重复增长threadIdx.z
divergent分支 :warp存在分支发散。
Not all ALUs do useful work! Worst case 1/8 peak performance
给定warpSize == 32, 以下代码是否有哪个warp存在分支发散
if (threadIdx.x > 15 )
{
}
以15为界,前一半和后一半做不同的事情
任意warpSize > 1,一下代码是否有哪个warp存在分支发散,不存在分支发散
if (threadIdx.x > warpSize - 1)
{
}
if(t % (2 * stride) == 0)//stride=1、2、4; partialSum[t] += partialSum[t + stride]; | if(t < stride) partialSum[t] += partialSum[t + stride]; |
stride=1、2、4 | stride=4、2、1 |
两种分割方式哪一种并行规约更好一些
加入warpSize = 2;
第一轮:有四个发散分支 | 团结的
第二轮:有两个分支发散 | 没有
第三轮:有一个分支发散 | 有一个发散分支
当剩余元素小于warpSize时,必然要存在分支发散
好的 分割可以让warp尽早的完成,50%的性能提升。