Conway's Game of Life is embarrassingly parallel, so one CUDA thread per cell gives a 70x speedup over Python

Video. CUDA Game of Life | Hardcore CUDA Hackathon Demos at AGI House · 06:59 · YouTube

I walk through porting Conway's Game of Life from Python to C to CUDA over a day at the AGI House Hardcore CUDA Hackathon, and compare the four implementations across 100,000 generation steps. The code is on GitHub.

The idea

Game of Life is a zero-player cellular automaton on a 2D grid. Each cell's next state depends only on its current state and its eight immediate neighbours, so every cell in a generation can be evaluated independently of every other cell. That makes the state transition embarrassingly parallel: instead of a CPU sweeping the grid one cell at a time, a GPU can assign one thread per cell and compute the entire next generation in a single kernel launch.

The project also looks at what it costs to watch the simulation. Streaming every intermediate frame back to host memory with cudaMemcpyAsync lets you render the grid in the terminal in real time, but the transfer dominates the runtime once the kernel itself is fast.

The mechanism

For a grid $S \in \{0, 1\}^{N \times M}$ , each cell $(i, j)$ sums its Moore neighbourhood:

n(i, j) = \sum_{\substack{dx, dy \in \{-1, 0, 1\} \\ (dx, dy) \neq (0, 0)}} S_{i + dx,\, j + dy}

Cells outside the grid count as dead (no wraparound). The update follows the standard rules:

S^{(t+1)}_{i,j} = \begin{cases} 1 & \text{if } S^{(t)}_{i,j} = 1 \text{ and } n(i, j) \in \{2, 3\} \\ 1 & \text{if } S^{(t)}_{i,j} = 0 \text{ and } n(i, j) = 3 \\ 0 & \text{otherwise} \end{cases}

The third case covers underpopulation ( $n < 2$ ) and overpopulation ( $n > 3$ ).

CUDA layout

One block, one thread per cell. The 32×32 grid maps onto a single thread block of 32×32 = 1,024 threads, launched as kernel<<<1, dim3(32, 32)>>>. Each thread reads its cell coordinates directly from threadIdx.x and threadIdx.y.
Neighbour count. Each thread calls a __device__ function that loops over the 3×3 window with a bounds check, then subtracts the centre cell.
Double buffering. The thread writes its result into a separate next_state buffer so that no thread reads a cell that has already been updated this generation.
In-kernel swap. After a __syncthreads() barrier, each thread copies next_state back into state. Because the whole grid lives in one block, the block-wide barrier is sufficient to guarantee every thread has finished reading before anyone writes.

__global__ void fast_update_next_state(int a[][GRID_SIZE], int next_state[][GRID_SIZE])
{
    int x = threadIdx.x;
    int y = threadIdx.y;
    int n = count_neighbors(a, x, y);

    if (n < 2 || n > 3)      next_state[x][y] = 0;
    else if (n == 3)         next_state[x][y] = 1;
    else                     next_state[x][y] = a[x][y];

    __syncthreads();
    a[x][y] = next_state[x][y];
}

The host loop is then just 100,000 kernel launches, optionally followed by a cudaMemcpyAsync of the grid back to the host for rendering.

Performance

100,000 steps on a 32×32 grid, from the repository README:

Implementation	Seconds / 100,000 steps	Speedup vs Python
Python	~62.17	1.0x
C (CPU)	2.58	24x
CUDA, copying every frame to host	1.78	35x
CUDA, end state only	0.88	71x

The C port alone gives most of the win over Python. CUDA roughly triples that again, but only if you stop asking for every frame back.

Worth knowing

The host copy is the bottleneck, not the kernel. Copying the 4 KB grid back to the host after every step doubles the runtime (1.78 s vs 0.88 s). Even though the copy is asynchronous, 100,000 small transfers over PCIe and the associated host synchronisation cost more than the compute. If you only need the final state, or can render every $k$ -th frame, keep the state on the device.
The one-block design does not scale. __syncthreads() only synchronises threads within a block, and a block is capped at 1,024 threads. Anything larger than 32×32 needs multiple blocks, and there is no cheap grid-wide barrier inside a kernel. The standard approach is to tile the grid across blocks with a one-cell halo of overlapping boundary cells, and to swap buffers between kernel launches (the launch boundary acts as the global barrier) rather than inside the kernel.
Shared memory buys nothing at this size. A shared-memory tiling version was prototyped, but a 32×32 int grid is 4 KB and stays resident in cache, so the explicit staging step is pure overhead. Shared memory starts to matter when many blocks each re-read overlapping halos from global memory.
Dense sweeps waste work on dead space. Every step evaluates all $N \times M$ cells even when a handful are alive. On large, sparse universes a sparse representation of live cells, or a memoised algorithm like Hashlife, wins by orders of magnitude regardless of hardware.

Use it when / don't use it when

Use it when

Simulating large, dense cellular automata (2D or 3D) where most of the grid is active and you need millions of cell updates per generation.
Generating synthetic grid-transition datasets at scale, for example as training data for reasoning or reinforcement learning models on ARC-style tasks.
The automaton is embedded in a GPU pipeline already (rendering, physics, simulation) and the state never needs to leave the device.

Don't use it when

The grid is tiny. For a 32×32 grid the C version is already fast, and kernel launch overhead is a large fraction of each step.
The universe is very sparse. If live cells are a small fraction of the grid, Hashlife or a sparse set-based CPU implementation will outperform a dense GPU sweep.
You need every intermediate frame on the host. The device-to-host transfer will dominate, and the GPU advantage largely disappears.