OpenCL Reduction操作与group组内同步(barrier)操作的理解_综合

OpenCL Reduction操作与group同步
先解释一下什么是reduction操作和barrier操作

Reduction操作：规约操作就是由多个数生成一个数，如求最大值、最小值、向量点积、求和等操作，都属于这一类操作。
group同步：OpenCL只提供了工作组内的各线程之间的同步机制，并没有提供所有线程的同步。提供组内item-work同步的方法：

void barrier (cl_mem_fence_flags flags)

参数说明：
　　cl_mem_fence_flags 可以取CLK_LOCAL_MEM_FENCE、CLK_GLOBAL_MEM_FENCE

函数说明：
　　一个work-group中所有work-item遇到barrier方法，都要等待其他work-item也到达该语句，才能执行后面的程序；还可以组内的work-item对local or global memory的顺序读写操作。
　　在这里插入图片描述

代码示例讲解

__kernel void reduction_scalar(__global float* data, __local float* partial_sums, __global float* output) {
    int lid = get_local_id(0);int group_size = get_local_size(0);partial_sums[lid] = data[get_global_id(0)];barrier(CLK_LOCAL_MEM_FENCE);for(int i = group_size/2; i>0; i >>= 1) {
    if(lid < i) {
    partial_sums[lid] += partial_sums[lid + i];}barrier(CLK_LOCAL_MEM_FENCE);}if(lid == 0) {
    output[get_group_id(0)] = partial_sums[0];}
}

这是一个标准的reduction算法和goup同步的一个kernel 示例，很多的书上或者教程都是一这个示例来讲解的。我这边会详细的，从我的理解上讲解什么是reduction和barrier。
从第一张图片可以清楚的看到barrier 的作用，在组内，每个工作项想要执行下一条指令代码，必须等待同组内所有工作项都完成该操作。
在这里插入图片描述

   partial_sums[lid] = data[get_global_id(0)];barrier(CLK_LOCAL_MEM_FENCE);

该行代码是将所有全局工作项数据分配到组内的工作项后，开始下一步操作。实际的数据情况如下图。

   for(int i = group_size/2; i>0; i >>= 1) {
    if(lid < i) {
    partial_sums[lid] += partial_sums[lid + i];}barrier(CLK_LOCAL_MEM_FENCE);}

这段代码就是reduction,每次对半，循环右移一位。
在这里插入图片描述
每次循环，都会做一次barrier操作，只到i不满足条件，即突出，那么最后每组数组就会产生一个值。
这个就是典型的reduction. 输入为数组，输出为一个标量。

output[get_group_id(0)] = partial_sums[0];

代码最后，将数据从local memry 赋值到share memry。 cpu即可读到数据。