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

Question

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

假设您可以访问 16 个工作项并且您的 vector 大小为 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 个线程同时运行。现在 sum[0] 是所有线程共享的单个内存位置，行 sum[0] += input[idx] 由每个线程运行16 个线程，同时进行。

现在指令 sum[0] += input[idx](我认为)扩展执行读取 sum[0]，然后添加 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 .