ip / ivanpleshkov.dev
EN

02 Блог

Игра «Жизнь» на compute shader

· gpu ·wgsl ·compute-shader ·graphics

Игра «Жизнь» Конвэя — четыре строчки правил на 2D-сетке: живая клетка с 2 или 3 живыми соседями выживает, мёртвая клетка с ровно 3 становится живой, остальные умирают. Параллелизм огромный, логика тривиальная — естественный compute-shader exercise.

Интересно не сами правила. Интересно, что workgroup’ы не синхронизируются между собой внутри одного dispatch, а Life-обновление требует от каждой клетки прочитать 8 соседей — включая клетки, которыми владеют другие workgroup’ы. Если делать в лоб — гонка с самим собой.

Этот пост — про WGSL-реализацию, которая обходит гонку через паттерн «тайл + halo + staging-буфер», работая in-place на одном storage-буфере.

Открыть интерактивное демо →

(compute.toys запускает WGSL прямо в браузере. Кликаешь — шейдер компилируется — симуляция стартует.)

Почему наивный подход не работает

На CPU «Жизнь» — двойной цикл:

для каждой клетки (x, y):
    n = живые_соседи(x, y)  # читаем из input
    output[x, y] = шаг(input[x, y], n)

Прошёл один полный pass, поменял input и output местами, повторяешь. CPU-потоки (при правильной синхронизации) не запишут что-то, пока ты это читаешь. Естественный GPU-перевод — такой же: ping-pong между двумя буферами. Работает, но требует двойного объёма памяти под сетку.

Эта реализация делает по-другому — один storage-буфер, in-place. Внутри одного dispatch одни и те же клетки одновременно записываются (новое значение) и читаются (как чьи-то соседи) в разных workgroup’ах, без возможности глобально синхронизироваться. Это гонка, и GPU имеет полное право выдать мусор.

Решение: разбиваем работу на тайлы 16×16, копируем каждой workgroup её соседей в shared memory, а границы тайла пишем в отдельный staging-буфер, который в этом pass никто не читает.

Тайлы и halo

Каждая workgroup владеет патчем 16×16. Чтобы применить правила к этим 256 клеткам, workgroup нужны 8 соседей каждой клетки — то есть 1-пиксельное кольцо вокруг патча, принадлежащее соседним workgroup’ам.

Поэтому workgroup загружает в shared memory область 18×18: свои 16×16 плюс 1-пиксельный halo от соседей.

const TILE_SIZE = 16;
const TILE_WITH_HALO = TILE_SIZE + 2; // 18
var<workgroup> tile: array<u32, TILE_WITH_HALO * TILE_WITH_HALO>; // 324

Каждый поток грузит свою клетку. Потоки на границе тайла дополнительно грузят соседнюю halo-точку; угловые потоки — диагональную. После workgroupBarrier() все 324 ячейки заполнены, и каждый поток читает 8 соседей из tile[], не трогая глобальный буфер.

fn apply_game_of_life_rules(local_position: vec3i) -> u32 {
    let cell_value = tile[tile_local_index(local_position)];

    var alive_neighbors = 0u;
    for (var dy = -1; dy <= 1; dy++) {
        for (var dx = -1; dx <= 1; dx++) {
            if (dx == 0 && dy == 0) { continue; }
            let neighbor = vec3i(local_position.x + dx, local_position.y + dy, 0);
            alive_neighbors += tile[tile_local_index(neighbor)];
        }
    }

    if (cell_value > 0) {
        return select(0u, 1u, alive_neighbors == 2 || alive_neighbors == 3);
    } else {
        return select(0u, 1u, alive_neighbors == 3);
    }
}

Тут пока без сюрпризов. Самое интересное начинается, когда мы записываем результат обратно.

Гонка и staging-буфер

Тайл A считает новые значения для своих 16×16. То же самое делает тайл B справа от него. Halo тайла B на самой левой колонке — это самая правая колонка тайла A, его граница. Если A в этом dispatch запишет новое значение границы в storage-буфер, а B в этом же dispatch прочитает из storage-буфера для своего halo, порядок workgroup’ов не определён. B может увидеть новое значение A или старое — случайным образом.

Поэтому: не пишем границы тайла в storage-буфер в том же dispatch, где они читаются. Границы идут в отдельный staging-буфер.

#storage buffer array<u32>          // основная сетка
#storage tile_edges array<u32>      // staging для границ

Внутренние клетки — 14×14, которые не лежат на границе тайла — никогда не попадают в halo соседей, и их можно писать напрямую:

let is_border = write_tile_edge(global_position, local_position, screen_size, new_value);
if (!is_border) {
    // Внутренняя клетка — ничей halo нас не трогает, пишем прямо в основной буфер.
    buffer[global_index(global_position, screen_size)] = new_value;
}

Второй dispatch копирует staged-значения в основной буфер:

@compute @workgroup_size(16, 16)
fn updateBorders(...) {
    let staged_value = read_tile_edge(global_position, local_position, screen_size);
    if (staged_value < 100) {  // sentinel 1000 = «не граница»
        buffer[global_index(global_position, screen_size)] = staged_value;
    }
}

Между двумя dispatch’ами WebGPU-runtime ставит memory barrier. После второго dispatch storage-буфер полностью обновлён.

Layout tile_edges — tile_id * 64 + edge * 16 + offset, по 64 ячейки на тайл (4 ребра × 16 пикселей). Угловые пиксели лежат сразу на двух рёбрах, но write-функция выбирает первое подходящее ребро в цепочке if/else, а read-функция использует тот же порядок — углы оказываются в одной согласованной ячейке, не в двух.

Пайплайн

Четыре dispatch’а на кадр:

initState      (один раз, frame 0)  — инициализация сетки из текстуры
updateState                          — load tile + halo, применение правил,
                                       внутренности в буфер, границы в staging
updateBorders                        — копирование staging-значений в буфер
main_image                           — рендер буфера на экран

#dispatch_once (директива compute.toys) запускает initState только на первом кадре. Остальные три идут каждый кадр. Чтобы симулировать быстрее одного шага в кадр, диспатчим updateState и updateBorders несколько раз за кадр.

Где встречается этот паттерн

Архитектурный приём — разбить работу на тайлы, загрузить тайл + halo в shared memory, делать stencil из shared, аккуратно обработать запись на границах — повторяется в любом GPU-stencil’е:

  • Фильтры изображений: gaussian blur, edge detection, морфология.
  • Свёрточные слои в фреймворках инференса.
  • Fluid simulation (любой explicit-Euler шаг PDE).
  • Векторные индексы: GPU HNSW build в Qdrant использует тот же танец «load tile + stencil», только по нерегулярному графу вместо регулярной сетки.

Конкретно трюк с одним буфером + staging для границ — это специфика. Не каждому stencil’у он нужен; ping-pong c двумя буферами проще, но стоит вдвое больше памяти. Single-buffer + staging — выбор, когда память важнее количества dispatch’ей.

Исходник

Полный WGSL на compute.toys → — ~280 строк, включая хелперы, комментарии и инициализацию сетки из текстуры. Редактор compute.toys позволяет форкнуть и поковырять.