Параллельное суммирование данных OpenCL в переменную

Question

Можно ли использовать параллельное ядро ​​данных opencl для суммирования вектора размера N без выполнения частичного трюка?

Говорят, что если у вас есть доступ к 16 рабочих элементов и ваш вектор имеет размер 16. Не будет ли это не возможно, чтобы просто иметь ядро, выполнив следующие действия

__kernel void summation(__global float* input, __global float* sum) 
{ 
    int idx = get_global_id(0); 

    sum[0] += input[idx]; 
} 

Когда я попытался это , переменная sum не обновляется, а только перезаписывается. Я прочитал что-то об использовании барьеров, и я попытался вставить барьер перед суммированием выше, он каким-то образом обновляет переменную, но не воспроизводит правильную сумму.

Answer 1

Позвольте мне объяснить, почему sum[0] перезаписан, а не обновлен.

В вашем случае из 16 рабочих элементов имеется 16 потоков, которые работают одновременно. Теперь sum[0] представляет собой единую ячейку памяти, которая совместно используется всеми потоками, а строка sum[0] += input[idx] запускается каждым из 16 потоков одновременно.

Теперь инструкция sum[0] += input[idx] (я думаю) расширяет, выполняет чтение sum[0], а затем добавляет input[idx], прежде чем записывать результат обратно до sum[0].

Будет отображаться data race, поскольку несколько потоков считывают и записывают в одно и то же место общей памяти. Так что может случиться:

• Все нити могут считывать значение sum[0] перед любым другим потоком записывает свой обновленный результат обратно sum[0], в этом случае конечного результат sum[0] бы значение input[idx] из нить , которая выполняла самую медленную. Поскольку это будет отличаться каждый раз, , если вы запустите пример несколько раз, вы должны увидеть разные результаты .
• Или, один поток может выполнять несколько медленнее, и в этом случае другой поток может уже написал обновленный результат обратно sum[0] перед этой медленной нитью читает sum[0], и в этом случае будет прибавление, используя значения более одного потока, но не все темы.

Итак, как вы можете избежать этого?

Вариант 1 - Atomics (хуже вариант):

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

Вариант 2 - уменьшение (лучший вариант):

Лучшим решением было бы уменьшить массив, так как вы можете наиболее эффективно использовать параллелизм, и может дать O производительность (журнал (N)). Вот хороший обзор сокращения с помощью OpenCL: Reduction Example.

Answer 2

Вариант 3 (хуже всего)

__kernel void summation(__global float* input, __global float* sum) 
{ 
    int idx = get_global_id(0); 
    for(int j=0;j<N;j++) 
    { 
     barrier(CLK_GLOBAL_MEM_FENCE| CLK_LOCAL_MEM_FENCE); 
     if(idx==j) 
     sum[0] += input[idx]; 
     else 
     doOtherWorkWhileSingleCoreSums(); 

    } 
} 

с использованием основной GPU, это должно суммировать все они так медленно, как Pentium MMX. Это похоже на вычисление на одном ядре и предоставление другим ядрам других заданий, но медленнее.

Устройство cpu может быть лучше, чем gpu для такого рода.