OpenCL数据并行求和成一个变量

Question

是否可以使用 opencl 数据并行内核对大小为 N 的向量求和，而不使用部分求和技巧？

假设您可以访问 16 个工作项并且向量的大小为 16。难道不能让内核执行以下操作

    __kernel void summation(__global float* input, __global float* sum)
{
    int idx = get_global_id(0);

    sum[0] += input[idx];
}

当我尝试这样做时，sum 变量没有得到更新，而只是被覆盖。我读过一些关于使用障碍的文章，我尝试在上面的求和之前插入一个障碍，它确实以某种方式更新了变量，但它没有重现正确的总和。

Answer 1

让我试着解释一下为什么 sum[0] 被覆盖而不是更新。

在您的 16 个工作项的情况下，有 16 个线程同时运行ning。现在 sum[0] 是一个由所有线程共享的内存位置，行 sum[0] += input[idx] 是 16 个线程中的每一个同时运行。

现在指令 sum[0] += input[idx]（我认为）扩展执行 sum[0] 的读取，然后将 input[idx] 添加到 input[idx]，然后再将结果写回 sum[0]。

将出现 data race，因为多个线程正在读取和写入同一共享内存位置。所以可能会发生的是：

所有线程都可以在任何其他线程之前读取 sum[0] 的值将更新后的结果写回 sum[0]，在这种情况下，最终的 sum[0] 的结果将是线程的 input[idx] 的值哪个执行最慢。因为每次都不一样，如果你运行这个例子多次你应该看到不同的结果。
或者，一个线程可能执行得稍微慢一些，在这种情况下另一个线程可能已经将更新的结果写回 sum[0] 在这个慢线程读取 sum[0] 之前，在这种情况下将是一个使用多个线程的值的加法，但不是所有线程。

那么如何避免这种情况呢？

选项 1 - 原子（更差的选项）：

如果另一个线程正在共享内存位置上执行操作，您可以使用 atomics 强制所有线程阻塞，但这显然会导致性能损失，因为您正在使并行进程串行（并产生并行化的成本——例如在主机和设备之间移动内存并创建线程）。

选项 2 - 减少（更好的选项）：

最好的解决方案是减少数组，因为您可以最有效地使用并行性，并且可以提供 O(log(N)) 的性能。下面是使用 OpenCL 进行缩减的一个很好的概述：Reduction Example.

Answer 2

选项 3（也是最糟糕的）

    __kernel void summation(__global float* input, __global float* sum)
{
    int idx = get_global_id(0);
    for(int j=0;j<N;j++)
    {
        barrier(CLK_GLOBAL_MEM_FENCE| CLK_LOCAL_MEM_FENCE);
        if(idx==j)
         sum[0] += input[idx];
        else
         doOtherWorkWhileSingleCoreSums();

    }
}

使用主流 gpu，这应该将它们加起来和奔腾 mmx 一样慢。这就像在单个内核上进行计算并为其他内核提供其他工作，但速度较慢。

A cpu 设备在这种情况下可能比 gpu 更好。

OpenCL数据并行求和成一个变量

OpenCL data parallel summation into a variable

c++

arrays

sum

opencl