OpenCL 数据并行求和到一个变量中

OpenCL data parallel summation into a variable

本文关键字：一个 变量 数据 并行 求和 OpenCL      更新时间：2023-10-16

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

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

吗 
    __kernel void summation(__global float* input, __global float* sum)
{
    int idx = get_global_id(0);
    sum[0] += input[idx];
}

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

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

在 16 个

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

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

当多个线程读取和写入同一共享内存位置时，将会出现数据竞赛。所以可能发生的情况是：

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

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

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

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

选项

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

最好的解决方案是减少数组，因为您可以最有效地使用并行性，并且可以提供 O（log（N）） 性能。以下是使用 OpenCL 进行缩减的良好概述：缩减示例。

选项 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一样慢。这就像在单个内核上进行计算，并为其他内核提供其他工作，但速度较慢。

对于这种设备，CPU设备可能比gpu更好。

相关文章：