Write a SYCL-accelerated version of the Game of Life cellular automaton. The Game of Life is a zero-player game that simulates the evolution of a grid of cells based on a set of simple rules. Each cell can be either alive (`1`) or dead (`0`), and its state in the next generation is determined by the states of its eight neighbors using the following rules:

1.  Any live cell with fewer than two live neighbors dies (underpopulation).
2.  Any live cell with more than three live neighbors dies (overpopulation).
3.  Dead cells with exactly three live neighbors become alive (reproduction).

In all other cases, the cell's state remains unchanged (i.e., a live cell with two or three live neighbors stays alive, and a dead cell with any number of live neighbors other than three stays dead).

The cell update is performed on the entire grid in each iteration and the number of iterations is specified as a configuration parameter. A copy of the grid is needed when performing the update to avoid overwriting the current state of the cells while calculating the next state (even when using a serial implementation); however, the final result should end in the data structure that was used as the input.

The input grid has a one-cell-wide halo of dead cells around the actual grid. That is, all cells with coordinates `(0,*)`, `(*,0)`, `(W-1,*)`, and `(*,H-1)` are always `0` (and must stay `0`). The updates are performed only on the `(1,1)` to `(W-2,H-2)` range of the grid. Hence, the implementation does not have to check for out-of-bounds accesses when reading the neighbors of a cell.

To save memory (both in space and bandwidth), the grid is already encoded using bit-packing (one cell is represented by one bit). A cell grid of size `W x H` (width times height) is represented by a `Grid<cellpack_t, 2>` class, where `cellpack_t` is a 64-bit unsigned integer type. Blocks of aligned adjacent cells on a row are packed into a single `cellpack_t` value, so the dimensions of the `Grid` class are `W/64 x H` (the width is divided by 64, and the height remains unchanged). Accessing the value of the cell at position `(x, y)` can be done by loading the `cellpack_t` value at `(x/64, y)` and checking the bit at position `x % 64` (the bits are indexed from the least significant bit).

The `getCell()` and `setCell()` functions, which extract and encode the cell value from/to a given `cellpack_t` value, are already provided in `gol.hpp`. More details on how to use them can be found in the serial implementation.

### Technical details

The starter pack with the serial implementation and makefile is available to you on the `hpc.troja.mff.cuni.cz` server in `/home/teaching/NPRG042/labtest`. The performance will be tested on the `volta02` machine, using two V100 GPUs (same as in the SYCL home assignment).

The referential serial implementation can be found in the `serial.hpp` file (the `operator()` of the `SerialImplementation` class). You need to implement the SYCL version in the `sycl.hpp` file (the interface of the `SyclImplementation` class must be preserved), and the `sycl.hpp` file **is the only file you should submit to ReCodEx**.

The testing datasets are available in the `/home/teaching/NPRG042/labtest/data` directory. There are two input files (`stargate.lif` and `guns.lif`) that are used for real tests. The framework generates random initial states in case no input file is provided. Let us note that the grid size should be specified (and it is independent of the input file).

### Evaluation

Your solutions will be tested on 4 configurations/datasets:

-   `stargate.lif` with grid size 1024×1024
-   `guns.lif` with grid size 2048×2048
-   randomly generated grid with size 6400x1920
-   randomly generated grid with size 3392x5760

The smallest (stargate) configuration will be measured for `1000` iterations, while the remaining three configurations use `100` iterations. The binary snapshots of the output grids are provided as `.out` files (e.g., `stargate-1k-1k.out`); you may use them with the `--verify` argument to check the correctness of your implementation.

Furthermore, you may use the associated `test.sh` script, which uses `srun` to compile the code, run the serial and parallel versions, verify correctness, and measure the speedup. The script takes two or three arguments: `./test.sh 1k 1k stargate` invokes the first configuration, `./test.sh 6400 1920` (without the third argument) invokes the third configuration (a randomly generated grid). You may use it for debugging as well (for smaller grid sizes); it will use the serial output as a reference for correctness verification.

Your solution will be evaluated on the four aforementioned configurations, and the final speedup will be calculated as the harmonic mean of the individual speedups. The points will be awarded according to the following table:

-   100× or less: **0 points**
-   100× to 250×: **1 point**
-   250× to 400×: **2 points**
-   400× to 600×: **3 points**
-   600× or more: **4 points**

### Hints

First of all, do not forget that the grid width (size in dimension 0) is already divided by 64.

Please note that both the `Grid` class and SYCL multi-dimensional buffers use row-major order. **However, they differ in the order in which the dimension coordinates are specified.** In the case of our `Grid`, dimension 0 is horizontal (`x`, `width`), and dimension 1 is vertical (`y`, `height`). When creating the corresponding 2D SYCL buffer, one needs to specify the dimensions in the opposite order (`sycl::range<2>{height, width}`). If you use `sycl::id<2>` to index into the buffer, you also need to write the dimensions in reverse order (i.e., `id[1]` is the `x` dimension, and `id[0]` is the `y` dimension).

If you are not comfortable with SYCL multi-dimensional buffers, you can also use 1D buffers and calculate the corresponding 1D index from the 2D coordinates yourselves (and vice versa). The cell `(x, y)` is stored in the cell-pack with index `y * W/64 + x/64` (bit `x % 64`).

Although the serial version iterates over the cell-grid (i.e., without any concern for bit-packing), it might be worth considering iterating over the cell-packs instead of individual cells in the SYCL version. The main reason is that the instructions cannot operate on individual bits, but only on the whole 64-bit values. So, dividing the workload to perform 64 updates (of the whole cell-pack) in a single work-item would not require any additional synchronization.