目录备注Reduction规约求和Reduce朴素版本Reduce1——使用共享内存Reduce2——去除分支发散和除法操作bank conflict计算Reduce3——减少Bank ConflictsReduce4——减少空闲线程Reduce5——warp展开Reduce6——for循环进一步展开Reduce7——网格步幅循环加载Reduce8——warp shuffle参考文献备注这个笔记是在自学过程的一些记录中间包含了很多个人的理解以及问AI的回答因为找不到很系统的教程官方文档有时候不会讲的特别细比方说某个指标的具体计算过程网上也没找到什么资料所以在这里记录下后续自己有了新的理解或者发现写的不对的地方都会在这不断更新如有不对的地方还请指正。Reduction规约求和Reduction就是对一组数据经过一些操作变换成一个结果这里以求和为例进行记录。每个数加上它右边间隔s的数开始s1每迭代一轮s*2迭代log₂ⁿ轮n是数据量这里假设n是2的幂次。迭代完后第一个数就是求和结果。测试过程数据是4000000个int的求和然后默认blocksize512显卡是RTX 4070 Super。Reduce朴素版本kernel代码如下__global__voidreduce_naive(int*g_idata,int*g_odata){inttidthreadIdx.x;intiblockIdx.x*blockDim.xthreadIdx.x;for(ints1;sblockDim.x;s*2){if(tid%(2*s)0){g_idata[i]g_idata[is];}__syncthreads();}if(tid0)g_odata[blockIdx.x]g_idata[blockIdx.x*blockDim.x];}这段代码相当于在每个block的512个thread里面做规约g_odata里面是每个block里512个数据的和最后把g_odata求和就得到最终结果。Reduce1——使用共享内存我们都知道共享内存能够帮助我们减少内存搬运耗时因为它是片上内存我们可以先把数据搬运到共享内存上再去计算。kernel代码如下__global__voidreduce1(int*g_idata,int*g_odata){extern__shared__intsdata[];inttidthreadIdx.x;intiblockIdx.x*blockDim.xthreadIdx.x;sdata[tid]g_idata[i];__syncthreads();for(ints1;sblockDim.x;s*2){if(tid%(2*s)0){sdata[tid]sdata[tids];}__syncthreads();}if(tid0)g_odata[blockIdx.x]sdata[0];}这里我用Nsight Compute分析对比上面两个kernel时发现耗时没有明显降低然后上网搜索了一下说是Naive的全局显存数据请求到数据后会暂存在L2 cache中复用所以相对于共享内存版本耗时差异不大。从下表中可以看到Naive版本的L2 cache throughput很高。PS这里其实没有完全理解在这里Mark一下。我当前的猜测是因为大部分操作数据在内存上都是相邻的一次事务可以取32个4字节的int数所以内存事务复用率较高所以和共享内存版本差异不大Reduce2——去除分支发散和除法操作Reduce1有两个问题一个是分支发散比较严重另一个是’%操作比较耗时需要优化一下。__global__voidreduce2(int*g_idata,int*g_odata){extern__shared__intsdata[];inttidthreadIdx.x;intiblockIdx.x*blockDim.xthreadIdx.x;sdata[tid]g_idata[i];__syncthreads();for(ints1;sblockDim.x;s*2){intidx2*s*tid;if(idxblockDim.x){sdata[idx]sdata[idxs];}__syncthreads();}if(tid0)g_odata[blockIdx.x]sdata[0];}从下表可以看出我们实现了规约操作且减少了分支发散和除法操作但是这里带来了一个新的问题就是引入了Bank conflict。例如第一轮迭代tid0的线程会读取idx0的数据tid16的线程会读取idx32的数据而这两个地址的数据在同一Bank0idx%320里面这就相当于不同线程在读取同一Bank的不同地址的数据这就是Bank conflict导致并行变成串行增加耗时。不过相比于前面的kernel还是能看到耗时明显降低说明之前的kernel分支发散和除法操作比bank conflict更耗时。bank conflict计算先按一个512个int、blocksize512的demo计算bank conflicts指标定义参考Nvidia Nsight Compute官方文档对于512个线程每个线程处理一个int数据可以分为512/3216个warpShare Store1、对于sdata[tid] g_idata[i];这是在给共享内存写入所以对应Share Store总共16个warp对应16个Instructions和Requests这里由于都是连续写入所以不存在bank conflictswavefronts162、对于sdata[idx] sdata[idx s]第一次迭代也就是s1的时候是只在idx——0~255的索引下写入所以只有8个warp总共9次迭代依次是8、4、2、1、1、1、1、1、1个warp在做共享内存写入所以对应84211111120个Instructions和Requests这里我举了s1和s16时的索引可以看到s1时idx和idxs都是间隔2的索引就是一个2-way的bank conflict一个warp就需要2次transactions依次类推s2时4-way的bank conflict需要4次transactions而到了s16时注意tid16时idx已经到了512索引已经超出了超出的这部分是无效的虽然这个warp是32-way的但是我只需要发送16次transactions就满足了后续迭代以此类推。最终得到1次load或1次store的wavefronts1616161616842195综合下来Share Store Instructions和Requests 16 20 36wavefronts9516111Share Load1、对于sdata[idx] sdata[idx s]分析思路和Share Store一致不过有个区别是这里每个线程读取是有两个地址idx和idxs所以Share Load Instructions和Requests 20 20 40wavefronts参考前面的Store分析wavefronts95*2190乘2是因为做了两次读取2、if (tid 0) g_odata[blockIdx.x] sdata[0];这行语句也有一个读取共享内存操作不过只做一次 sdata[0]相当于做了一次内存读取也就相当于一次wavefront综合下来Share Load Instructions和Requests 40 1 41wavefronts1901191而bank conflicts就相当于额外的wavefronts结果也就等于总的wavefronts减去请求的requestsShare Load Bank Conflicts 191 - 41 150Share Store Bank Conflicts 111 - 36 75注意到这里Load的冲突数刚好是Store的两倍这是因为发生冲突的语句sdata[idx] sdata[idx s]刚好是两次读取一次写入。我后面又尝试了不同尺寸的输入发现Share Load wavefronts不能通过单个block的wavefronts * 总的blocks计算而且每次跑可能会不一样虽然相比于总数不会差太多其他几个都可以。问了下AI说是Wavefronts 是流水线层面的“物理动作”。在 GPU 的硬件流水线中当某几个 Warp 的执行由于某些外部原因例如抢占、上下文切换、或规约中的 __syncthreads() 同步导致暂停再恢复时LSU 阶段可能为了确保数据的危害Hazard检测通过在极个别周期发生了物理重复请求Retry。这种极低概率的硬件自发抖动会被物理计数器记录下来。Reduce3——减少Bank Conflicts我们可以将求和的间隔由小到大改成由大到小如下图所示这样就可以避免绝大多数的Bank conflict可以看到Bank Conflits明显减少Reduce4——减少空闲线程前面的Reduce方法有比较多的线程冗余比如说我们每个block开了512个线程但是后面256个线程只做了一次数据搬运就没有其他操作了第一轮迭代只有前面256个线程在做加法运算运行后面256个线程就都空置了后续每次迭代空置的线程也就越来越多。为了把线程利用起来可以在把数据往共享内存搬运的时候就做一次加法比方说我一个block还是512个线程但是我用这512个线程处理1024个数据在把数据从全局内存搬运到共享内存的时候我读取第0个数据和第512个数据把他们相加然后写入共享内存然后一直到第511个数据和第1023个数据把他们相加然后写入共享内存这样就相当于每个线程都做了一次加法操作相比上面的Reduce把空闲线程也做了利用。__global__voidreduce4(int*g_idata,int*g_odata,intsize){extern__shared__intsdata[];inttidthreadIdx.x;intiblockIdx.x*(blockDim.x*2)threadIdx.x;if(isize)sdata[tid]g_idata[i]g_idata[iblockDim.x];elsesdata[tid]0;__syncthreads();for(intsblockDim.x/2;s0;s1){if(tids){sdata[tid]sdata[tids];}__syncthreads();}if(tid0)g_odata[blockIdx.x]sdata[0];}这里需要注意由于我们每个block是用512个线程处理1024个数据所以总的block数相比于之前的Reduce要减半。可以看到kernel耗时显著降低Reduce5——warp展开我们知道一个warp里面有32个线程且是同步执行的所以当间隔s从大间隔降到32的时候我们可以不用 __syncthreads()命令直接把内部规约展开注意kenel中for循环中止条件从s0变成了s32__device__voidwrapReduce5(volatileint*sdata,inttid){sdata[tid]sdata[tid32];sdata[tid]sdata[tid16];sdata[tid]sdata[tid8];sdata[tid]sdata[tid4];sdata[tid]sdata[tid2];sdata[tid]sdata[tid1];}__global__voidreduce5(int*g_idata,int*g_odata,intsize){extern__shared__intsdata[];inttidthreadIdx.x;intiblockIdx.x*(blockDim.x*2)threadIdx.x;if(isize)sdata[tid]g_idata[i]g_idata[iblockDim.x];elsesdata[tid]0;__syncthreads();for(intsblockDim.x/2;s32;s1){if(tids){sdata[tid]sdata[tids];}__syncthreads();}if(tid32)wrapReduce5(sdata,tid);if(tid0)g_odata[blockIdx.x]sdata[0];}可以看到kernel耗时进一步降低Reduce6——for循环进一步展开可以进一步展开kernel中的for循环templateunsignedintblockSize__device__voidwrapReduce(volatileint*sdata,inttid){if(blockSize64)sdata[tid]sdata[tid32];if(blockSize32)sdata[tid]sdata[tid16];if(blockSize16)sdata[tid]sdata[tid8];if(blockSize8)sdata[tid]sdata[tid4];if(blockSize4)sdata[tid]sdata[tid2];if(blockSize2)sdata[tid]sdata[tid1];}templateunsignedintblockSize__global__voidreduce6(int*g_idata,int*g_odata,intsize){extern__shared__intsdata[];inttidthreadIdx.x;intiblockIdx.x*(blockDim.x*2)threadIdx.x;if(isize)sdata[tid]g_idata[i]g_idata[iblockDim.x];elsesdata[tid]0;__syncthreads();// 展开循环if(blockSize1024){if(tid512){sdata[tid]sdata[tid512];}__syncthreads();}if(blockSize512){if(tid256){sdata[tid]sdata[tid256];}__syncthreads();}if(blockSize256){if(tid128){sdata[tid]sdata[tid128];}__syncthreads();}if(blockSize128){if(tid64){sdata[tid]sdata[tid64];}__syncthreads();}// Warp 级优化if(tid32){// 强制转换为 volatile 指针传递确保读写可见性wrapReduceblockSize((volatileint*)sdata,tid);}if(tid0)g_odata[blockIdx.x]sdata[0];}不过耗时并没有太大变化但是Mark Harris当时测试时有一定提升猜测是这些年硬件有一定的迭代优化Reduce7——网格步幅循环加载Mark Harris文档里最后还提了一个网格步幅循环加载Grid-Stride Loop的规约原理是考虑数据量很大然后以有限的block有限的线程去做运算就相当于每个线程处理ceil(size/gridsize)个数据templateunsignedintblockSize__global__voidreduce7(int*g_idata,int*g_odata,intsize){extern__shared__intsdata[];inttidthreadIdx.x;intiblockIdx.x*(blockSize*2)threadIdx.x;intgridSizeblockSize*2*gridDim.x;sdata[tid]0;while(isize){sdata[tid]g_idata[i];if(iblockSizesize){// 防止越界sdata[tid]g_idata[iblockSize];}igridSize;}__syncthreads();// 展开循环if(blockSize1024){if(tid512){sdata[tid]sdata[tid512];}__syncthreads();}if(blockSize512){if(tid256){sdata[tid]sdata[tid256];}__syncthreads();}if(blockSize256){if(tid128){sdata[tid]sdata[tid128];}__syncthreads();}if(blockSize128){if(tid64){sdata[tid]sdata[tid64];}__syncthreads();}// Warp 级优化if(tid32){// 强制转换为 volatile 指针传递确保读写可见性wrapReduceblockSize((volatileint*)sdata,tid);}if(tid0)g_odata[blockIdx.x]sdata[0];}这里我就没去调节block数了直接和上面的kernel用的一样的block数Reduce8——warp shuffle现在的CUDA有个shuffle的功能支持warp间不同线程做一些数据处理等价于把之前的取数相加赋值改成__shfl_down_sync函数调用templateunsignedintblockSize__global__voidreduce8(constint*__restrict__ g_idata,int*__restrict__ g_odata,intsize){extern__shared__intsdata[];unsignedinttidthreadIdx.x;unsignedintidxblockIdx.x*(blockSize*2)tid;intsum0;// 处理不能整除情况if(idxsize)sumg_idata[idx];if(idxblockSizesize)sumg_idata[idxblockSize];sdata[tid]sum;__syncthreads();//--------------------------------------------------// Block Reduction//--------------------------------------------------if(blockSize1024){if(tid512)sdata[tid]sdata[tid512];__syncthreads();}if(blockSize512){if(tid256)sdata[tid]sdata[tid256];__syncthreads();}if(blockSize256){if(tid128)sdata[tid]sdata[tid128];__syncthreads();}if(blockSize128){if(tid64)sdata[tid]sdata[tid64];__syncthreads();}//--------------------------------------------------// Warp Reduction using shuffle//--------------------------------------------------if(tid32){intvalsdata[tid];if(blockSize64)valsdata[tid32];for(intoffset16;offset0;offset1){val__shfl_down_sync(0xffffffff,val,offset);}if(tid0)g_odata[blockIdx.x]val;}}不过后面几种Reduce方法耗时都比较接近了参考文献Optimizing Parallel Reduction in CUDA