如何在 OpenCL 中正确地将结果从本地内存汇总到全局内存答案

【问题标题】：How to correclty sum results from local to global memory in OpenCL如何在 OpenCL 中正确地将结果从本地内存汇总到全局内存
【发布时间】：2014-07-03 16:18:37
【问题描述】：

我有一个 OpenCL 内核，其中每个工作组在本地内存中生成一个结果向量。然后，我需要将所有这些结果汇总到全局内存中，以便以后检索到主机。
为了测试这一点，我创建了以下内核代码：

//1st thread in each workgroup initializes local buffer
if(get_local_id(0) == 0){
    for(i=0; i<HYD_DIM; i++){
        pressure_Local[i] = (float2){1.0f, 0.0f};
    }
}

//wait for all workgroups to finish accessing any memory
barrier(CLK_GLOBAL_MEM_FENCE | CLK_LOCAL_MEM_FENCE);

/// sum all the results into global storage
for(i=0; i<get_num_groups(0); i++){

    //1st thread in each workgroup writes the group's local buffer to global memory
    if(i == get_group_id(0) && get_local_id(0) == 0){
        for(j=0; j<HYD_DIM; j++){
            pressure_Global[j] += pressure_Local[j];
            // barrier(CLK_GLOBAL_MEM_FENCE);
        }
    }

    //flush global memory buffers:
    barrier(CLK_GLOBAL_MEM_FENCE);
}

本质上，我希望全局内存中向量的所有元素都等于工作组的数量（在我的例子中是 128 个）。实际上，它们通常在 60 到 70 之间变化，并且结果因运行而异。
谁能告诉我我错过了什么，或者如何正确地做到这一点？

【问题讨论】：

标签： opencl reduce

【解决方案1】：

您无法使用 opencl 在不同工作组之间进行同步。 CLK_GLOBAL_MEM_FENCE 不能那样工作。它只保证将保持内存操作（由工作组访问）的顺序。请参阅OCL 1.2 spec 中的“6.12.8 同步函数”部分。

我会通过为每个工作组使用不同的全局内存块来解决您的问题。你将数据写入全局，你的内核就完成了。然后，如果您想将数据减少到单个块，您可以让另一个内核从全局读取数据，并将其与其他结果块合并。您可以根据需要进行任意多层合并，但最终合并必须由单个工作组完成。

搜索 gpu/opencl 缩减算法。这是一个不错的开始。 Case Study: Simple Reductions

【讨论】：