Суммирование вложенных циклов в 2D с использованием OpenCL

Недавно я начал работать с OpenCL на C++ и пытаюсь полностью понять, как использовать 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]. Я понимаю, что мне, возможно, придется каким-то образом использовать параллельное сокращение суммы.


c++ opencl reduction
person leadhoarse    schedule 06.06.2015    source источник

Ответы (1)


arrow_upward
2
arrow_downward

При отправке ядра с двумерным глобальным рабочим пространством OpenCL создает сетку рабочих элементов. Каждый рабочий элемент выполняет ядро ​​и получает уникальные идентификаторы в обоих этих измерениях.

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

Рабочие элементы также делятся на группы и получают уникальные идентификаторы в этих рабочих группах. Например. для рабочих групп размера (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.

person pAIgn10    schedule 06.06.2015
comment
Спасибо большое за помощь! Сокращение строк помогло мне понять, что происходит на самом деле. В данный момент я работаю над его расширением, так как невозможно создать слишком большое количество nPoints. Наличие нескольких рабочих групп на строку — вариант, но тогда мне нужно синхронизировать запись с weightSum[i], и я не решаюсь использовать специальную реализацию atomic_add для double. - person leadhoarse; 06.06.2015
comment
Синхронизация существует только внутри рабочей группы. Также атомарность работает только с целочисленными типами (по крайней мере, для сложения). Способ приблизиться к этому следующий: для начала загрузите 2 расстояния (двойные) для каждого рабочего элемента или, что еще лучше, загрузите 2 двойных 2 или двойных 4 для каждого рабочего элемента, чтобы каждый рабочий элемент вычислял 4 (или 8) весов. Затем, если nPoints все еще слишком велико, вы уменьшаете веса в блоках. Вы сохраняете сумму от каждой рабочей группы в глобальной памяти, а во втором ядре вы уменьшаете эти суммы, как обычно. - person pAIgn10; 06.06.2015