Pull to refresh

CUDA: аспекты производительности при решении типичных задач

High performance *
Перед тем как начать переносить реализацию вычислительного алгоритма на видеокарту стоит задуматься — получим ли мы желаемый прирост производительности или только потеряем время. И несмотря на обещания производителей о сотнях GFLOPS, у современного поколения карт есть свои проблемы, о которых лучше знать заранее. Я не буду глубоко уходить в теорию и рассмотрю несколько существенных практических моментов и сформулирую некоторые полезные выводы.

Будем считать, что вы примерно разобрались, как работает CUDA и уже скачали стабильную версию CUDA Toolkit.

Я буду мучить теперь уже middle-end видеокарту GTX460 на Core Duo E8400.

Вызов функции

Да, если мы что-то хотим посчитать, то без вызова функции, выполняемой на карточке никак не обойтись. Для этого напишем простейшую тестовую функцию:

__global__ void stubCUDA(unsigned short* output)
{
  // the most valid function: yep, does nothing.
}

Напомню, что спецификатор __global__ позволяет выполнить функцию на GPU, вызвав ее с CPU:
cudaThreadSynchronize();
stubCUDA<<<GRID, THREADS>>>(0);
cudaThreadSynchronize();

Все вызовы функций по умолчанию асинхронны, поэтому вызовы cudaThreadSynchronize() необходимы для ожидания завершения вызванной функции.

Попробуем прогнать такой блок в цикле: получаем порядка 15000 вызовов в секунду для GRID=160, THREADS=96.

Скажем так, совсем не густо. Даже самая простейшая функция, которая ничего не делает, не может выполниться быстрее чем за 0.7 мс.

Первое предположение заключается в том, что бОльшая часть времени уходит на синхронизацию потоков и асинхронные вызовы отрабатывали бы значительно быстрее (хотя и применять их в конкретных задачах более специфично).

Проверим. Без синхронизации удалось запустить функцию 73100 раз в секунду. Результат, надо заметить, нисколько не впечатляющий.

И последний тест, запустим функцию с GRID=THREADS=1, казалось бы, это должно устранить накладные расходы на создание кучи потоков внутри карточки. Но это не так, получаем те же 73000-73500 вызовов в секунду.

Итак, мораль:
  • Абсолютно бессмысленно запускать на карточке те задачи, которые и на CPU считаются за миллисекунды.
  • Синхронизация потоков после вызова уменьшает производительность совсем незначительно на средних задачах.
  • Количество потоков и размер сетки не влияет на итоговое количество вызовов в секунду (разумеется это не так для «полезных» функций, которые делают что-то).

Доступ к памяти извне

Для того чтобы считать что-то полезное нам потребуются входные и выходные данные. Для этого надо понимать насколько быстро идет передача данных из/в видеокарту. Воспользуемся следующей функцией:
cudaMemcpy(data_cuda, image, data_cuda_size, cudaMemcpyHostToDevice);

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

Копируем большие блоки: как и в сторону cudaMemcpyHostToDevice, так и cudaMemcpyDeviceToHost получаем производительность порядка 2 Гбайт/c на больших блоках (более 100 мегабайт). В целом это очень даже неплохо.

Значительно хуже обстоят дела с совсем небольшими структурами. Передавая по 4 байта мы получаем не более 22000 вызовов в секунду, т.е. 88 кбайт/c.

Мораль:
  • Желательно группировать данные в большие блоки и передавать их одним вызовом функции cudaMemcpy.

Доступ к памяти изнутри

После того как мы передали данные на карточку, можно начинать с ними работать. Хочется оценить примерную скорость доступа к видеопамяти. Для этого напишем следующую функцию:
__global__ void accessTestCUDA(unsigned short* output, unsigned short* data, int blockcount, int blocksize)
{
  // just for test of max access speed: does nothing useful
  unsigned short temp;
  for (int i = blockIdx.x; i < blockcount; i += gridDim.x)
  {
    int vectorBase = i * blocksize;
    int vectorEnd = vectorBase + blocksize;

    for (int j = vectorBase + threadIdx.x; j < vectorEnd; j += blockDim.x)
    {
      temp = data[j];      
    }
  }
  output[0] = temp;
}


Здесь уже используются параметры GRID и THREADS, пока не буду объяснять зачем, но поверьте — все как следует. Придирчивые скажут, что результат пишется неправильно из-за отсутствия синхронизации, но нам-то он и не нужен.

Итак, получаем порядка 42 Гбайт/c для произвольного чтения. Вот это совсем неплохо.

Теперь модифицируем функцию, чтобы она копировала входные данные на выход. Бессмысленно, но позволяет оценить скорость записи в видеопамять (поскольку изменение совсем несложное, я не буду дублировать код).

Получаем порядка 30 Гбайт/с на ввод-вывод. Тоже неплохо.

Следует сделать поправку на то, что фактически мы использовали последовательный (с некоторыми отступлениями) доступ к памяти. Для произвольного цифры могут ухудшится до двух раз — но ведь и это не проблема?

Мораль:
  • За счет очень высокой скорости доступа к памяти на картах эффективно реализовывать алгоритмы, интенсивно ее использующие.

Арифметические операции

Совсем простые примеры опустим и сделаем что-то полезное. А именно — нормализацию изображения (pixel[t] := (pixel[t]-sub)*factor). Собственно код:
__global__ void normalizeCUDA(unsigned short* data, int blockcount, int blocksize, float sub, float factor)
{
  for (int i = blockIdx.x; i < blockcount; i += gridDim.x)
  {
    int vectorBase = i * blocksize;
    int vectorEnd = vectorBase + blocksize;

    for (int j = vectorBase + threadIdx.x; j < vectorEnd; j += blockDim.x)
    {
      register float d = (float)data[j];
      d = (d - sub) * factor;
      data[j] = (unsigned short)d;
    }
  }
}


