Compute-шейдери: масовопаралельна симуляція

1. Чому compute-шейдери, а не fragment-шейдери

До WebGPU браузерний GPGPU означав зловживання fragment-шейдером: закодувати стан симуляції як пікселі у текстурі з плаваючою крапкою, відрендерити повноекранний квад, а вихідну текстуру декодувати як наступний стан (техніка, що використовується у transform feedback / render-to-texture GPGPU). Це працює, але примушує кожне обчислення проходити через растеризатор, витрачає пропускну здатність на непотрібну логіку семплінгу текстур і обмежує вас 2D-сітками текселів.

Compute-шейдер — це універсальне ядро, яке читає й пише сирі storage-буфери — плоскі масиви структурованих даних, точнісінько як SSBO у Vulkan/OpenGL або kernel у CUDA. Ви диспетчеризуєте 1D-, 2D- чи 3D-сітку потоків напряму, без жодного етапу растеризації між ними. Кожен потік отримує свій global_invocation_id, використовує його для індексації в буфер, обчислює щось і записує результат назад. Це набагато природніше лягає на N-body фізику, солвери тканини, симуляцію рідини, сортування та інференс нейромереж, ніж колишній трюк з текстурою-квадом.

2. Запит GPU-пристрою

Усе в WebGPU починається з асинхронного рукостискання: спочатку запитуємо adapter (фізичний або віртуальний GPU), а потім device (логічне з'єднання, через яке ви фактично надсилаєте команди):

async function initGPU() {
  if (!('gpu' in navigator)) {
    throw new Error('WebGPU не підтримується в цьому браузері');
  }

  const adapter = await navigator.gpu.requestAdapter({
    powerPreference: 'high-performance',
  });
  if (!adapter) throw new Error('Не знайдено придатний GPU-адаптер');

  const device = await adapter.requestDevice({
    requiredLimits: {
      maxStorageBufferBindingSize: 512 * 1024 * 1024, // 512 МБ
    },
  });

  device.lost.then(info => {
    console.error(`GPU-пристрій втрачено: ${info.message}`);
  });

  return device;
}

На відміну від синхронного getContext('webgl2') у WebGL, тут це промісо-орієнтоване узгодження — браузеру може знадобитися підняти окремий GPU-процес і перевірити ліміти, перш ніж віддати вам пристрій. Запитуйте device лише один раз і повторно використовуйте його протягом усього життя застосунку.

3. Storage-буфери та bind group

Storage-буфер — це сирий блок GPU-пам'яті, який ваш шейдер може читати і, що важливо, записувати — на відміну від uniform-буфера, доступного лише для читання і значно меншого за розміром (зазвичай 64 КБ проти сотень МБ у storage). Буфери створюються з явними прапорцями використання:

const PARTICLE_COUNT = 500_000;
const FLOATS_PER_PARTICLE = 8; // pos.xyz, vel.xyz, life, pad

const particleBuffer = device.createBuffer({
  size: PARTICLE_COUNT * FLOATS_PER_PARTICLE * 4, // 4 байти на f32
  usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
  mappedAtCreation: true, // записуємо початкові дані до першого використання
});
new Float32Array(particleBuffer.getMappedRange()).set(initialData);
particleBuffer.unmap();

// bind group layout описує, до чого має доступ стадія шейдера
const bindGroupLayout = device.createBindGroupLayout({
  entries: [{
    binding: 0,
    visibility: GPUShaderStage.COMPUTE,
    buffer: { type: 'storage' },
  }],
});

// bind group — це конкретний ресурс, прив'язаний до цього layout
const bindGroup = device.createBindGroup({
  layout: bindGroupLayout,
  entries: [{ binding: 0, resource: { buffer: particleBuffer } }],
});

4. WGSL: мова шейдингу

WGSL (WebGPU Shading Language) статично типізована, має явні режими доступу read/read_write для storage-прив'язок і, на відміну від GLSL, вимагає явно оголошувати макет структури, що точно відповідає макету пам'яті вашого буфера:

struct Particle {
  pos  : vec3<f32>,
  vel  : vec3<f32>,
  life : f32,
  _pad : f32,
}

@group(0) @binding(0) var<storage, read_write> particles : array<Particle>;

struct Uniforms {
  dt         : f32,
  time       : f32,
  noiseScale : f32,
}
@group(0) @binding(1) var<uniform> u : Uniforms;

@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid : vec3<u32>) {
  let i = gid.x;
  if (i >= arrayLength(&particles)) { return; } // захист від зайвих потоків

  var p = particles[i];
  let curl = curlNoise(p.pos * u.noiseScale + u.time * 0.1);
  p.vel += curl * u.dt;
  p.pos += p.vel * u.dt;
  p.life -= u.dt;

  if (p.life <= 0.0) {
    p.pos  = vec3<f32>(0.0);
    p.vel  = vec3<f32>(0.0);
    p.life = 5.0;
  }
  particles[i] = p;
}

