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

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

__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 из входного массива в двумерный текст inputTex, связанный с массивом CUDA. Выходные данные выводятся из глобальной памяти на хост, а затем обрабатываются.

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

Ответы

Ответ 1

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

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

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

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

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

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

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

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

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

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

Ответ 3

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.

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

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

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