// Бросается в глаза 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));
};