5. Workgroup і виміри dispatch

Compute-ядро викликається як сітка workgroup, а кожен workgroup — сам є сіткою invocation (потоків), розмір якої ви оголошуєте у шейдері через @workgroup_size(x, y, z). Загальна кількість викликів дорівнює workgroup_size × dispatch_count:

const WORKGROUP_SIZE = 64;
const dispatchX = Math.ceil(PARTICLE_COUNT / WORKGROUP_SIZE);

const pass = encoder.beginComputePass();
pass.setPipeline(pipeline);
pass.setBindGroup(0, bindGroup);
pass.dispatchWorkgroups(dispatchX); // 1D-dispatch
pass.end();

Вибір workgroup_size має значення: 64 — безпечне, портативне значення за замовчуванням, яке ділиться без остачі на 32-широкі warp'и (NVIDIA) і 64-широкі wavefront'и (AMD), що фактично виконуються в лок-степі на реальному залізі. Розміри менші за 32 марнують доріжки виконання; розміри значно більші за 256 можуть шкодити occupancy, вичерпуючи бюджет спільної пам'яті та регістрів на workgroup. Тестуйте 64, 128 і 256 для вашого конкретного ядра — оптимум залежить від навантаження й GPU.

6. Побудова compute-конвеєра

GPUComputePipeline прив'язує точку входу шейдерного модуля до фіксованого макету bind group. Після побудови диспетчеризація щокадру дешева — дорога компіляція шейдера відбувається лише раз, при створенні конвеєра:

const shaderModule = device.createShaderModule({ code: wgslSource });

const pipeline = device.createComputePipeline({
  layout: device.createPipelineLayout({
    bindGroupLayouts: [bindGroupLayout],
  }),
  compute: { module: shaderModule, entryPoint: 'main' },
});

function frame() {
  const encoder = device.createCommandEncoder();
  const pass = encoder.beginComputePass();
  pass.setPipeline(pipeline);
  pass.setBindGroup(0, bindGroup);
  pass.dispatchWorkgroups(dispatchX);
  pass.end();
  device.queue.submit([encoder.finish()]);
  requestAnimationFrame(frame);
}

Зверніть увагу — у циклі кадру немає await: queue.submit() — «запустив і забув»: GPU виконує команди асинхронно, і ви ніколи не блокуєте головний потік очікуванням завершення compute-обчислень, на відміну від наївного циклу симуляції на CPU.

7. Кейс: поле течії з 500К частинок

Об'єднуючи все вищесказане: поле течії на основі curl-шуму, що переносить пів мільйона частинок, рендериться напряму з того ж storage-буфера через render-конвеєр WebGPU (зчитування на CPU не потрібне — vertex-шейдер читає particles[vertexIndex] напряму):

async function setupFlowField(device) {
  const N = 500_000;
  const initial = new Float32Array(N * 8);
  for (let i = 0; i < N; i++) {
    const o = i * 8;
    initial[o + 0] = (Math.random() - 0.5) * 10; // pos.x
    initial[o + 1] = (Math.random() - 0.5) * 10; // pos.y
    initial[o + 2] = (Math.random() - 0.5) * 10; // pos.z
    initial[o + 6] = Math.random() * 5; // life
  }

  const buffer = device.createBuffer({
    size: initial.byteLength,
    usage: GPUBufferUsage.STORAGE | GPUBufferUsage.VERTEX,
    mappedAtCreation: true,
  });
  new Float32Array(buffer.getMappedRange()).set(initial);
  buffer.unmap();
  return buffer;
}

