如何在 OpenCL 中使用并行归约实现求和？

Question

我正在尝试实现一个进行并行缩减的内核。下面的代码有时会起作用，我无法确定为什么它有时会出错。

__kernel void summation(__global float* input, __global float* partialSum, __local float *localSum){
int local_id = get_local_id(0);
int workgroup_size = get_local_size(0);
localSum[local_id] = input[get_global_id(0)];

for(int step = workgroup_size/2; step>0; step/=2){
    barrier(CLK_LOCAL_MEM_FENCE);

    if(local_id < step){
    localSum[local_id] += localSum[local_id + step];
    }
}
if(local_id == 0){
    partialSum[get_group_id(0)] = localSum[0];
}}

本质上我是对每个工作组的值求和并将每个工作组的总和存储到 partialSum 中，最后的求和在主机上完成。下面是设置求和值的代码。

size_t global[1];
size_t local[1];

const int DATA_SIZE = 15000;
float *input = NULL;
float *partialSum = NULL;
int count = DATA_SIZE;

local[0] = 2;
global[0] = count;
input = (float *)malloc(count * sizeof(float));
partialSum = (float *)malloc(global[0]/local[0] * sizeof(float));

int i;
for (i = 0; i < count; i++){
    input[i] = (float)i+1;
}

我在想，当输入的大小不是 2 的幂时，这与它有关吗？我注意到它开始在 8000 左右及以后的数字上消失。欢迎任何帮助。谢谢。

Answer 1

I'm thinking it has something to do when the size of the input is not a power of two?

是的。考虑一下当您尝试减少 9 个元素时会发生什么。假设您启动 1 个工作组，其中包含 9 个工作项：

for (int step = workgroup_size / 2; step > 0; step /= 2){
    // At iteration 0: step = 9 / 2 = 4
    barrier(CLK_LOCAL_MEM_FENCE);

    if (local_id < step) {
        // Branch taken by threads 0 to 3
        // Only 8 numbers added up together! 
        localSum[local_id] += localSum[local_id + step];
    }
}

您从来没有对第 9 个元素求和，因此归约是不正确的。一个简单的解决方案是用足够多的零填充输入数据，使工作组大小成为紧邻的下一个二次幂。

如何在 OpenCL 中使用并行归约实现求和？

How to implement summation using parallel reduction in OpenCL?

c

parallel-processing

opencl