本篇教程中,我们学习一下如何用opencl有效实现数组求和,也就是通常所说的reduction问题。
在程序中,我们设置workgroup size为256,kernel的输入、输出缓冲参数都用uint4的格式,这样我们原始求和的数组大小为256*4的倍数,数据类型为uint。我们设定每个workgroup处理处理512个uint4,即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[i];
printf("gpu reduction result:%d\n", output);
if(refOutput==output) printf("passed\n");
程序执行后结果如下:
完整的代码请参考:
工程文件gclTutorial11
代码下载:
稍后提供