CUDA(十) 周斌 CUDA程序基本优化

 

 

目录

  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%的性能提升。