Какое ядро OpenCL 2 является оптимальным для суммирования чисел?
C ++ 17 представил ряд новых алгоритмов для поддержки параллельного выполнения, в частности станд :: уменьшить это параллельная версия станд :: аккумулируют что позволяет non-deterministic поведение для non-commutative операции, такие как сложение с плавающей запятой. Я хочу реализовать алгоритм сокращения, используя OpenCL 2.
У Intel есть пример Вот который использует OpenCL 2 work group функции ядра для реализации станд :: exclusive_scan Ядро OpenCL 2. Ниже приведено ядро для суммирования чисел на основе Intel exclusive_scan пример:
kernel void sum_float (global float* sum, global float* values)
{
float sum_val = 0.0f;
for (size_t i = 0u; i < get_num_groups(0); ++i)
{
size_t index = get_local_id(0) + i * get_enqueued_local_size(0);
float value = work_group_reduce_add(values[index]);
sum_val += work_group_broadcast(value, 0u);
}
sum[0] = sum_val;
}
Ядро выше работает (или кажется!). Тем не мение, exclusive_scan требуется work_group_broadcast функция для передачи последнего значения одного work group к следующему, тогда как это ядро требует только добавления результата work_group_reduce_add к sum_valтак atomic add более уместно.
OpenCL 2 обеспечивает atomic_int который поддерживает atomic_fetch_add, Целочисленная версия вышеупомянутого ядра с использованием atomic_int:
kernel void sum_int (global int* sum, global int* values)
{
atomic_int sum_val;
atomic_init(&sum_val, 0);
for (size_t i = 0u; i < get_num_groups(0); ++i)
{
size_t index = get_local_id(0) + i * get_enqueued_local_size(0);
int value = work_group_reduce_add(values[index]);
atomic_fetch_add(&sum_val, value);
}
sum[0] = atomic_load(&sum_val);
}
OpenCL 2 также предоставляет atomic_float но это не служба поддержки atomic_fetch_add,
Каков наилучший способ реализации ядра OpenCL2 для суммирования чисел?
Решение
kernel void sum_float (global float* sum, global float* values)
{
float sum_val = 0.0f;
for (size_t i = 0u; i < get_num_groups(0); ++i)
{
size_t index = get_local_id(0) + i * get_enqueued_local_size(0);
float value = work_group_reduce_add(values[index]);
sum_val += work_group_broadcast(value, 0u);
}
sum[0] = sum_val;
}
это имеет условие состязания для записи данных в элемент с нулевым индексом суммы, все рабочие группы выполняют одно и то же вычисление, которое делает это O (N * N) вместо O (N) и занимает более 1100 миллисекунд для завершения суммы массива 1M-элемента ,
Для того же массива элементов 1-M это (global = 1M, local = 256)
kernel void sum_float2 (global float* sum, global float* values)
{
float sum_partial = work_group_reduce_add(values[get_global_id(0)]);
if(get_local_id(0)==0)
sum[get_group_id(0)] = sum_partial;
}
с последующим (global = 4k, local = 256)
kernel void sum_float3 (global float* sum, global float* values)
{
float sum_partial = work_group_reduce_add(sum[get_global_id(0)]);
if(get_local_id(0)==0)
values[get_group_id(0)] = sum_partial;
}
делает то же самое в течение нескольких миллисекунд, кроме третьего шага. Первый получает каждую группу сумм в свой элемент, связанный с идентификатором группы, а второе ядро суммирует их в 16 значений, и эти 16 значений могут быть легко суммированы ЦП (микросекунды или меньше) (в качестве третьего шага).
Программа работает так:
values: 1.0 1.0 .... 1.0 1.0
sum_float2
sum: 256.0 256.0 256.0
sum_float3
values: 65536.0 65536.0 .... 16 items total to be summed by cpu
если вам нужно использовать атомику, вы должны делать это как можно реже. Простейшим примером может быть использование локальной атомики для суммирования многих значений по каждой группе, а затем выполнение последнего шага с использованием одной глобальной атомарной функции на группу для добавления всех. У меня пока нет готовой установки C ++ для OpenCL, но я полагаю, что атомы OpenCL 2.0 лучше, когда вы используете несколько устройств с тем же ресурсом памяти (возможно, потоковым режимом или в SVM) и / или ЦПУ используя функции C ++ 17. Если у вас нет нескольких устройств, работающих в одной и той же области одновременно, то я полагаю, что эти новые атомы могут быть микрооптимизацией только поверх уже работающих атомов OpenCL 1.2. Я не использовал эти новые атомы, поэтому воспринимайте их как зерно соли.
Другие решения
Других решений пока нет …