是否可以在 OpenCL 中并行运行总和计算？答案

【问题标题】：Is it possible to run the sum computation in parallel in OpenCL?是否可以在 OpenCL 中并行运行总和计算？
【发布时间】：2013-02-18 11:00:19
【问题描述】：

我是 OpenCL 的新手。但是，我了解 C/C++ 基础知识和 OOP。我的问题如下：是否可以并行运行求和计算任务？理论上可行吗？下面我将描述我尝试做的事情：

任务是，例如：

double* values = new double[1000]; //let's pretend it has some random values inside
double sum = 0.0;

for(int i = 0; i < 1000; i++) {
    sum += values[i];
}

我在 OpenCL 内核中尝试做的事情（我觉得这是错误的，因为它可能同时从不同的线程/任务访问相同的“sum”变量）：

__kernel void calculate2dim(__global float* vectors1dim,
                            __global float output,
                            const unsigned int count) {
    int i = get_global_id(0);
    output += vectors1dim[i];
}

此代码错误。如果有人回答我是否理论上可以并行运行这些任务，如果可以的话，我将非常感激 - 如何！

【问题讨论】：

这是一个经典的归约问题。查看here 逐步解释为多核架构优化此过程（它是CUDA，但原理完全相同，可能除了关于模板的部分）。尽管有关该主题的更多介绍性材料可能会更有帮助，但我将其留给正确的答案。
非常感谢！现在我知道这是一个普遍的问题，并且会学习如何解决它！

标签： c++ c parallel-processing opencl

【解决方案1】：

如果您想以并行方式对数组的值求和，则应确保减少争用并确保线程之间没有数据依赖关系。

数据依赖会导致线程必须相互等待，从而产生争用，这是你想要避免的，以获得真正的并行化。

您可以这样做的一种方法是将您的数组拆分为 N 个数组，每个数组都包含原始数组的一些子部分，然后使用每个不同的数组调用您的 OpenCL 内核函数。

最后，当所有内核都完成了艰苦的工作时，您可以将每个数组的结果汇总为一个。 CPU 可以轻松完成此操作。

关键是在每个内核中完成的计算之间没有任何依赖关系，因此您必须相应地拆分数据和处理。

我不知道您的数据是否与您的问题有任何实际依赖关系，但这是您自己弄清楚的。

【讨论】：

感谢您的回答。可能，我应该将我的数组分成几个单独的！您知道将二维数组（如 double[][]）传递给内核的任何方法吗？因为pointet-to-pointer不能用作函数参数。
您不需要传递二维数组，只需传递您的缓冲区并像这样 myarr[y*WIDTH + x] 那样寻址即可。

【解决方案2】：

我提供的供参考的代码应该可以完成这项工作。

例如你有 N 个元素，你的工作组的大小是 WS = 64。我假设 N 是 2*WS 的倍数（这很重要，一个工作组计算 2*WS 元素的总和）。然后你需要运行内核指定：

globalSizeX = 2*WS*(N/(2*WS));

因此，sum 数组将包含 2*WS 元素的部分总和。（例如 sum[1] - 将包含索引从 2*WS 到 4*WS-1 的元素的总和）。 p>

如果您的 globalSizeX 为 2*WS 或更小（这意味着您只有一个工作组），那么您就完成了。只需使用 sum[0] 作为结果。如果不是 - 您需要重复过程，这次使用 sum 数组作为输入数组并输出到其他数组（创建 2 个数组并在它们之间进行乒乓球）。以此类推，直到您只有一个工作组。

同时搜索 Hilli Steele / Blelloch 并行算法。 This 文章也很有用

这是一个实际的例子：

__kernel void par_sum(__global unsigned int* input, __global unsigned int* sum)
{
    int li = get_local_id(0);
    int groupId = get_group_id(0);

    __local int our_h[2 * get_group_size(0)];
    our_h[2*li + 0] = hist[2*get_group_size(0)*blockId + 2*li + 0];
    our_h[2*li + 1] = hist[2*get_group_size(0)*blockId + 2*li + 1];

    // sweep up
    int width = 2;
    int num_el = 2*get_group_size(0)/width;
    int wby2 = width>>1;

    for(int i = 2*BLK_SIZ>>1; i>0; i>>=1)
    {

        barrier(CLK_LOCL_MEM_FENCE);

        if(li < num_el)
        {
            int idx = width*(li+1) - 1;
            our_h[idx] = our_h[idx] + our_h[(idx - wby2)];
        }

        width<<=1;
        wby2 = width>>1;
        num_el>>=1;
    }

        barrier(CLK_LOCL_MEM_FENCE);

    // down-sweep
    if(0 == li)
        sum[groupId] = our_h[2*get_group_size(0)-1]; // save sum
}

【讨论】：