Проблема производительности OpenCL на GPU - PullRequest
Новая
       10

Проблема производительности OpenCL на GPU

3 голосов
/ 13 марта 2019

Я использую OpenCL для оптимизации кода в Raspberry Pi GPU (Videocore IV). Я использую реализацию VC4CL, которая предлагает максимальный размер рабочей группы 12.

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

Например, для следующего ядра: 

#define GLOBAL_SIZE 12
#define LOCAL_SIZE  1
#define WIDTH       12*12*12
#define NTIMES      1000

__attribute__ ((reqd_work_group_size(LOCAL_SIZE, LOCAL_SIZE, LOCAL_SIZE)))
__kernel void int_sum(global const uint* A, global const uint* B, global uint* C)
{
  int g_id0 = get_global_id(0);
  int g_id1 = get_global_id(1);
  int g_id2 = get_global_id(2);

  int index = g_id0 + g_id1 * GLOBAL_SIZE + g_id2 * GLOBAL_SIZE * GLOBAL_SIZE;

  for(int k = 0; k < NTIMES; k++)
    C[index + k * WIDTH] = A[index + k * WIDTH] + B[index + k * WIDTH];
}

, где суммируются два массива из 1e6 позиций, производительность процессора намного лучше ... Я пытался изменить рабочую группу на одномерное, а также использовать другие комбинации, такие как (6x6x6 -> глобальный размер, 2x2x2 -> локальный размер). Любой намек на то, что я могу делать неправильно?

Заранее спасибо.

Ответы [ 2 ]

1 голос
/ 14 марта 2019

Кроме того, что говорили все остальные в комментариях, согласно автору реализации OpenCL RPi,"Скорость доступа к памяти" для GPU (копия памяти CPU-GPU?) Намного медленнее, чем для CPU.Таким образом, «арифметически легкое» ядро, такое как сумма массива, будет ограничено пропускной способностью памяти и может быть намного медленнее, чем на процессоре.Кроме того, теоретические GFlops для GPU не намного выше, чем для GPU (24 против 6).

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

1 голос
/ 13 марта 2019

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

  • Это целочисленный код ALU, а не операции с плавающей запятой,Многие графические процессоры вообще не оптимизированы для этого.
  • Я бы не стал полагаться на то, что компилятор оптимизирует вычисления смещения массива;особенно глупый компилятор может выдавать 3 целочисленных умножения для C[index + k * WIDTH] = A[index + k * WIDTH] + B[index + k * WIDTH]; на каждой итерации цикла.Я бы сохранял смещение в переменной и добавлял к нему на каждой итерации, умножение не требуется.
  • 1000-итерация цикла обычно выделяется как потенциальный источник лучшего параллелизма.Многие графические процессоры плохо работают с долго работающими ядрами.
  • Шаблоны доступа к памяти кажутся неоптимальными.Попробуйте расположить так, чтобы смежные рабочие элементы в группе обращались к смежным ячейкам памяти.Локальный размер 2x2x2 кажется особенно плохим выбором для этого.Вы пробовали 12x1x1?
  • Почему вы даже так упорядочиваете рабочие элементы?Похоже, вы буквально рассчитываете C[i] = A[i] + B[i] для i = 0..1000*12*12*12.Как насчет написания вашего ядра именно так и представления 1728000 рабочих элементов в одном измерении?Это экономит на всех сложных расчетах индекса.

Если вы можете получить какие-либо отзывы от драйверов о том, с чем связан графический процессор (ALU, загрузка памяти, планирование потоков и т. Д.), Это очень поможет при выборе места для поиска.способы ускорить его.

...