Всем привет!
Недавняя статья на Хабре в очередной раз показала неостывающий интерес к игре «Жизнь» в частности и всевозможным оптимизациям в общем. Статья и комментарии к ней, особенно любопытство к вычислениям на 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
Memory 8 GB 1600 MHz DDR3
Graphics Intel Iris 1536 MB
| 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 |
| 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
Memory 16 GB 2133 MHz LPDDR3
Graphics Intel HD Graphics 630 1536 MB
| 1024 | 2048 | 4096 | 8192 | 16384 | 32768 | |
|---|---|---|---|---|---|---|
| Низкий уровень (min) | 0 | 0 | 0 | 2 | 8 | 38 |
| Высокий уровень (min) | 0 | 0 | 0 | 0 | 3 | 13 |
| 100 шагов | 35 | 67 | 163 | 450 | 1451 | 5886 |
| Время на шаг (avg) | 0 | 1 | 2 | 5 | 15 | 59 |
| 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 реализации нижнего уровня
Код верхнего уровня
maxzhurkin
24 не является степенью двойки, но, тем не менее, делится без остатка на 8. И таких чисел много
P.S. А, только в частности!