На дискретному GPU середнього класу це ядро оновлює й рендерить 500 000 частинок значно швидше ніж за 2 мс на кадр — порівняно з десятками мілісекунд для еквівалентного циклу на CPU у JavaScript, і масштабується майже лінійно з кількістю ядер GPU, а не обмежується продуктивністю однопотокового JS. Та сама архітектура напряму узагальнюється на флокінг boids, сусідські суми SPH-симуляції рідини та розв'язання обмежень тканини — усе, що виражається як «для кожного елемента прочитати стан, записати новий стан».

8. Бар'єри та гонки read-after-write

Потоки всередині workgroup можуть ділитися швидкою пам'яттю на чипі — workgroup memory (var<workgroup>), але читання й запис у неї потрібно явно синхронізувати через workgroupBarrier() — без цього один потік може прочитати застарілі дані, які інший потік ще не встиг записати, оскільки виклики всередині workgroup не виконуються покроково синхронно на всіх виробниках заліза:

var<workgroup> tile : array<f32, 64>;

@compute @workgroup_size(64)
fn main(
  @builtin(local_invocation_id) lid : vec3<u32>,
  @builtin(global_invocation_id) gid : vec3<u32>
) {
  tile[lid.x] = particles[gid.x].pos.x;
  workgroupBarrier(); // чекаємо: кожен потік має завершити запис вище
                       // перш ніж будь-хто читатиме tile[]

  var sum = 0.0;
  for (var j = 0u; j < 64u; j += 1u) {
    sum += tile[j]; // безпечно: бар'єр гарантує, що всі 64 записи вже виконані
  }
}

Між різними dispatch-викликами (наприклад, двома послідовними викликами dispatchWorkgroups() в одному проході, чи окремими compute-проходами) WebGPU автоматично гарантує порядок — пізніший dispatch, що читає storage-буфер, завжди побачить записи попереднього dispatch, який його записував. Вам ніколи не потрібні ручні бар'єри буферів між проходами, як у сирому Vulkan, — рівень валідації браузера сам вставляє необхідну синхронізацію.

9. Продуктивність: occupancy та memory coalescing

На практиці пропускну здатність compute-шейдера визначають три речі:

• Occupancy: скільки workgroup можуть виконуватись паралельно на одному обчислювальному блоці, обмежене використанням спільної пам'яті та регістрів на потік. Тримайте масиви var<workgroup> невеликими й уникайте надмірної кількості локальних змінних у гарячих циклах.
• Memory coalescing: сусідні потоки (gid.x, gid.x+1, ...) мають читати/писати сусідні адреси пам'яті. Макет struct-of-arrays (окремі буфери positions[], velocities[]) часто коалесциться краще за array-of-structs для широких ядер з уніфікованим доступом, хоча AoS зазвичай простіше в розумінні й цілком підходить, коли поля кожного потоку читаються разом.
• Розбіжне галуження: if, що йде різними шляхами для різних потоків одного workgroup, змушує залізо виконувати обидва шляхи послідовно для всієї групи. Тримайте гілки уніфікованими там, де можливо, або перебудовуйте з select() для дешевого вибору між двома варіантами.

10. Визначення підтримки і резервний WebGL

Приблизно п'ята частина браузерного трафіку станом на середину 2026 року досі не підтримує WebGPU (старіші версії Safari/iOS, деякі корпоративні збірки Chrome з вимкненою функцією за політикою). Продакшн-симуляція має визначати можливості й плавно деградувати до техніки WebGL2 transform feedback:

async function createSimBackend(canvas) {
  if ('gpu' in navigator) {
    const adapter = await navigator.gpu.requestAdapter();
    if (adapter) {
      return new WebGPUBackend(await adapter.requestDevice());
    }
  }
  console.warn('WebGPU недоступний — переходимо на WebGL2 transform feedback');
  return new WebGL2Backend(canvas.getContext('webgl2'));
}

Структуруйте ядро симуляції за невеликим інтерфейсом — step(dt), getPositionBuffer(), dispose() — щоб рівні рендерингу й UI ніколи не знали, який бекенд активний. Це той самий патерн, що використовується в туторіалі про WebGL compute-частинки на цьому сайті, і він дозволяє випустити одну симуляцію, яка працює всюди, отримуючи при цьому весь виграш у продуктивності від WebGPU там, де він доступний.