Game of Life с битовой магией, многопоточностью и на GPU

Game of Life с битовой магией, многопоточностью и на GPU

Публикации

Истории

Работа

Ближайшие события

Всем привет!

О себе

Масштабы

Архитектура

Низкий уровень

Высокий уровень

Бенчмарки

Итог

Всем привет!

О себе

Масштабы

Архитектура

Низкий уровень

Высокий уровень

Бенчмарки

Итог

Moximko 3 июл 2020 в 19:44

8 мин

14K

C++*Алгоритмы*ВидеокартыВысокая производительность*

Технотекст 2020

Недавняя статья на Хабре в очередной раз показала неостывающий интерес к игре «Жизнь» в частности и всевозможным оптимизациям в общем. Статья и комментарии к ней, особенно любопытство к вычислениям на GPU, вдохновили меня на то, чтобы поделиться своими изысканиями на данном поприще и, забегая вперёд, скажу, что повествование пойдёт о расчётах на GPU, битовой магии, многопоточности и огромных полях для игры «Жизнь», порядка миллиарда клеток.

Пару слов о себе и проекте. Вот уже несколько лет моим хобби и pet-проектом является написание игры «Жизнь», в ходе которого я разбираюсь в интересующих меня темах. Так, сперва я сделал её с помощью библиотеки glut и прикрутил многопользовательский режим, будучи вдохновлённым архитектурой peer-to-peer. А около двух лет назад решил начать всё с нуля, используя Qt и вычисления на GPU.

Идея нового проекта заключалась в том, чтобы сделать кроссплатформенное приложение, в первую очередь для мобильных устройств, в котором пользователи смогут окунуться в игру «Жизнь», познакомиться с разнообразием всевозможных паттернов, наблюдать за причудливыми формами и комбинациями этой игры, выбрав скорость от 1 до 100 мс на шаг и задав размер игрового поля вплоть до 32'768 x 32'768, то есть 1'073'741'824 клетки.

На всякий случай напомню об игре «Жизнь». Игра происходит на плоскости, поделённой на множество клеток. Клетки квадратные, соответственно, для каждой клетки есть 8 соседей. Каждая клетка может быть либо живой, либо мёртвой, то есть пустой. В рамках игры пользователь заполняет клетки жизнью, выставляя на поле заинтересовавшие его комбинации клеток — паттерны.

Далее процесс шаг за шагом развивается по заданным правилам:

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

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

В моей реализации поле зациклено, то есть после 32'767-й клетки вновь следует 0-я. Таким образом, перемещающиеся паттерны способны обогнуть поле и оказаться в точке, где они были размещены изначально. Забавный эффект случается с паттерном «планёрное ружьё», когда выпущенные им планёры в конечном итоге врезаются в само ружьё и разрушают его, этакое самоубийство в мире клеточных автоматов:

Что касается реализации, то в первую очередь хочется отметить, что каждая клетка в игровом поле представляет собой всего лишь 1 бит в памяти. Так, при выборе размера игрового поля в памяти выделяется буфер размером n^2 / 8 байт. Это улучшает локальность данных и снижает объём потребляемой памяти, а главное позволяет применить ещё более хитроумные оптимизации, о которых поговорим ниже. Игровое поле всегда квадратное и его сторона всегда степени 2, в частности затем, чтобы без остатка осуществлять деление на 8.

Архитектурно выделяются два слоя, ответственные за вычисления:

низкий уровень — платформозависимый; на текущий момент существует реализация на Metal API — графическое API от Apple позволяет задействовать GPU на устройствах от Apple, а также fallback-реализация на CPU. Одно время существовала версия на OpenCL. Планируется реализация на Vulkan API для запуска на Android
высокий уровень — кроссплатформенный; по сути класс-шаблонный метод, делегирующий реализацию нужных методов на низкий уровень

Задача низкого уровня заключается непосредственно в расчёте состояния игрового поля и устроена следующим образом. В памяти хранятся два буфера n^2 / 8 байт. Первый буфер служит как input — текущее состояние игрового поля. Второй буфер — как output, в который записывается новое состояние игрового поля в процессе расчётов. По завершении вычислений достаточно сделать swap буферов и следующий шаг игры готов. Такая архитектура позволяет распараллелить расчёты в силу константности input. Дело в том, что каждый поток может независимо обрабатывать клетки игрового поля.

За счёт многопоточности GPU-подход добивается максимальной эффективности. На CPU также происходит распараллеливание, но, соответственно, с гораздо меньшим числом потоков и меньшей эффективностью, что отражено в секции бенчмарков ниже. Сам алгоритм идентичен для всех реализаций и заключается в следующем.

Назовём атомарной операцией расчёт одного байта буфера, хранящего игровое поле, то есть расчёт восьми клеток. Этот расчёт совершается одним потоком и требует информации о соседних восьми байтах. Попробуем схематично изобразить один байт следующим образом, где точка — это бит:

[........]

