🖥️ Графіка та рендеринг · Веб-API
📅 Квітень 2026⏱ 15 хв🔴 Просунутий рівень

Обчислювальні шейдери WebGPU: програмування GPU загального призначення в браузері

Обчислювальний конвеєр WebGPU відкриває програмування GPU загального призначення безпосередньо в браузері — без потреби в графічному конвеєрі. Фізика частинок, виведення нейронних мереж, обробка зображень та процедурна генерація, що зайняли б секунди на CPU, можуть виконатися за мілісекунди, коли тисячі ядер GPU працюють паралельно. Ця стаття охоплює повний робочий процес обчислень: шейдери WGSL, топологію робочих груп, буфери зберігання, синхронізацію та практичні шаблони для навантажень масштабу симуляції.

1. WebGPU проти WebGL: чому важливі обчислення

WebGL надає OpenGL ES 2.0/3.0 — суто графічний API, де шейдери можуть лише читати/писати через графічний конвеєр (вершина → растеризація → фрагмент). Немає способу записати назад у довільну пам'ять із шейдера. Завдання обчислювального типу доводилося обходити через рендеринг у текстуру й зчитування кольорів пікселів як виходу.

WebGPU надає семантику Vulkan/Metal/D3D12 з повноцінним обчислювальним конвеєром, окремим від конвеєра рендерингу. Обчислювальний шейдер може читати з довільних буферів зберігання й писати в них, уперше уможливлюючи повноцінні GPGPU-навантаження в браузері:

Обхідні прийоми обчислень WebGL проти нативних обчислень WebGPU: WebGL «GPGPU»: 1. Упакувати дані у float-текстуру RGBA 2. Відмалювати повноекранний чотирикутник із шейдером, що читає цю текстуру 3. Записати результат в іншу текстуру через FBO 4. За потреби зчитати назад через readPixels() → Потребує пакування/розпакування даних у RGBA4, обмежено компонентами float32×4 Обчислення WebGPU: 1. Покласти дані в буфер зберігання (довільне розташування структур) 2. Запустити обчислювальний шейдер 3. Шейдер читає/пише буфер напряму 4. За потреби відобразити буфер назад на CPU → Довільні типи (u32, f32, vec3f, власні структури), довільний доступ, атомарні операції
Підтримка браузерами (2025): WebGPU постачається в Chrome 113+ та Edge 113+. Firefox постачає його за прапорцем (firefox.webgpu). Safari постачає його на macOS 14+ та iOS 18+. Використовуйте navigator.gpu для виявлення доступності. Для старіших браузерів резервом є WebGL2 із transform feedback або обчисленнями через розширення.

2. Архітектура GPU для обчислень

Розуміння моделі виконання GPU є необхідним для написання ефективних обчислювальних шейдерів. GPU — це не швидший CPU, а масово паралельний SIMD-процесор, спроєктований виконувати тисячі потоків одночасно, приховуючи затримку пам'яті через перемикання потоків.

// Спрощена ієрархія GPU (назви від виробників різняться) GPU └── потокові мультипроцесори (SM) / обчислювальні блоки (CU) — наприклад, 80 на RTX 4080 └── ядра CUDA / шейдерні процесори — 128 на SM в Ampere └── Warp / Wave = 32 (NVIDIA) або 64 (AMD) потоки └── Усі потоки у warp виконують ОДНУ І ТУ Ж інструкцію Термінологія WebGPU: робоча група = група потоків, що спільно використовують швидку локальну пам'ять і можуть синхронізуватися бар'єрами Інвокація = окремий потік (одне виконання шейдера) Глобальний ID = унікальний ID серед усіх інвокацій цього диспетчу Локальний ID = ID у межах робочої групи (наприклад, 0..63) Виклик диспетчу: dispatchWorkgroups(x, y, z) → Запускає x×y×z робочих груп → Кожна робоча група виконує workgroup_size потоків → Усього потоків = x × y × z × workgroup_size

Ключовий принцип продуктивності: тримати потоки зайнятими корисною роботою й уникати розгалуження потоків (коли різні потоки беруть різні гілки if/else). Розбіжні потоки в одному warp виконують обидві гілки послідовно, удвічі знижуючи пропускну здатність.

3. WGSL: мова шейдерів WebGPU

WGSL (WebGPU Shading Language) — це статично типізована, орієнтована на безпеку мова шейдерів. На відміну від GLSL/HLSL, вона не має неявних перетворень типів, невизначеної поведінки через неініціалізовані змінні та арифметики вказівників — це навмисно, щоб дозволити безпечне виконання в пісочниці браузера.

