Перед тем как начать переносить реализацию вычислительного алгоритма на видеокарту стоит задуматься — получим ли мы желаемый прирост производительности или только потеряем время. И несмотря на обещания производителей о сотнях 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, морфологию изображений и поисковые индексы и покажу несколько интересных структур данных, которые работают достаточно быстро и не создают лишних проблем с синхронизацией.