Способ выполнения окончательной суммы с уменьшением

Question

Я занимаюсь продолжением своей первой проблемы, поясненной на this link.

Напоминаю, что я хотел бы применить метод, который может выполнять множественные сокращения суммы с помощью OpenCL (мое устройство GPU поддерживает только OpenCL 1.2). Мне нужно вычислить суммарное уменьшение массива для проверки критерия сходимости для каждой итерации основного контура,

В настоящее время я сделал версию только для одного уменьшения суммы (например, одной итерации ). В этой версии и для простоты я использовал последовательный цикл ЦП, чтобы вычислить сумму каждой частичной суммы и получить окончательное значение суммы.

Из ваших рекомендаций в моем прецеденте моя проблема в том, что я не знаю, как выполнить окончательную сумму, вызвав вторую функцию NDRangeKernel (т. Е. Выполняя второй раз код ядра).

Действительно, со вторым вызовом я всегда сталкиваюсь с той же проблемой для получения суммы частичных сумм (сам рассчитывается при первом вызове NDRangeKernel): это кажется рекурсивной проблемой.

Давайте возьмем пример из приведенного выше рисунка: если размер входного массива 10240000 и WorkGroup size16 это, мы получаем 10000*2^10/2^4 = 10000*2^6 = 640000 WorkGroups.

Итак, после первого звонка я получаю 640000 partial sums: как бороться с окончательной суммой всех этих частичных сумм? Если я вызову еще раз код ядра с, например, WorkGroup size = 16 и глобальным size = 640000, я получу nWorkGroups = 640000/16 = 40000 partial sums, поэтому мне нужно еще раз вызвать код ядра и повторить этот процесс до nWorkGroups < WorkGroup size.

Может быть, я не очень хорошо понимаю, второй этап, в основном это часть кода ядра от «снижения двухступенчатой» (on this link, I think this is the case of searching for minimum into input array)

__kernel 
void reduce(__global float* buffer, 
      __local float* scratch, 
      __const int length, 
      __global float* result) { 

    int global_index = get_global_id(0); 
    float accumulator = INFINITY; 
    // Loop sequentially over chunks of input vector 
    while (global_index < length) { 
    float element = buffer[global_index]; 
    accumulator = (accumulator < element) ? accumulator : element; 
    global_index += get_global_size(0); 
    } 

    // Perform parallel reduction 
    ... 

Если кто-то может объяснить, что это выше фрагмент кода ядра код делает.

Есть ли связь со вторым этапом редукции, т. Е. Окончательная сумма?

Не стесняйтесь спрашивать меня, если вы не поняли мою проблему.

Благодаря

Answer 1

Как уже упоминалось в комментариях: Заявление

, если размер входного массива 10240000 и размер WorkGroup составляет 16, мы получаем 10000 * 2^10/2^4 = 10000 * 2^6 = 640000 рабочих групп.

Неправильная версия. Вы можете выбрать «произвольный» размер рабочей группы и «произвольное» количество рабочих групп. Номера, которые можно выбрать здесь, могут быть настроены для целевого устройства. Например, устройство может иметь определенный размер локальной памяти.Это может быть запрошена с clDeviceGetInfo:

cl_ulong localMemSize = 0; 
clDeviceGetInfo(device, CL_DEVICE_LOCAL_MEM_SIZE, 
    sizeof(cl_ulong), &localMemSize, nullptr); 

Это может быть использовано для вычисления размера локальной рабочей группы, принимая во внимание тот факт, что каждая рабочая группа будет требовать

sizeof(cl_float) * workGroupSize 

байт локальной памяти.

Аналогично, количество рабочих групп может быть получено из других параметров, специфичных для устройства.

---

Ключевой момент в отношении самого сокращения, что размер рабочей группы не ограничивает размер массива, который может быть обработан. Я также имел некоторые трудности с пониманием алгоритма в целом, так что я пытался объяснить это здесь, в надежде, что несколько изображений может стоить тысячи слов:

Как вы можете видеть, количество рабочие группы и размер рабочей группы являются фиксированными и независимыми от длины входного массива. Несмотря на то, что я использую 3 рабочих группы с размером 8 в примере (давая глобальный размер 24), массив длиной 64 может быть обработанный. Это в основном связано с первым циклом, который просто проходит через входной массив с размером шага, равным глобальному размеру работы (здесь 24). Результатом будет один накопленное значение для каждого из 24 потоков. Затем эти уменьшаются параллельно.