Какое ядро ​​OpenCL 2 является оптимальным для суммирования чисел?

Question

C++ 17 представил ряд новых алгоритмов для поддержки параллельного выполнения, в частности, std:: является параллельной версией std:: накапливать, которая разрешает non-deterministic поведение для non-commutative операции, такие как сложение с плавающей запятой. Я хочу реализовать алгоритм сокращения, используя OpenCL 2.

У Intel есть пример, который использует OpenCL 2 work group функции ядра для реализации ядра std:: 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 для суммирования чисел?

Answer 1

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. Я не использовал эти новые атомы, поэтому воспринимайте их как зерно соли.