Как оптимизировать игру жизни Конвея для CUDA? - PullRequest
Новая
       6

Как оптимизировать игру жизни Конвея для CUDA?

11 голосов
/ 03 января 2011

Я написал это ядро ​​CUDA для игры жизни Конвея: 

__global__ void gameOfLife(float* returnBuffer, int width, int height) {  
    unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;  
    unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;  
    float p = tex2D(inputTex, x, y);  
    float neighbors = 0;  
    neighbors += tex2D(inputTex, x+1, y);  
    neighbors += tex2D(inputTex, x-1, y);  
    neighbors += tex2D(inputTex, x, y+1);  
    neighbors += tex2D(inputTex, x, y-1);  
    neighbors += tex2D(inputTex, x+1, y+1);  
    neighbors += tex2D(inputTex, x-1, y-1);  
    neighbors += tex2D(inputTex, x-1, y+1);  
    neighbors += tex2D(inputTex, x+1, y-1);  
    __syncthreads();  
    float final = 0;  
    if(neighbors < 2) final = 0;  
    else if(neighbors > 3) final = 0;  
    else if(p != 0) final = 1;  
    else if(neighbors == 3) final = 1;  
    __syncthreads();  
    returnBuffer[x + y*width] = final;  
}

Я ищу ошибки / оптимизации. Параллельное программирование является для меня совершенно новым, и я не уверен, правильно ли я это понял.

Остальная часть представляет собой memcpy из входного массива для 2D-текстуры inputTex, привязанного к массиву CUDA. Выходные данные записываются из глобальной памяти на хост, а затем обрабатываются.

Как видите, поток имеет дело с одним пикселем. Я не уверен, что это самый быстрый способ, так как некоторые источники предлагают сделать строку или более на поток. Если я правильно понимаю, NVidia сами говорят, что чем больше тем, тем лучше. Я хотел бы получить совет по этому вопросу от кого-то с практическим опытом.

Ответы [ 3 ]

11 голосов
/ 03 января 2011

Мои два цента.

Вероятно, все это будет связано с задержкой связи между мультипроцессорами и памятью графического процессора.У вас есть код, для выполнения которого требуется около 30-50 тактов, и он генерирует как минимум 3 обращения к памяти, каждый из которых занимает более 200 тактов, если необходимые данные не находятся в кеше.

Использование текстурной памяти является хорошим способом решения этой проблемы, но это не обязательно оптимальный способ.

По крайней мере, попробуйте по 4 пикселя за раз (по горизонтали) на поток.Глобальная память может быть доступна по 128 байтов за раз (если у вас есть перекос, пытающийся получить доступ к любому байту с интервалом в 128 байтов, вы также можете использовать всю строку кэша практически без дополнительных затрат).Поскольку деформация состоит из 32 потоков, работа каждого потока на 4 пикселях должна быть эффективной.

Кроме того, вы хотите, чтобы вертикально смежные пиксели обрабатывались одним и тем же мультипроцессором.Причина в том, что смежные строки используют одни и те же входные данные.Если у вас есть пиксель (x = 0, y = 0), обработанный одним MP, а пиксель (x = 0, y = 1) обрабатывается другим MP, оба MP должны выполнить три глобальных запроса памяти каждый.Если они оба обрабатываются одним и тем же MP, и результаты должным образом кэшируются (неявно или явно), вам нужно всего четыре.Это можно сделать, если каждый поток будет работать с несколькими вертикальными пикселями, или с помощью blockDim.y> 1.

В более общем случае, вы, вероятно, захотите, чтобы каждый 32-ниточный деформирующий модуль загружал столько памяти, сколько выиметь доступ к MP (16-48 кбайт или, по крайней мере, блок 128x128), а затем обрабатывать все пиксели в этом окне.

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

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

Нет веских причин использовать float в качестве типа буфера.Вы не только в четыре раза увеличиваете пропускную способность памяти, но и код становится ненадежным и подверженным ошибкам.(Например, вы уверены, что if(neighbors == 3) работает правильно, поскольку вы сравниваете число с плавающей точкой и целое число?) Используйте unsigned char.А еще лучше, используйте uint8_t и typedef, чтобы обозначать беззнаковый символ, если он не определен.

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

3 голосов
/ 06 февраля 2016

TL; DR: см .: http://golly.sourceforge.net

Проблема в том, что большинство реализаций CUDA следуют безмозглой идее ручного подсчета соседей. Это настолько медленно, что любая интеллектуальная последовательная реализация ЦП превзойдет его.

Единственный разумный способ выполнить вычисления в GoL - использовать справочную таблицу.
Самые быстрые в настоящее время реализации на ЦП используют поиск квадратного блока 4x4 = 16 бит, чтобы увидеть, как будут помещаться будущие ячейки 2x2.

в этой настройке ячейки расположены так: 

 01234567
0xxxxxxxx //byte0
1xxxxxxxx //byte1 
2  etc
3
4
5
6
7

Некоторое смещение бит используется для того, чтобы блок 4x4 вписался в слово, и это слово ищется с помощью таблицы поиска. В таблицах поиска также содержатся слова, поэтому в таблице поиска можно сохранить 4 разные версии результата, поэтому вы можете минимизировать количество сдвигов, необходимых для ввода и / или вывода.

Кроме того, разные поколения находятся в шахматном порядке, так что вам нужно смотреть только на 4 соседних плиты, а не на 9. Вот так: 

AAAAAAAA 
AAAAAAAA   BBBBBBBB
AAAAAAAA   BBBBBBBB
AAAAAAAA   BBBBBBBB
AAAAAAAA   BBBBBBBB
AAAAAAAA   BBBBBBBB
AAAAAAAA   BBBBBBBB
AAAAAAAA   BBBBBBBB
           BBBBBBBB
//odd generations (A) are 1 pixel above and to the right of B,
//even generations (B) are 1 pixels below and to the left of A.

Это само по себе приводит к ускорению в 1000 раз + по сравнению с реализациями глупого подсчета.

Затем происходит оптимизация не вычисляемых слябов, которые являются статическими или имеют периодичность 2.

А потом есть HashLife , но это совсем другой зверь.
HashLife может генерировать шаблоны жизни за время O (log n) вместо обычного времени O (n). Это позволяет рассчитать генерацию: 6 366 548 773 467 669 985 195 496 000 (6 октиллионов) за считанные секунды.
К сожалению, Hashlife требует рекурсии, и поэтому это сложно для CUDA.

3 голосов
/ 10 января 2011

посмотрите на эту ветку, мы внесли множество улучшений ...

http://forums.nvidia.com/index.php?showtopic=152757&st=60

...