如何在 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

我认为当输入的大小不是 2 的幂时它有什么用？

是的。考虑一下当您尝试减少 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 次幂。