OpenCL：多个工作项将结果保存到同一个全局内存地址

Question

我正在尝试进行类似 reduce 的累积计算，其中需要根据特定条件存储 4 个不同的值。我的内核接收长数组作为输入，只需要存储 4 个值，这些值是 "global sums" 从输入的每个数据点获得的。例如，我需要存储满足特定条件的所有数据值的总和，以及满足该条件的数据点的数量。内核在下面以使其更清楚：

__kernel void photometry(__global float* stamp, 
                         __constant float* dark,
                         __global float* output)
{
int x = get_global_id(0);
int s = n * n;

if(x < s){
    float2 curr_px = (float2)((x / n), (x % n));
    float2 center = (float2)(centerX, centerY);
    int dist = (int)fast_distance(center, curr_px);
    if(dist < aperture){
        output[0] += stamp[x]-dark[x];
        output[1]++;
    }else if (dist > sky_inner && dist < sky_outer){
        output[2] += stamp[x]-dark[x];
        output[3]++;
    }
 }
}

所有未在内核中声明的值都是以前由宏定义的。 s 是输入数组 stamp 和 dark 的长度，它们是 nxn 矩阵扁平化为一维。

我得到了结果，但它们与我的 CPU 版本不同。当然我想知道：这是做我想做的事情的正确方法吗？我可以确定每个像素数据只被添加一次吗？我想不出任何其他方法来保存累积结果值。

Answer 1

您的情况需要原子操作，否则数据竞争会导致结果不可预测。

问题在这里：

output[0] += stamp[x]-dark[x];
output[1]++;

你可以想象同一个 wave 中的线程可能仍然遵循相同的步骤，因此，对于同一个 wave 中的线程来说可能是可以的。因为它们使用全局加载指令（广播）读取相同的 output[0] 值。然后，当它们完成计算并尝试将数据存储到相同的内存地址（output[0]）时，写入操作将被序列化。至此，您可能仍会得到正确的结果（对于同一 wave 中的工作项）。

但是，由于您的程序极有可能启动不止一波（在大多数应用程序中都是如此）。不同的波浪可能以未知的顺序执行；然后，当它们访问相同的内存地址时，行为变得更加复杂。例如，wave0 和 wave1 可能在任何其他计算发生之前就开始访问 output[0]，这意味着它们从 output[0] 获取相同的值；然后他们开始计算。计算后，他们将累积的结果保存到 output[0]；显然，其中一波的结果将被另一波覆盖，就好像只有写入内存的人后来被执行了一样。试想一下，您在实际应用程序中有更多的波浪，因此出现错误结果并不奇怪。

Answer 2

您可以在 O(log2(n)) 中同时执行此操作。一个概念想法：

你有 16 (1,2,3,4,5,6,7,8,9,10,11,12,13,14,15,16) 个输入，你想要的总和是这些输入同时进行。

你可以同时对2合1、4合3、6合5、8合7、10合9、12合11、14合13、16合15

然后你同时求和4合2，8合6，12合10，16合14

然后总是同时8个4个，16个10个

最后 16 中有 8 个

O(log2(n)) 在我们的案例中 4 篇文章.

中完成的所有事情

OpenCL：多个工作项将结果保存到同一个全局内存地址

OpenCL: multiple work items saving results to the same global memory address

gpu

gpgpu

opencl

pyopencl