// Мінімальний обчислювальний шейдер — поелементно додає два масиви @group(0) @binding(0) var<storage, read> a : array<f32>; @group(0) @binding(1) var<storage, read> b : array<f32>; @group(0) @binding(2) var<storage, read_write> result : array<f32>; @compute @workgroup_size(64) fn main(@builtin(global_invocation_id) gid : vec3u) { let i = gid.x; if (i >= arrayLength(&a)) { return; } // захист меж result[i] = a[i] + b[i]; } Ключові можливості WGSL: @group(g) @binding(b) — слот групи прив'язки var<storage, read> — буфер зберігання лише для читання var<storage, read_write> — буфер зберігання для читання-запису var<workgroup> — спільна пам'ять робочої групи var<uniform> — uniform-буфер (малі сталі лише для читання) @builtin(global_invocation_id) — vec3u: (x + gx*ws, y + gy*ws, z) @builtin(local_invocation_id) — vec3u у межах робочої групи @builtin(workgroup_id) — vec3u: яка це робоча група

Система типів WGSL

// Скалярні типи bool, i32, u32, f32, f16 // f16 потребує прапорця можливості shader-f16 // Векторні типи vec2f / vec2<f32>, vec3f, vec4f vec3i, vec3u // вектори int / uint // Матричні типи mat4x4f = 4 стовпці vec4f (за стовпцями) mat3x4f = 3 стовпці vec4f (3 стовпці, 4 рядки — незвично) // Структура (для розташування буфера) struct Particle { pos: vec3f, vel: vec3f, mass: f32, _pad: f32, // явне вирівнювання під вирівнювання std140/std430 } // Масив у буфері зберігання (розмір під час виконання) var<storage, read_write> particles : array<Particle>; // Доступ: let p = particles[gid.x]; particles[gid.x].vel += dt * force / p.mass;

4. Налаштування обчислювального конвеєра

Робочий процес обчислень WebGPU на JavaScript складається з п'яти кроків: ініціалізація пристрою, компіляція шейдера, створення конвеєра, прив'язка ресурсів і диспетч.

// Крок 1: отримуємо GPU-пристрій const adapter = await navigator.gpu.requestAdapter(); const device = await adapter.requestDevice(); // Крок 2: компілюємо обчислювальний шейдер const shaderModule = device.createShaderModule({ code: wgslString }); // Крок 3: створюємо обчислювальний конвеєр const pipeline = device.createComputePipeline({ layout: 'auto', // автогенерація розташування групи прив'язки з шейдера compute: { module: shaderModule, entryPoint: 'main', }, }); // Крок 4: створюємо буфери зберігання const N = 1_000_000; const bufA = device.createBuffer({ size: N * 4, // N значень float32 = N * 4 байтів usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST, }); const bufResult = device.createBuffer({ size: N * 4, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, }); // Завантажуємо дані в bufA device.queue.writeBuffer(bufA, 0, new Float32Array(N).fill(1.0)); // Крок 5: створюємо групу прив'язки й запускаємо диспетч const bindGroup = device.createBindGroup({ layout: pipeline.getBindGroupLayout(0), entries: [ { binding: 0, resource: { buffer: bufA } }, { binding: 1, resource: { buffer: bufB } }, { binding: 2, resource: { buffer: bufResult } }, ], }); const encoder = device.createCommandEncoder(); const pass = encoder.beginComputePass(); pass.setPipeline(pipeline); pass.setBindGroup(0, bindGroup); pass.dispatchWorkgroups(Math.ceil(N / 64)); // workgroup_size=64 pass.end(); device.queue.submit([encoder.finish()]);

Зчитування результатів назад на CPU

