是否有可能在OpenCL中并行运行求和计算?

Is it possible to run the sum computation in parallel in OpenCL?

本文关键字：运行 求和 计算 并行 有可能 OpenCL 是否      更新时间：2023-10-16

我是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];
}

这个代码是错误的。如果有人回答我，我将非常感激，如果它是理论上可能并行运行这样的任务，如果它是-如何!

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

数据依赖将导致线程必须等待彼此，创建争用，这是您想要避免的，以获得真正的并行化。

一种方法是将数组分成N个数组，每个数组包含原始数组的一部分，然后用每个不同的数组调用OpenCL内核函数。

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

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

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

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

。您有N个元素，并且您的工作组的大小是WS = 64。我假设N是2*WS的倍数(这很重要，一个工作组计算2*WS元素的总和)。然后你需要运行内核指定:

globalSizeX = 2*WS*(N/(2*WS));
因此，sum数组将包含2*WS元素的部分和。(例如sum[1] -将包含索引从2*WS到4*WS-1的元素的和)

如果您的globalSizeX是2*WS或更少(这意味着您只有一个工作组)，那么您就完成了。只使用sum[0]作为结果。如果没有，您需要重复这个过程，这次使用sum array作为输入数组并输出到其他数组(创建2个数组并在它们之间来回切换)。以此类推，直到你只有一个工作组。

也搜索Hilli Steele/Blelloch并行算法。这篇文章可能也很有用

下面是实际的例子:

__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
}
相关文章：