并行编程实战——CUDA编程的内核循环展开-育师

一、循环展开

开发经验相对丰富一些的程序员应该对循环展开并不陌生，特别是有过循环优化方面的经历的可能了解的会更深刻一些。循环是对CPU占用比较多的一种情况，如果在每次循环中再有大量的计算情况下，可能效果会更差。此时可以通过一定的方法手段缩减循环次数，而在每次循环中把多次的计算代码重复执行缩减的次数即可。其实这也有点时间和空间转换的味道。
循环展开可以手动展开也可以通过编译优化自动展开，这就看开发者想如何操作了。其下为示意代码：

//1:手动循环展开// 原始循环for(inti=0;i<16;i++){arr[i]=i*2;}// 展开5次for(inti=0;i<4;i+=4){arr[i]=i*2;arr[i+1]=(i+1)*2;arr[i+2]=(i+2)*2;arr[i+3]=(i+3)*2;arr[i+4]=(i+4)*2;}//2：自动循环展开// 原始循环for(inti=0;i<16;i++){arr[i]=i*2;}//编译器展开// GCC#pragmaGCC unroll4for(inti=0;i<n;i++){arr[i]=i*2;}

循环展开说着可能有点不好理解，但是看到代码估计一眼就看明白了。另外在学习C++的元编程时，还接触过类似于循环展开的例子，此处不展开了，有兴趣自己的去查看一下。

二、CUDA的循环展开

其实好的优化技术或方法基本都是通用的，只是对实际场景的匹配度大小罢了。CUDA线程中也是可以通过减少或消除循环控制来提高任务的运行效率。单纯的循环计算还是相对容易展开的，比较麻烦的可能是会有一些分支惩罚的情况。
如果开发者能够显式的控制循环迭代（比如使用#pargma unroll）或者CUDA本身可以识别循环的迭代次数，那么就可以进行分支循环的展开（即自动展开小循环）。当然，如果能够直接在代码中拆分循环会更方便，只不过这种情景可能很少遇到。看一下代码示意：

template<typenamegroup_t>__inline__ __device__floatwarp_test(group_tg,floatnum){//未指定参数，完全展开#pragmaunrollfor(intcount=g.size()/2;ofcountfset>0;count>>=1)num+=g.shfl_down(num,count);returnnum;}

不过还是那句话，循环展开如果应用场景不好，除了上面的分支惩罚，还有可能展开的代码占用了过多的寄存器导致效率的降低，也就是指令缓存未命中的惩罚。总之，开发者对于各种优化技术要灵活应用，千万不要教条。

三、例程

下面看一个循环展开优化的例程：

#include"cuda_runtime.h"#include"device_launch_parameters.h"#include<stdio.h>#include<iostream>#include<vector>#include<algorithm>// 手动展开__global__voidvecAddUnroll4(constfloat*A,constfloat*B,float*C,intn){inti=(blockIdx.x*blockDim.x+threadIdx.x)*4;if(i+3<n){// 处理4个元素C[i]=A[i]+B[i];C[i+1]=A[i+1]+B[i+1];C[i+2]=A[i+2]+B[i+2];C[i+3]=A[i+3]+B[i+3];}else{// 处理边界for(intj=0;j<4&&(i+j)<n;j++){C[i+j]=A[i+j]+B[i+j];}}}template<intUNROLL>__global__voidvecAddUnrollSet(constfloat*A,constfloat*B,float*C,intn){inttid=blockIdx.x*blockDim.x+threadIdx.x;inti=tid*UNROLL;#pragmaunrollfor(intj=0;j<UNROLL;j++){intid=i+j;if(id<n){C[id]=A[id]+B[id];}}}__global__voidvecAddUnroll(constfloat*A,constfloat*B,float*C,intn){inti=blockIdx.x*blockDim.x+threadIdx.x;if(i<n){floatsum=0.0f;// 编译器展开这个循环#pragmaunrollfor(intj=0;j<4;j++){intid=i+j*(blockDim.x*gridDim.x);if(id<n){sum+=A[id]+B[id];}}C[i]=sum;}}// pragma__global__voidvecAddUnrollSetPragma(constfloat*A,constfloat*B,float*C,intn){inti=blockIdx.x*blockDim.x+threadIdx.x;#pragmaunroll4for(intj=0;j<4;j++){intid=i*4+j;if(id<n){C[id]=A[id]+B[id];}}}intmain(){constintN=1<<20;//1Mconstsize_tsize=N*sizeof(float);std::vector<float>hA(N),hB(N),hC(N);std::generate(hA.begin(),hA.end(),[](){returnrand()/(float)RAND_MAX;});std::generate(hB.begin(),hB.end(),[](){returnrand()/(float)RAND_MAX;});float*dA,*dB,*dC;cudaMalloc(&dA,size);cudaMalloc(&dB,size);cudaMalloc(&dC,size);// 拷贝数据到设备cudaMemcpy(dA,hA.data(),size,cudaMemcpyHostToDevice);cudaMemcpy(dB,hB.data(),size,cudaMemcpyHostToDevice);// 设置执行配置intbSize=256;intgSize=(N+bSize-1)/bSize;vecAddUnroll4<<<gSize/4,bSize>>>(dA,dB,dC,N);cudaDeviceSynchronize();vecAddUnrollSet<4><<<gSize/4,bSize>>>(dA,dB,dC,N);cudaDeviceSynchronize();// pragma unrollvecAddUnroll<<<gSize,bSize>>>(dA,dB,dC,N);cudaDeviceSynchronize();vecAddUnrollSetPragma<<<gSize,bSize>>>(dA,dB,dC,N);cudaDeviceSynchronize();// 清理cudaFree(dA);cudaFree(dB);cudaFree(dC);return0;}