c - 如何在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 个元素求和，因此减少是不正确的。一个简单的解决方案是用足够的零填充输入数据，以使工作组大小为紧邻的下一个 2 的幂。