// GPU-буфери не можна відобразити напряму — треба скопіювати в проміжний буфер const readBuf = device.createBuffer({ size: N * 4, usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST, }); const enc2 = device.createCommandEncoder(); enc2.copyBufferToBuffer(bufResult, 0, readBuf, 0, N * 4); device.queue.submit([enc2.finish()]); await readBuf.mapAsync(GPUMapMode.READ); const data = new Float32Array(readBuf.getMappedRange()); console.log(data[0]); // використовуємо результати readBuf.unmap();
Шаблон проміжного буфера: результати обчислень GPU містяться у VRAM (відеопам'яті на GPU). CPU не може напряму читати VRAM. Завжди потрібен проміжний буфер з прапорцем MAP_READ. Виклик mapAsync чекає, доки GPU завершить усю поставлену в чергу роботу й буфер буде перенесено в пам'ять, доступну для читання CPU — це головна точка синхронізації CPU-GPU, і вона зупиняє CPU.

5. Буфери зберігання та групи прив'язки

Буфери зберігання — це основні контейнери даних для обчислювальних шейдерів. Вони містять довільні структури й масиви, підтримують довільний доступ і дозволяють як читання, так і запис. Групи прив'язки організовують зв'язок між буферами, виділеними в JavaScript, і слотами прив'язки шейдера.

// Прапорці використання буфера (можна об'єднувати через OR) GPUBufferUsage.STORAGE // читання/запис обчислювальними/рендер-шейдерами GPUBufferUsage.UNIFORM // uniform-буфер (швидкий, малий, макс. 64 КБ на деякому обладнанні) GPUBufferUsage.COPY_SRC // може бути джерелом copyBufferToBuffer GPUBufferUsage.COPY_DST // може бути приймачем writeBuffer / copyBuffer GPUBufferUsage.MAP_READ // CPU може mapAsync для читання (шаблон проміжного буфера) GPUBufferUsage.MAP_WRITE // CPU може mapAsync для запису (проміжне завантаження) GPUBufferUsage.VERTEX // буфер атрибутів вершин для рендер-конвеєра GPUBufferUsage.INDEX // буфер індексів для рендер-конвеєра // Спільний буфер обчислень+рендерингу (позиції частинок для обох): const particleBuf = device.createBuffer({ size: N * 32, // 2 × vec3f + вирівнювання f16 = 32 байти на Particle usage: GPUBufferUsage.STORAGE // обчислення можуть читати/писати | GPUBufferUsage.VERTEX, // рендер-конвеєр читає як атрибут вершини });

Розташування групи прив'язки

За layout: 'auto' WebGPU виводить розташування групи прив'язки з рефлексії шейдера. Для продакшн-коду з кількома конвеєрами, що використовують спільні ресурси, явні розташування дозволяють повторно використовувати групу прив'язки між конвеєрами — це важливо для зменшення навантаження на CPU за кадр.

6. Робочі групи, спільна пам'ять і синхронізація

Робоча група — це основна одиниця співпраці в обчислювальному шейдері. Потоки в межах робочої групи можуть обмінюватися даними через спільну пам'ять робочої групи — невеликий швидкий буфер (зазвичай ~32 КБ), спільний для всіх інвокацій групи. Доступ набагато швидший, ніж до буфера зберігання (≈у 100 разів менша затримка).

// Паралельна редукція зі спільною пам'яттю (обчислюємо мінімум N float) @group(0) @binding(0) var<storage, read> data : array<f32>; @group(0) @binding(1) var<storage, read_write> result : array<f32>; const WS = 256u; var<workgroup> shared : array<f32, WS>; // спільна пам'ять робочої групи @compute @workgroup_size(WS) fn reduce_min( @builtin(global_invocation_id) gid : vec3u, @builtin(local_invocation_id) lid : vec3u, @builtin(workgroup_id) wid : vec3u, ) { let i = gid.x; let n = arrayLength(&data); // Фаза 1: завантаження з буфера зберігання у швидку спільну пам'ять shared[lid.x] = select(1e38, data[i], i < n); // заповнюємо вихід за межі +inf workgroupBarrier(); // чекаємо, доки ВСІ потоки завершать завантаження // Фаза 2: деревоподібна редукція у спільній пам'яті var stride = WS / 2u; loop { if (stride == 0u) { break; } if (lid.x < stride) { shared[lid.x] = min(shared[lid.x], shared[lid.x + stride]); } workgroupBarrier(); stride /= 2u; } // Потік 0 записує результат робочої групи у вихід if (lid.x == 0u) { result[wid.x] = shared[0]; } }

Бар'єри

workgroupBarrier() // бар'єр пам'яті + виконання в межах робочої групи // ВСІ потоки мають досягти цього виклику, перш ніж будь-який продовжить // синхронізує спільну пам'ять робочої групи ТА пам'ять зберігання storageBarrier() // синхронізує лише записи в буфер зберігання в межах робочої групи // легший — корисний, коли пишеться лише сховище // ⚠️ Виклик workgroupBarrier() усередині умови чи циклу, де // деякі потоки його пропускають, спричиняє невизначену поведінку або зависання GPU. // Усі інвокації в робочій групі МАЮТЬ виконувати ті самі бар'єри.

Вибір розміру робочої групи

Оптимальний розмір робочої групи балансує завантаження обладнання й використання спільної пам'яті. Поширені варіанти:

Операції підгруп (2025): розширення WebGPU subgroups (постачається 2025 року) надає примітиви рівня warp, як-от subgroupAdd, subgroupMin, subgroupBallot. Вони швидші за редукції зі спільною пам'яттю, бо використовують апаратні інструкції shuffle у межах warp, повністю усуваючи записи в спільну пам'ять для найвнутрішнішого рівня редукції.

7. Поширені обчислювальні шаблони

Паралельне відображення масиву

Найпростіший шаблон: одна інвокація на елемент, кожна виконує незалежну роботу. Ідеально для поелементних операцій — досконало паралельно, без потреби в обміні даними.

// диспетч: Math.ceil(N / workgroup_size) робочих груп @compute @workgroup_size(64) fn map(@builtin(global_invocation_id) gid: vec3u) { let i = gid.x; if (i >= arrayLength(&data)) { return; } output[i] = expensive_function(data[i]); }

Префіксна сума (сканування)

Префіксна сума використовується для ущільнення, потокового ущільнення та побудови гістограм. Вона не є тривіально паралельною — кожен елемент залежить від усіх попередніх елементів. Паралельний алгоритм сканування розкладає залежність на O(log N) проходів:

// Паралельна виключна префіксна сума — двофазний підхід // Фаза 1: локальне сканування робочої групи → часткові суми у спільній пам'яті // Фаза 2: сканування часткових сум (рекурсивно) → зсуви на робочу групу // Фаза 3: додавання зсуву робочої групи до кожного елемента у вихідній робочій групі // Це ядро потокового ущільнення на GPU: // «залишити лише елементи, що задовольняють предикат, і записати їх компактно» // Використовується в: смерті/народженні частинок, ущільненні активних променів у трасуванні, // побудові списку колізій широкої фази у фізиці

Атомарні операції

// Атомарні операції в WGSL (лише цілі — у WebGPU немає атомарних float ) var<storage, read_write> counter : atomic<u32>; var<storage, read_write> histogram : array<atomic<u32>, 256>; // У шейдері: atomicAdd(&counter, 1u); // потокобезпечне збільшення let old = atomicMin(&histogram[bin], val); // повертає старе значення atomicStore(&counter, 0u); // скидання в нуль let v = atomicLoad(&counter); // потокобезпечне читання // Немає атомарного float: використовуйте цілочислове кодування // encoded_f32 = bitcast<u32>(f32_value * scale) // Тоді цілочислові atomicMax / atomicMin на закодованих значеннях

8. Сценарії використання: фізика, ML та обробка зображень

Симуляція фізики (системи частинок)

Симуляція частинок WebGPU використовує два буфери зберігання (ping-pong між кадрами) та обчислювальний шейдер на кожен крок фізики. Зчитування на CPU не потрібне, коли той самий буфер прив'язаний як буфер вершин для рендерингу:

// Обчислювальний шейдер на частинку (спрощена гравітація N тіл) @compute @workgroup_size(64) fn simulate(@builtin(global_invocation_id) gid: vec3u) { let i = gid.x; var p = particles_in[i]; var force = vec3f(0.0); for (var j = 0u; j < num_particles; j++) { if (j == i) { continue; } let diff = particles_in[j].pos - p.pos; let d2 = dot(diff, diff) + 0.01; // згладжування force += diff * (G * p.mass * particles_in[j].mass / (d2 * sqrt(d2))); } p.vel += dt * force / p.mass; p.pos += dt * p.vel; particles_out[i] = p; } // 100 000 частинок, 64 ітерації: ~8 мс на RTX 3080 // Та сама симуляція на CPU (однопотоковий JS): ~600 мс

Виведення машинного навчання

Виведення трансформерів і CNN виконується безпосередньо в браузері за допомогою обчислювальних шейдерів WebGPU. Бібліотеки на кшталт TensorFlow.js та ONNX Runtime Web використовують бекенди WebGPU, що реалізують множення матриць (GEMM), згортку й функції активації як обчислювальні шейдери. Квантований LLM на 7 млрд параметрів може працювати зі швидкістю 10-20 токенів/с на ігрових GPU середнього класу через WebGPU.

Обробка зображень

// 2D-згортка з плитками робочої групи (наприклад, розмиття Гаусса) const TILE = 16u; var<workgroup> tile : array<array<vec4f, TILE + 4>, TILE + 4>; @compute @workgroup_size(TILE, TILE) fn blur( @builtin(global_invocation_id) gid: vec3u, @builtin(local_invocation_id) lid: vec3u, ) { // Завантажуємо плитку + кайму (граничні пікселі для перекриття ядра) у спільну пам'ять // Застосовуємо сепарабельне ядро Гаусса 5×5 у спільній пам'яті // Записуємо результат у вихідну текстуру (текстура зберігання лише для запису) } // Перевага над розмиттям у фрагментному шейдері: // - Доступ за плитками уникає надлишкових вибірок текстури між потоками // - Кайма у спільній пам'яті означає, що кожен піксель читається з VRAM рівно один раз // - Без налаштування рендер-конвеєра: лише диспетч + зчитування результату
Перевірка реальної продуктивності: обчислення WebGPU не завжди швидші за CPU. GPU має дуже високу пропускну здатність пам'яті, але й високу затримку запуску роботи — подання командного кодувальника має навантаження 0.1–1 мс. Для навантажень менших за ~10 000 елементів або зі складним розгалуженням, залежним від даних, CPU може бути швидшим. Оптимальна ніша для обчислень WebGPU — великі, регулярні, паралелізовні навантаження: фізика частинок, N тіл, згортка, сортування, редукція.