c - 如何在OpenCL中使用并行约简实现求和？

我正在尝试实现并行还原的内核。下面的代码有时会起作用，但我无法确定为什么有时会出错。

__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左右及以后的数字。欢迎任何帮助。谢谢。

最佳答案

我在想输入的大小不是2的幂时有事吗？


是。考虑一下当您尝试减少9个元素时会发生什么。假设您启动了一个包含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的幂。

关于c - 如何在OpenCL中使用并行约简实现求和？，我们在Stack Overflow上找到一个类似的问题：https://stackoverflow.com/questions/28736948/