使用 OpenCL 在 2D 中嵌套循环求和

Question

我最近开始在 C++ 中使用 OpenCL，我正在尝试完全了解如何使用 2D 和 3D NDRange。我目前正在 OpenCL 中实现反距离加权，但我的问题很普遍。

下面是计算权重的串行函数，它由一个嵌套循环组成。

void computeWeights(int nGrids, int nPoints, double *distances, double *weightSum, const double p) {

    for (int i = 0; i < nGrids; ++i) {
        double sum = 0;
        for (int j = 0; j < nPoints; ++j) {
            double weight = 1 / pow(distances[i * nPoints + j], p);
            distances[i * nPoints + j] = weight;
            sum += weight;
        }
        weightSum[i] = sum;
    }
}

我想要的是使用 2D NDRange 来实现上述功能，第一个是在 nGrids 上，第二个是在 nPoints 上。但是，我不明白的是如何将权重的总和处理为 weightSum[i]。我知道我可能不得不以某种方式使用并行求和。

Answer 1

在使用 2D 全局工作空间调度内核时，OpenCL 会创建一个工作项网格。每个工作项都执行内核并在这两个维度上获得唯一的 ID。

(x,y)|________________________
     | (0,0) (0,1) (0,2) ...
     | (1,0) (1,1) (1,2)
     | (2,0) (2,1) (2,2)
     | ...

工作项也被分成组，并在这些工作组中获得唯一的 ID。例如。对于大小为 (2,2) 的工作组：

(x,y)|________________________
     | (0,0) (0,1) (0,0) ...
     | (1,0) (1,1) (1,0)
     | (0,0) (0,1) (0,0)
     | ...

您可以安排工作组，以便每个工作组执行缩减。

您的 SDK 可能有示例，并行缩减就是其中之一。

为了帮助您入门，这里有一个内核可以解决您的问题。它是最简单的形式，每行适用于一个工作组。

// cl::NDRange global(nPoints, nGrids);
// cl::NDRange local(nPoints, 1);
// cl::Local data(nPoints * sizeof (double));
kernel
void computeWeights(global double *distances, global double *weightSum, local double *data, double p)
{
    uint nPoints = get_global_size(0);

    uint j = get_global_id(0);
    uint i = get_global_id(1);
    uint lX = get_local_id(0);

    double weight = 1.0 / pow(distances[i * nPoints + j], p);

    distances[i * nPoints + j] = weight;

    data[lX] = weight;

    for (uint d = get_local_size(0) >> 1; d > 0; d >>= 1)
    {
        barrier(CLK_LOCAL_MEM_FENCE);
        if (lX < d)
            data[lX] += data[lX + d];
    }

    if (lX == 0) 
        weightSum[i] = data[0];
}

每一行工作项（即每个工作组）计算grid i 的权重（及其总和）。每个工作项计算一个权重，将其存储回distances，并将其加载到本地内存中。然后每个工作组对本地内存进行缩减，最后将结果存储在weightSum。