GPGPU OpenCL Reduction操作如何与group同步2015-05-11Reduction操作:规约操作就是由多个数生成一个数,如求最大值、最小值、向量点积、求和等操作,都属于这一类操作。有大量数据的情况下,使用GPU进行任务并行与数据并行,可以收到可好的效果。
group同步:OpenCL只提供了工作组内的各线程之间的同步机制,并没有提供所有线程的同步。提供组内item-work同步的方法:void barrier (cl_mem_fence_flags flags)
参数说明:cl_mem_fence_flags 可以取CLK_LOCAL_MEM_FENCE、CLK_GLOBAL_MEM_FENCE
函数说明:(1)一个work-group中所有work-item遇到barrier方法,都要等待其他work-item也到达该语句,才能执行后面的程序;(2)还可以组内的work-item对local or global memory的顺序读写操作。如下图中每个大框表示任务并行、每个group线程;框中的计算是数据并行、每个item-work线程:

作为练习,给出个完整的使用OpenCL计算整数序列求和,在数据并行中使用Local Memory 加速,group组内并行同步使用CLK_LOCAL_MEM_FENCE。
程序实例(整数序列求和):1.核函数(Own_Reduction_Kernels.cl):
__kernelvoidreduce(__global uint4* input, __global uint4* output, int NUM){NUM = NUM / 4;//每四个数为一个整体uint4。unsigned int tid = get_local_id(0);unsigned int localSize = get_local_size(0);unsigned int globalSize = get_global_size(0);uint4 res=(uint4){0,0,0,0};__local uint4 resArray[64];unsigned int i = get_global_id(0);while(i < NUM){res+=input[i];i+=globalSize;}resArray[tid]=res;//将每个work-item计算结果保存到对应__local memory中barrier(CLK_LOCAL_MEM_FENCE);// do reduction in shared memfor(unsigned int s = localSize >> 1; s > 0; s >>= 1) {if(tid < s) {resArray[tid] += resArray[tid + s];}barrier(CLK_LOCAL_MEM_FENCE);}// write result for this block to global memif(tid == 0) output[get_group_id(0)] = resArray[0];}