Беды с медленными скоростями в opencl - PullRequest
4 голосов
/ 15 мая 2019

Я пытаюсь использовать opencl в первый раз, цель состоит в том, чтобы вычислить argmin каждой строки в массиве.Поскольку операция в каждом ряду не зависит от других, я подумал, что это будет легко установить на видеокарту.

Мне кажется, что при использовании этого кода производительность хуже, чем когда я просто запускаю код на процессоре с внешним циклом, любая помощь будет признательна.

Вот код:

#pragma OPENCL EXTENSION cl_khr_fp64 : enable

int argmin(global double *array, int end)
{
  double minimum = array[0];
  int index;
  for (int j = 0; j < end; j++)
  {
    if (array[j] < minimum)
    {
      minimum = array[j];
      index = j;
    }
  }
  return index;
}

kernel void execute(global double *dist, global long *res, global double *min_dist)
{
  int row_size = 0;
  int i = get_global_id(0);

  int row_index = i * row_size;
  res[i] = argmin(&dist[row_index], row_size);
  min_dist[i] = dist[res[i] + row_index];

}

1 Ответ

0 голосов
/ 15 мая 2019

Комментаторы делают несколько правильных точек, но я постараюсь быть немного более конструктивным и организованным:

  1. Ваши данные состоят из double точных значений с плавающей точкой.В зависимости от вашего GPU, это само по себе может быть плохой новостью.Графические процессоры потребительского уровня обычно не оптимизированы для работы с double с, часто достигая лишь 1/32 или 1/16 пропускной способности по сравнению с float операциями с одинарной точностью.Многие графические процессоры профессионального уровня (Quadro, Tesla, FirePro, некоторые карты Radeon Pro), тем не менее, работают с ними, достигая пропускной способности 1/2 или 1/4 по сравнению с float.Поскольку вы выполняете только тривиальную арифметическую операцию (сравнение), и есть большая вероятность того, что в вашей среде выполнения преобладает доступ к памяти, это может также подойти и на потребительском оборудовании.
  2. Я полагаю, что row_size нефактически 0, это помогло бы узнать, что является истинным (типичным) значением, и является ли оно фиксированным, переменным за строкой или переменной за цикл, но одинаковым для каждой строки.В любом случае, если row_size не является очень маленьким, тот факт, что вы запускаете последовательный цикл for над ним, может сдерживать ваш код.
  3. Насколько велик размер вашей работы?Другими словами, сколько строк в вашем массиве (укажите типичный диапазон, если он меняется)?Если оно очень маленькое, вы увидите небольшую выгоду от параллелизма графических процессоров: графические процессоры имеют большое количество процессоров и могут планировать несколько потоков на процессор.Таким образом, для достижения достойного использования оборудования ваши рабочие элементы должны исчисляться сотнями или даже тысячами.
  4. Вы читаете очень большой массив из (предположительно) системной памяти и не выполняете каких-либо интенсивных операций с ней.Это означает, что ваше узкое место, как правило, будет на стороне доступа к памяти - для дискретных графических процессоров доступ к системной памяти должен проходить через PCIe, поэтому скорость этого соединения будет устанавливать верхнюю границу вашей производительности.Кроме того, ваш шаблон доступа к памяти далек от идеала для графических процессоров - вы обычно хотите, чтобы рабочие элементы считывали соседние ячейки памяти одновременно, поскольку блок памяти обычно извлекает 64 байта или более сразу.

Предложения по улучшению:

  • Профилирование.Если это вообще возможно, используйте инструменты профилирования вашего производителя графических процессоров, чтобы определить свои истинные узкие места.В противном случае мы просто догадываемся.
  • Для (4) - если это вообще возможно, старайтесь не перемещать большие объемы данных вокруг слишком большого количества.Если вы можете сгенерировать входные массивы на графическом процессоре, сделайте так, чтобы они никогда не покидали VRAM.
  • For (4) - Оптимизируйте доступ к памяти.У AMD, NVidia и Intel есть руководства по оптимизации OpenCL GPU, в которых объясняется, как это сделать.По сути, перестройте структуру данных или ядро ​​так, чтобы смежные рабочие элементы считывали смежные фрагменты памяти.В идеале вы хотите, чтобы рабочий элемент 0 считывал элемент массива 0, рабочий элемент 1 считывал элемент массива 1 и т. Д. Возможно, вам потребуется использовать локальную память для координации между рабочими элементами.Другим вариантом является чтение фрагментов данных векторного размера для каждого рабочего элемента.(например, каждый рабочий элемент читает double8 за один раз) В этом случае следите за выравниванием.
  • Для (2) & (3) - Если row_size не очень маленький (и фиксированный), попробуйтеразделите ваш цикл по нескольким рабочим элементам и координируйте, используя локальную память (алгоритмы сокращения) и атомарные операции в глобальной памяти.
  • For (1): Если вы оптимизировали все остальное, и профилирование говорит вам, что сравнение double s на потребительском оборудовании работает слишком медленно, либо проверьте, можете ли вы генерировать данные как float s без потери точности (это также сократит проблемы с пропускной способностью вашей памяти), либо проверьте, можете ли вы каким-то образом добиться большего успеха, например, с помощьюобрабатывая double как long и распаковывая вручную и сравнивая экспоненту и мантиссу, используя целочисленные операции.
Добро пожаловать на сайт PullRequest, где вы можете задавать вопросы и получать ответы от других членов сообщества.
...