А участок игрового поля с интересующим нас байтом и его восемью соседними байтами изобразим так:

[........][........][........]
[........][********][........]
[........][........][........]

Исходя из рисунка, можно догадаться, что все соседние клетки можно уложить в один uint64_t (назовём его neighbours), а ещё в одном uint8_t (назовём его self), будет храниться информация о соседях в самом обрабатываемом байте.

Рассмотрим на примере расчёт 0-го бита целевого байта. Звёздочкой отметим интересующий бит, а нулём соседние для него биты:

[.......0][00......][........]
[.......0][*0......][........]
[.......0][00......][........]

Пример посложнее. Здесь звёздочкой отмечены 0-й, 3-й и 7-й биты целевого байта и соответствующими числами соседние биты:

[.......0][00333.77][7.......]
[.......0][*03*3.7*][7.......]
[.......0][00333.77][7.......]

Думаю, кто-то из читателей уже догадался, в чём заключается магия. Имея указанную информацию, остаётся лишь составить битовые маски для каждого бита целевого байта и применить их к neighbours и self соответственно. Результатом станут 2 значения, сумма единичных бит которых будет характеризовать число соседей, что можно интерпретировать как правила игры «Жизнь»: 2 или 3 бита для поддержания жизни в живой клетке и 3 бита для зарождения новой жизни в пустой клетке. В противном случае клетка остаётся/становится пустой.

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

После заполнения всего выходного буфера игровое поле считается рассчитанным.

Код Metal-шейдера по обработке 1 байта

// Бросается в глаза C++-подобный синтаксис
#include <metal_stdlib>
#include <metal_integer>

using namespace metal;

// Вспомогательные функции для вычисления позиции клетки
ushort2 pos(uint id) 
{ 
  return ushort2(id % WIDTH, id / HEIGHT); 
}
uint idx(ushort2 pos) 
{ 
  return pos.x + pos.y * HEIGHT; 
}
ushort2 loopPos(short x, short y) 
{ 
  return ushort2((x + WIDTH) % WIDTH, (y + HEIGHT) % HEIGHT); 
}

// Битовые маски для вычисления соседей интересующего бита
template<uint Bit> 
struct Mask 
{ 
  constant constexpr static uint c_n_e_s_w = 0x70007 << (Bit - 1); 
  constant constexpr static uint c_nw_ne_se_sw = 0x0; 
  constant constexpr static uint c_self = 0x5 << (Bit - 1); 
};
template<> 
struct Mask<0> 
{ 
  constant constexpr static uint c_n_e_s_w = 0x80030003; 
  constant constexpr static uint c_nw_ne_se_sw = 0x80000080; 
  constant constexpr static uint c_self = 0x2; 
};
template<> 
struct Mask<7> 
{ 
  constant constexpr static uint c_n_e_s_w = 0xC001C0; 
  constant constexpr static uint c_nw_ne_se_sw = 0x10100; 
  constant constexpr static uint c_self = 0x40; 
};

// Для указанного бита функция вычисляет состояние клетки в зависимости от её соседей, применяя соответствующие биту маски
template<uint Bit>
uint isAlive(uint self, uint n_e_s_w, uint nw_ne_se_sw) 
{ 
  /*
  [.......0][00333.77][7.......]
  [.......0][*03*3.7*][7.......]
  [.......0][00333.77][7.......]
  */
  // До определённой версии в Metal не было 64-битного целого, поэтому составляются две маски
  uint neighbours = popcount(Mask<Bit>::c_n_e_s_w & n_e_s_w) 
    + popcount(Mask<Bit>::c_nw_ne_se_sw & nw_ne_se_sw) 
    + popcount(Mask<Bit>::c_self & self); 

  return static_cast<uint>((self >> Bit & 1) == 0 
    ? neighbours == 3 
    : neighbours == 2 || neighbours == 3) << Bit;
}

// Язык Metal даже умеет в шаблонную магию
template<uint Bit>
uint calculateLife(uint self, uint n_e_s_w, uint nw_ne_se_sw) 
{ 
  return isAlive<Bit>(self, n_e_s_w, nw_ne_se_sw) 
    | calculateLife<Bit - 1>(self, n_e_s_w, nw_ne_se_sw); 
}
template<>
uint calculateLife<0>(uint self, uint n_e_s_w, uint nw_ne_se_sw)
{
  return isAlive<0>(self, n_e_s_w, nw_ne_se_sw); 
}

