OpenCL 学习step by step (11) 数组求和(reduction)

迈克老狼2012

OpenCL 学习step by step (11) 数组求和(reduction)

<div class="postbody"><div id="cnblogs_post_body">     本篇教程中，我们学习一下如何用opencl有效实现数组求和，也就是通常所说的reduction问题。
       在程序中，我们设置workgroup size为256，kernel的输入、输出缓冲参数都用uint4的格式，这样我们原始求和的数组大小为256*4的倍数，数据类型为uint。我们设定每个workgroup处理处理512个uint，即2048个uint
       为了简便期间，我们输出数组长度定为4096，即需要2个workgruop来处理。
     
  kernel代码如下：
  __kernel void reduce(__global uint4* input, __global uint4* output, __local uint4* sdata)     
{      
    // 把数据装入lds      
    unsigned int tid = get_local_id(0);      
    unsigned int bid = get_group_id(0);      
    unsigned int gid = get_global_id(0);
      unsigned int localSize = get_local_size(0);     
    unsigned int stride = gid * 2;      
    sdata[tid] = input[stride] + input[stride + 1];
      barrier(CLK_LOCAL_MEM_FENCE);     
    // 在lds中进行reduction操作，得到数组求和的结果      
    for(unsigned int s = localSize >> 1; s > 0; s >>= 1)      
    {      
        if(tid < s)      
        {      
            sdata[tid] += sdata[tid + s];      
        }      
        barrier(CLK_LOCAL_MEM_FENCE);      
    }
     // 把一个workgroup计算的结果输出到输出缓冲，是一个uint4，还需要在host端再进行一次reduction过程      
    if(tid == 0) output[bid] = sdata[0];      
}
       在程序中，global和local的NDRange，我们都用一维的形式。下面以图的方式看下kernel代码是如何执行的：
  
          对第一个workgroup中的第一个thread的来说，它首先进行一次reduction操作，把两个uint4相加，放到lds(shared memory)中，然后再在lds中进行reduction操作，此时要从global memory中取数据，可以看出连续的thread访问连续的global memory，这时可以利用合并读写。
        申请的shared memory大小为groupsize*sizeof(uint4)，相加后uint4放入32bank的lds中，放置的方式应该是如下图所示，因为放入的是uint4，所以会放入连续的4个bank中(每个bank都是dword宽)，可见只能同时有8个thread访问lds，所以会有一定程序的bank conflit。从App profiler session，我们可以看到：
  
  
            接下来，kernel会通过一个for循环迭代执行reduction操作，求得一个workgroup中的uint4的和。
  迭代的第一次s=128,这时会执行如下图的两两相加，workgroup中同时执行的thread为128,thread local id大于等于128的线程都不会做什么事情，在每个循环的末尾，有一个barrier来同步所有thread，以便所有thread都完成这次循环后再进入下一次循环。
  
         第二次迭代的时候，只剩下前面128个uint4，workgroup中同时执行的thread为64。最后，当s=1时候，完成迭代reduction操作，然后把thread0（第一个thread)的结果输出。
       在host段，我们还要做一次相加操作，把不同workgroup得到的uint4，拆分成uint，并相加求得最终的结果。
  //在cpu reduction各个workgroup的结果以及uint4分量 reduction      
output = 0;      
for(int i = 0; i < numBlocks * VECTOR_SIZE; ++i)      
    output += outMapPtr;      
printf("gpu reduction result:%d\n", output);      
if(refOutput==output) printf("passed\n");
  
  程序执行后结果如下：
  
     

  完整的代码请参考：
  工程文件gclTutorial11
  代码下载：
  稍后提供