Здесь используется аж три казалось бы затратных вычислительных процедуры: приведение к вещественным числам, ADDMUL и приведение к целым. На форумах пугают, что приведение целые-вещественные работает из рук вон плохо. Может быть это было верно для старых поколений карточек, но сейчас это не так.

Итоговая скорость обработки: 26 Гбайт/c. Три операции ухудшили производительность относительно прямого ввода-вывода всего на 13%.

Если внимательно посмотреть код, то нормализует он не совсем верно. Перед записью в целые числа, вещественное необходимо округлить, например функцией round(). Но не делайте так, и постарайтесь ее никогда не использовать!

round(d): 20 Гбайт/c, еще минус 23%.
(unsigned short)(d + 0.5): 26 Гбайт/с, собственно время в пределах погрешности измерений даже не поменялось.

Мораль:
  • Арифметические операции работают и правда быстро!
  • Для простейших алгоритмов обработки изображений можно рассчитывать на скорость в 10-20 Гбайт/c.
  • Лучше избегать использования функции round().

Логические операции

Попробуем оценить скорость работы логических операций, а заодно сделаем еще одно доброе дело: найдем минимум и максимум значений в массиве. Данный этап обычно предваряет нормализацию (и именно для этого и писался), но у нас все будет наоборот — т.к. он сложнее. Вот рабочий код:
__global__ void getMinMaxCUDA(unsigned short* output, unsigned short* data, int blockcount, int blocksize)
{
  __shared__ unsigned short sMins[MAX_THREADS];
  __shared__ unsigned short sMaxs[MAX_THREADS];
  
  sMins[threadIdx.x] = data[0];
  sMaxs[threadIdx.x] = data[0];
  
  for (int i = blockIdx.x; i < blockcount; i += gridDim.x)
  {
    int vectorBase = i * blocksize;
    int vectorEnd = vectorBase + blocksize;

    for (int j = vectorBase + threadIdx.x; j < vectorEnd; j += blockDim.x)
    {
      register unsigned short d = data[j];
      if (d < sMins[threadIdx.x])
        sMins[threadIdx.x] = d;
      if (d > sMaxs[threadIdx.x])
        sMaxs[threadIdx.x] = d;    
    }
  }

  __syncthreads();

  if (threadIdx.x == 0)
  {
    register unsigned short min = sMins[0];
    for (int j = 1; j < blockDim.x; j++)
      if (sMins[j] < min)
        min = sMins[j];
    if (min < output[0])
      output[0] = min;
  }
  
  if (threadIdx.x == 1)
  {
    register unsigned short max = sMaxs[0];
    for (int j = 1; j < blockDim.x; j++)
      if (sMaxs[j] > max)
        max = sMaxs[j];
    if (max > output[1])
      output[1] = max;
  }

  __syncthreads();
}


Здесь уже не обойтись без синхронизации потоков и shared memory.

Итоговая скорость: 29 Гбайт/c, даже быстрее нормализации.

Почему я объединил код минимума и максимума — обычно нужны оба, а вызовы по отдельности теряют время (см. первый абзац).

В общем, киньте камнем в того, кто сказал что на видеокартах плохо с условным операциями: искусственно удалось замедлить этот фрагмент практически в 2 раза, но для этого потребовалось увеличить глубину условий аж до 4! if () if () if () if () else if ()…

Мораль:
  • На современных карточках в общем-то не так и плохо с логикой, но следует избегать большой глубины вложенных условий.

Сложные структуры данных

Руководствуясь идеей о том что алгоритмы и структуры данных сильно связаны (хотя бы вспомнить Н. Вирта), следует проверить как же обстоят дела с некоторыми сложными структурами данных.

Вот тут возникает проблема, при передаче данных в функции мы можем использовать только два вида объектов — константные интегральные типы (чиселки) и ссылки на блоки видеопамяти.

Идея строить например деревья, основанные на ссылках накрывается сразу же:
  • мы не можем выделять память из функции, работающей на карточке;
  • любое выделение и копирование небольшого объема данных работает очень медленно (см. раздел 2).

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

Расплата за подобные ухищрения — необходимость применения двойной индексации:
    for (int j = vectorBase + threadIdx.x; j < vectorEnd; j += blockDim.x)
    {
      temp = data[index[j]+i];      
    }

Данный фрагмент работает со скоростью от 10 до 30 Гбайт/c в зависимости от наполнения и размеров индекса и данных. Использование памяти можно пытаться соптимизировать но даже в лучшем случае мы теряем 25% скорости доступа. Тройные индексы ведут себя еще хуже, теряя 40%-60% производительности.

Сегодня мы многое поняли

При грамотном использовании возможностей видеокарты можно получить небывалую производительность в задачах, скажем обработки изображений, звука, видео — везде где есть большие объемы данных, необходимость хитрой арифметики и отсутствие сложных структур данных.

Если топик вам понравится, то расскажу про то как обсчитывать на видеокарте несколько полезных объектов: Distance Map, морфологию изображений и поисковые индексы и покажу несколько интересных структур данных, которые работают достаточно быстро и не создают лишних проблем с синхронизацией.
Tags: высокопроизводительные вычислениявидеокартыCUDAоптимизацияbenchmark
Hubs: High performance
Total votes 86: ↑84 and ↓2 +82
Comments 67
Comments Comments 67

Popular right now