// Главная функция compute-шейдера. На вход подаются два буфера, о которых речь шла выше - константный input и output, а также id - координата целевого байта
kernel void lifeStep(constant uchar* input [[buffer(0)]], 
                            device uchar* output [[buffer(1)]], 
                            uint id [[thread_position_in_grid]]) 
{ 
  ushort2 gid = pos(id * 8); 
  // Вычисляем соседние байты
  uint nw = idx(loopPos(gid.x - 8, gid.y + 1)); 
  uint n  = idx(loopPos(gid.x,     gid.y + 1)); 
  uint ne = idx(loopPos(gid.x + 8, gid.y + 1)); 
  uint e  = idx(loopPos(gid.x + 8, gid.y    )); 
  uint se = idx(loopPos(gid.x + 8, gid.y - 1)); 
  uint s  = idx(loopPos(gid.x    , gid.y - 1)); 
  uint sw = idx(loopPos(gid.x - 8, gid.y - 1)); 
  uint w  = idx(loopPos(gid.x - 8, gid.y    )); 

  // Вычисляем байт с целевым битом
  uint self = static_cast<uint>(input[id]); 

  // Подготавливаем битовые маски с соседями
  // north_east_south_west
  uint n_e_s_w = static_cast<uint>(input[n >> 3]) << 0 * 8 
    | static_cast<uint>(input[e >> 3]) << 1 * 8 
    | static_cast<uint>(input[s >> 3]) << 2 * 8 
    | static_cast<uint>(input[w >> 3]) << 3 * 8; 

  // north-west_north-east_south-east_south-west
  uint nw_ne_se_sw = static_cast<uint>(input[nw >> 3]) << 0 * 8 
    | static_cast<uint>(input[ne >> 3]) << 1 * 8 
    | static_cast<uint>(input[se >> 3]) << 2 * 8 
    | static_cast<uint>(input[sw >> 3]) << 3 * 8; 
  
  // В этой строчке рассчитываются все 8 клеток обрабатываемого байта
  output[id] = static_cast<uchar>(calculateLife<7>(self, n_e_s_w, nw_ne_se_sw)); 
};

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

Отрисовка клеток выполняется средствами Qt, а именно посредством заполнения пикселей в QImage. Интерфейс выполнен в QML. Пиксели заполняются лишь для небольшой области видимого игроку игрового поля. Таким образом удаётся избежать лишних затрат ресурсов на отрисовку.

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

Результаты замеров скорости работы игры для разных машин в зависимости от размеров игрового поля приведены в таблицах. Полученные значения в таблицах указаны в миллисекундах. При проведении тестов игровое поле заполнялось жизнью на 50% с помощью генератора случайных чисел.

В таблицах есть замеры скорости работы нижнего и верхнего уровня, причём берутся минимальные значения. А для более общей картины приведены замеры скорости вычисления 100 шагов игры и среднего времени выполнения шага. Последнее более показательно, так как включает в себя время на отрисовку и прочие сопутствующие операции.

MacBook Pro 2014

Processor 2,6 GHz Dual-Core Intel Core i5
Memory 8 GB 1600 MHz DDR3
Graphics Intel Iris 1536 MB

**GPU реализация**
	1024	2048	4096	8192	16384	32768
Низкий уровень (min)	0	0	2	9	43	170
Высокий уровень (min)	0	0	0	1	12	55
100 шагов	293	446	1271	2700	8603	54287
Время на шаг (avg)	3	4	13	27	86	542

**CPU реализация**
	1024	2048	4096	8192	16384	32768
Низкий уровень (min)	3	25	117	552	2077	21388
Высокий уровень (min)	0	0	0	1	4	51
100 шагов	944	3894	15217	65149	231260	-
Время на шаг (avg)	9	39	152	651	2312	-

MacBook Pro 2017

Processor 2.8 GHz Intel Core i7
Memory 16 GB 2133 MHz LPDDR3
Graphics Intel HD Graphics 630 1536 MB

**CPU реализация**
	1024	2048	4096	8192	16384	32768
Низкий уровень (min)	1	7	33	136	699	2475
Высокий уровень (min)	0	0	0	0	3	18
100 шагов	434	1283	4262	18377	79656	264711
Время на шаг (avg)	4	13	43	184	797	2647

Видно, что на стареньком макбуке процессор совсем не справляется с огромным игровым полем и приложение становится неюзабельным. Даже процессор нового макбука едва вытягивает такой объём вычислений. Зато видеокарта значительно превосходит по производительности процессор как на старом, так и на новом железе. С использованием GPU новый макбук предлагает вполне комфортный пользовательский опыт даже на огромных игровых пространствах.

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

В отличие от оригинальной статьи я не ставил целью провести бенчмарк различных оптимизаций, поэтому сравнивать две статьи один к одному нельзя, зато материал отлично раскрывает тему сравнения потенциала GPU и CPU со значительным превосходством первого.

В будущем хочется провести бенчмарки на мобильных устройствах. Проект находится на ранней стадии разработки, а потому Apple-Developer аккаунта для проведения таких тестов сейчас у меня нет.

Спасибо за внимание! Буду рад любым комментариям к статье и к проекту:
Код на GitHub.

Код Metal реализации нижнего уровня

Код CPU реализации нижнего уровня

Код верхнего уровня

Теги:

Хабы:

Низкий уровень (min)

Высокий уровень (min)

Время на шаг (avg)