Ця стаття спробує навчити вас основам WebGPU.

WebGPU - це API, яке дозволяє робити дві основні речі:

1. Малювати трикутники/точки/лінії на текстури

2. Запускати обчислення на графічному процесорі

Оце і все!

Усе інше, в технології WebGPU, залежить від вас. Це як вивчення комп’ютерної мови. Наприклад JavaScript, Rust або C++. Спочатку ви вивчаєте основи, а потім можете творчо використовувати ці основи для вирішення своєї проблеми.

WebGPU — це API надзвичайно низького рівня. Хоча ви і можете створити кілька невеликих прикладів, але для багатьох програм, імовірно, знадобиться великий обсяг коду та серйозна організація даних. Наприклад, three.js, який підтримує WebGPU, складається з ~600 тисяч мінімізованих рядків JavaScript коду, і це лише його основна частина, яка не включає в себе завантажувачі, елементи керування, постобробку та багато інших функцій. Подібним чином бібліотека TensorFlow із серверною частиною на WebGPU становить ~500 тисяч мінімізованого JavaScript коду.

Справа в тому, що якщо ви хочете просто намалювати щось на екрані, то в такому разі краще обрати бібліотеку, яка уже надає цю величезну кількість коду, яку вам довелося б писати, роблячи це самостійно.

З іншого боку, якщо у вас дуже особливий випадок, чи ви хочете змінити існуючу бібліотеку, або вам просто цікаво, як усе це працює, тоді в цьому разі продовжуйте читати!

Починаємо

Досить важко вирішити з чого почати. На певному рівні, WebGPU дуже проста система. Все, що вона робить це запускає 3 типи функцій. Вершинний шейдер, фрагментний шейдер та обчислювальний шейдер.

Вершинний шейдер обчислює вершини. Цей шейдер повертає розташування цих вершин. Для кожної групи з 3 вершин, вершинний шейдер повертає трикутник намальований між цих 3 позицій. ^[1]

Фрагментний шейдер обчислює кольори ^[2]. Коли трикутник намальовано, для кожного пікселя графічний процесор викликає фрагментний шейдер. Цей шейдер в свою чергу повертає значення кольору.

Обчислювальний шейдери це функції більш загального призначення. Фактично це просто функція, яку ви викликаєте і кажете “виконай це N кількість разів”. Графічний процесор передає номер ітерації кожного разу, коли він викликає вашу функцію, щоб ви могли використати це число для чогось унікального для кожної ітерації.

Якщо сильно примружитись, можна вважати ці функції схожими на ті, які ми передаємо в array.forEach чи array.map. Функції, які ви запускаєте на графічному процесорі це такі ж звичайні функції, як і JavaScript функції. Частина, яка відрізняється, це те, що вони запускаються саме на графічному процесорі, тож для їхнього запуску вам потрібно скопіювати усі необхідні дані на графічний процесор у вигляді буферів та текстур та повертати з них дані у тих же буферах та текстурах. Вам потрібно зазначати в цих функціях в яких прив’язках чи розташуваннях ця функція буде шукати за даними. А в самому JavaScript коді вам потрібно прив’язати буфери та текстури, які містять ваші дані, до точок прив’язки та розташувань. Зробивши це, ви просите графічний процесор виконати цю функцію.

Можливо тут допоможе малюнок. Ось спрощена діаграма налаштувань WebGPU для малювання трикутника з допомогою вершинного шейдера та фрагментного шейдера.

На що варто звернути увагу на цій діаграмі

• Наявність Пайплайну (Pipeline). Він містить в собі вершинний та фрагментний шейдер, які будуть виконані графічним процесором. Пайплайн також може містити обчислювальний шейдер.

• Шейдери посилаються на ресурси (буфери, текстури, семплери) опосередковано через Групи прив’язки (Bind Groups).

• Пайплайн визначає атрибути, які посилаються на буфери опосередковано через внутрішній стан.

• Атрибути витягують дані з буферів і передають їх у вершинний шейдер.

• Вершинний шейдер може передавати дані в фрагментний шейдер.

• Фрагментний шейдер записує текстури опосередковано через render pass description.

Для запуску шейдерів на графічному процесорі, вам необхідно створити усі ці ресурси і налаштувати стан. Створення ресурсів відносно просте завдання. Одна важлива річ полягає в тому, що більшість ресурсів WebGPU не можна змінити після створення. Ви можете змінити їхній вміст, але не їхній розмір, використання, формат, тощо… Якщо вам потрібно змінити якусь з цих речей, ви мусите створити новий ресурс і знищити старий.

Деякі з цих станів налаштовується створенням буферів команд. Буфери команд це буквально те, про що нам говорить їхня назва. Це буфери, які містять в собі команди. Ви створюєте кодеки. Ці кодеки перетворюють команди в буфери команд. Далі ви завершуєте (finish) ваш кодек, а він повертає вам буфер команд, який він створив. Далі ви можете надіслати (submit) ций буфер на виконання для того, щоб WebGPU запустив ці команди.

Ось псевдокод для кодування буфера команд, поряд з відображенням цього буфера.

encoder = device.createCommandEncoder()
// намалювати дещо
{
  pass = encoder.beginRenderPass(...)
  pass.setPipeline(...)
  pass.setVertexBuffer(0, …)
  pass.setVertexBuffer(1, …)
  pass.setIndexBuffer(...)
  pass.setBindGroup(0, …)
  pass.setBindGroup(1, …)
  pass.draw(...)
  pass.end()
}
// намалювати дещо інше
{
  pass = encoder.beginRenderPass(...)
  pass.setPipeline(...)
  pass.setVertexBuffer(0, …)
  pass.setBindGroup(0, …)
  pass.draw(...)
  pass.end()
}
// обчислити щось
{
  pass = encoder.beginComputePass(...)
  pass.beginComputePass(...)
  pass.setBindGroup(0, …)
  pass.setPipeline(...)
  pass.dispatchWorkgroups(...)
  pass.end();
}
commandBuffer = encoder.finish();

Створивши буфер команд, ви можете надіслати (submit) його на виконання.

device.queue.submit([commandBuffer]);

Діаграма вище показує позиції команд draw в буфері команд. Виконання цих команд встановить внутрішній стан (internal state), а потім вкаже графічному процесору виконати вершинний шейдер (і опосередковано фрагментний шейдер). Команда dispatchWorkgroup скаже графічному процесору виконати обчислювальний шейдер.

Сподіваюся, це дало вам певне уявлення про стан, який вам потрібно створити. Як згадувалося вище, WebGPU має 2 основні функції:

1. Малювати трикутники/точки/лінії на текстури

2. Запускати обчислення на графічному процесорі

Ми розглянемо невеликий приклад виконання кожної з цих функцій. Інші статті покажуть різні способи надсилання даних до цих функцій. Зверніть увагу, що цей приклад буде дуже спрощений. Нам потрібно побудувати фундамент з цих основ. Пізніше ми покажемо, як використовувати це для речей, які люди зазвичай роблять із графічним процесором, як-от 2D-графіка, 3D-графіка тощо…

Малюємо трикутник на текстурі

WebGPU може малювати трикутники на текстурах. Для цілей цієї статті, текстура це двовимірний чотирикутник з пікселів.^[3] Елемент <canvas> репрезентує собою текстуру на вебсторінці. В WebGPU ми можемо попросити в цього елемента його текстуру і тоді намалювати щось на ції текстурі.

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

1. Вершинні шейдери

   Вершинні шейдери - це функції, які вираховують позиції вершин для малювання трикутників/ліній/точок.

2. Фрагментні шейдери

   Фрагментні шейдери - це функції, які вираховують колір (або щось інше) для кожного пікселя, який буде намальований під час малювання трикутників/ліній/точок.

Давайте почнемо з дуже маленької WebGPU програми для малювання трикутника.

Спершу нам потрібне полотно для відображення цього трикутника.

<canvas></canvas>

Далі нам потрібен <script> тег для для зберігання нашого JavaScript коду.

<canvas></canvas>
+<script type="module">

... тут буде javascript ...

+</script>

Увесь JavaScript код, який буде нижче, буде зберігатись в цьому тезі.

WebGPU - це асинхронний API, тож його легше використовувати з асинхронними функціями. Ми починаємо з того, що робимо запит за адаптером, а потім уже за пристроєм з цього адаптера.

async function main() {
  const adapter = await navigator.gpu?.requestAdapter();
  const device = await adapter?.requestDevice();
  if (!device) {
    fail('цей код потребує браузера, який підтримує WebGPU');
    return;
  }
}
main();

Наведений вище код досить зрозумілий. Спершу, ми робимо запит за адаптером використовуючи ?. оператора опціонального ланцюжка для того, щоб якщо об’єкт navigator.gpu не існує, то в adapter буде встановлене значення undefined. Якщо він таки існує, то ми викликаємо requestAdapter. Цей метод повертає свій результат асинхронно тож нам потрібно використати ключове слово await. Цей адаптер уособлює собою конкретний графічний процесор. Деякі пристрої можуть мати декілька графічних процесорів.

Ми робимо запит за пристроєм з цього адаптера і знову використовуємо оператор ?. для того, щоб якщо адаптера немає, то в змінну device буде встановлене значення undefined.

Якщо значення device дорівнює undefined, то швидше за все у користувача старий браузер.

Далі ми знаходимо наше полотно і створюємо webgpu контекст для нього. Це дозволить нам отримати текстуру для візуалізації. Ця текстура буде використана для відображення полотна на вебсторінці.

  // Отримуємо WebGPU контекст з полотна і налаштовуємо його
  const canvas = document.querySelector('canvas');
  const context = canvas.getContext('webgpu');
  const presentationFormat = navigator.gpu.getPreferredCanvasFormat();
  context.configure({
    device,
    format: presentationFormat,
  });

Знову ж таки, наведений вище код досить зрозумілий. Ми отримуємо контекст "webgpu" з полотна. Запитуємо у системи, якому формату полотна вона надає перевагу. Це буде або "rgba8unorm", або "bgra8unorm". Насправді не настільки важливо, що це за формат, але запитуючи його ми робимо певні речі швидшими для користувацької системи.

Ми передаємо це значення як format в webgpu контекст нашого полотна роблячи виклик метода configure. Ми також передаємо туди змінну device, яка пов’язує це полотно з пристроєм, який ми щойно створили.

Далі ми створюємо шейдерний модуль. Шейдерний модуль містить в собі один або більше шейдерів. В нашому випадку, ми створюємо 1 вершинний шейдер і один фрагментний шейдер.

  const module = device.createShaderModule({
    label: 'our hardcoded red triangle shaders',
    code: `
      @vertex fn vs(
        @builtin(vertex_index) vertexIndex : u32
      ) -> @builtin(position) vec4f {
        let pos = array(
          vec2f( 0.0,  0.5),  // top center
          vec2f(-0.5, -0.5),  // bottom left
          vec2f( 0.5, -0.5)   // bottom right
        );

        return vec4f(pos[vertexIndex], 0.0, 1.0);
      }

      @fragment fn fs() -> @location(0) vec4f {
        return vec4f(1.0, 0.0, 0.0, 1.0);
      }
    `,
  });

Шейдери пишуться мовою програмування, яка називається WebGPU Shading Language (WGSL) і часто вимовляється як віґ-сіль (wig-sil). WGSL - це строго типізована мова, яку ми спробуємо описати в більшій кількості деталей в іншій статті. Наразі я сподіваюся, що з невеликими поясненнями ви зможете зрозуміти деякі основи.

Вище ми бачимо, що функція під назвою vs оголошена з атрибутом @vertex. Це позначає її як функцію вершинного шейдера.

      @vertex fn vs(
        @builtin(vertex_index) vertexIndex : u32
      ) -> @builtin(position) vec4f {
         ...

Ця функція приймає один параметр, який ми назвали vertexIndex. vertexIndex має тип u32, що означає 32-бітне беззнакове число. Цей параметр отримує своє значення з вбудованої змінної, яку називають vertex_index. vertex_index це наче номер ітерації, схожий на змінну index в цьому JavaScript коді Array.map(function(value, index) { ... }). Якщо ми скажемо графічному процесору виконати цю функцію 10 разів викликаючи метод draw, то першого разу vertex_index буде 0, другого разу - 1, третього разу - 2 і так далі^[4]

Ми визначили нашу vs функцію, як таку, що повертає значення типу vec4f, яке в свою чергу є вектором з чотирьох 32-бітних чисел з рухомою комою. Думаєте про це значення, як про масив з чотирьох чисел, або обʼєкт з 4 властивостями, схожий на цей {x: 0, y: 0, z: 0, w: 0}. Це значення буде присвоєне вбудованій змінній position. В режимі “triangle-list”, кожних 3 рази виконання цього вершинного шейдера буде намальовано трикутник, що з’єднує ті 3 значення, які ми повертаємо в вбудовану змінну position.

Позиції в WebGPU потрібно повертати у значеннях простору відсікання (clip space), де X починається з -1.0 в лівій частині простору і закінчується значенням +1.0 в правій частині простору, а Y починається з -1.0 внизу і закінчується +1.0 зверху. Ці не залежить від розміру текстури, на якій ми малюємо.

Функція vs оголошує масив з трьох значень типу vec2f. Кожне таке значення містить два 32-бітних числа з рухомою комою.

        let pos = array(
          vec2f( 0.0,  0.5),  // top center
          vec2f(-0.5, -0.5),  // bottom left
          vec2f( 0.5, -0.5)   // bottom right
        );

Нарешті, вона використовує vertexIndex для того, щоб повернути одне з 3 значень цього масиву. Оскільки функція вимагає повернути саме 4 числа та через те, що значення нашого масиву мають тип vec2f, наш код надає два додаткових значення 0.0 та 1.0 до значення, яке ми повертаємо.

        return vec4f(pos[vertexIndex], 0.0, 1.0);

Шейдерний модуль також оголошує функцію fs з атрибутом @fragment, який робить її фрагментним шейдером.

      @fragment fn fs() -> @location(0) vec4f {

Ця функція не приймає жодного параметра і повертає в змінну location(0) значення з типом vec4f. Це означає, що вона буде записувати значення в першу ціль візуалізації. Пізніше, цією ціллю візуалізації ми зробимо текстуру нашого полотна.

        return vec4f(1, 0, 0, 1);

Код повертає 1, 0, 0, 1, що позначає червоний колір. Кольори в WebGPU зазвичай описуються числовим значенням в межах від 0.0 до 1.0. Чотири значення вище відповідають червоному, зеленому, синьому та альфа каналу.

Під час того, як графічний процесор растеризує певний трикутник (малює його з допомогою пікселів), він викликає фрагментний шейдер, щоб отримати значення кольору для кожного пікселя. В нашому випадку, ми просто повертаємо червоний колір.

Ще одна річ, на яку варто звернути увагу, це поле label. Майже кожному об’єкту, який ви можете створити в WebGPU, можна передавати значення поля label. Це значення не обов’язкове, але додавання цієї позначки вважається кращою практикою. Причина полягає в тому, що коли ви отримуєте помилку, більшість реалізацій WebGPU виведе повідомлення про помилку, яке міститиме цю позначку та інші речі пов’язані з помилкою.

У звичайних програмах у вас було би від 100 до 1000 буферів, текстур, шейдерних модулів, пайплайнів, та іншого… Якщо б ви отримала повідомлення про помилку, схоже на це "WGSL syntax error in shaderModule at line 10", то в якому зі 100 шейдерних модулів ви б шукали причину цієї помилки? Натомість, якщо ви передасте в поле label якесь значення, то отримаєте помилку схожу на цю "WGSL syntax error in shaderModule('our hardcoded red triangle shaders'), яка набагато корисніша та збереже вам купу часу у відстеженні проблеми.

Тепер, коли ми створили шейдерний модуль, ми можемо створити пайплайн візуалізації:

  const pipeline = device.createRenderPipeline({
    label: 'our hardcoded red triangle pipeline',
    layout: 'auto',
    vertex: {
      module,
      entryPoint: 'vs',
    },
    fragment: {
      module,
      entryPoint: 'fs',
      targets: [{ format: presentationFormat }],
    },
  });

У цьому випадку немає багато деталей, щоб детально їх розглянути. Ми встановили для параметра layout значення 'auto', що вказує WebGPU отримати макет даних із шейдерів. Проте ми не використовуємо тут жодних даних.

Потім ми вказуємо пайплайну візуалізації, що потрібно використовувати функцію vs з нашого шейдерного модуля для вершинного шейдера та функцію fs для нашого фрагментного шейдера. В випадку фрагментного шейдера ми також повідомляємо йому формат першої цілі візуалізації. “Ціль візуалізації” позначає текстуру, яку ми будемо рендерити. Коли ми створюємо пайплайну ми маємо вказати формат текстури, яку ми будемо використовувати для рендерингу цього пайплайну.

Елемент під індексом 0 з масиву targets відповідає змінній @location(0), яку ми зазначили в типі повернення нашого фрагментного шейдера. Пізніше, ми встановимо в цю ціль текстуру для нашого полотна.

Далі ми підготуємо GPURenderPassDescriptor, який опише, в які текстури ми хочемо малювати і як їх використати.

  const renderPassDescriptor = {
    label: 'our basic canvas renderPass',
    colorAttachments: [
      {
        // view: <- to be filled out when we render
        clearValue: [0.3, 0.3, 0.3, 1],
        loadOp: 'clear',
        storeOp: 'store',
      },
    ],
  };  

GPURenderPassDescriptor містить в собі масив colorAttachments, який перелічує текстури, в які ми будемо рендерити і як їх до них ставитись. Поки ми відкладемо на потім опис тієї текстури, в яку ми хочемо рендерити. Наразі ми встановили в clearValue значення напівтемного сірого, та визначили loadOp і storeOp. loadOp: 'clear' уточнює, що текстура повинна бути очищена до значення clearValue перед малюванням. Іншим значенням цього поля може бути 'load', що вказує на те, що поточний вміст текстури повинен бути завантажений на графічний процесор, щоб малювати поверх того, що уже в цій текстурі. storeOp: 'store' вказує на те, що потрібно зберігати результати малювання. Ми можемо передати сюди також значення 'discard', що буде викидати будь-які результати малювання. Ми пояснимо те, чому ми можемо захотіти таку поведінку в іншій статті.

Тепер настав час рендерингу.

  function render() {
    // отримуємо поточну текстуру з контексту нашого полотна
    // і встановлюємо її, як текстуру в яку потрібно малювати
    renderPassDescriptor.colorAttachments[0].view =
        context.getCurrentTexture().createView();

    // вказуємо кодеку команд почати їхнє кодування 
    const encoder = device.createCommandEncoder({ label: 'our encoder' });

    // make a render pass encoder to encode render specific commands
    const pass = encoder.beginRenderPass(renderPassDescriptor);
    pass.setPipeline(pipeline);
    pass.draw(3);  // викликаємо наш вершинний шейдер 3 рази
    pass.end();

    const commandBuffer = encoder.finish();
    device.queue.submit([commandBuffer]);
  }

  render();

Спершу ми викликаємо context.getCurrentTexture(), щоб отримати текстуру, яка з’явиться на полотні. Виклик createView встановлює область відображення в конкретну частину текстури, але викликаючи цей метод без параметрів ми отримаємо стандартну область відображення, що нам і потрібно в цьому випадку. Наразі, єдиним елементом масиву colorAttachments є область відображення текстури з нашого полотна, яке ми отримала з контексту створеного на початку цієї статті. Знову ж таки, елемент під індексом 0 з масиву colorAttachments відповідає змінній @location(0), яку ми зазначили в типі повернення нашого фрагментного шейдера.

Далі ми створюємо кодек команд. Кодек команд використовується, щоб створити буфер команд. Ми використовуємо його для кодування команд і “надсилання” буфера команд, який він створив, на виконання.

Потім ми використовуємо кодек команд, щоб створити кодек render-pass’у викликаючи beginRenderPass. Кодек render-pass’у - це особливий кодек, для створення команд, які пов’язані з рендерингом. Ми передаємо йому наш renderPassDescriptor, щоб вказати йому на текстуру, в яку ми плануємо рендерити.

Ми викликаємо setPipeline для встановлення нашого пайплайну і далі вказуємо йому виконати наш вершинний шейдер 3 рази, викликаючи метод draw. Стандартною поведінкою в цьому випадку є: кожних три рази, коли виконується наший вершинний шейдер малюється один трикутник, який сполучає 3 значення, які повернулись в цей момент з нашого вершинного шейдера.

Ми завершуємо роботу render-pass’у і після цього завершуємо процес кодування команд. Цим ми отримуємо буфер команд, який уособлює ті кроки, які ми щойно вказали. Зрештою, ми надсилаємо цей буфер команд в чергу на виконання.

Коли командна draw буде виконана, то стан нашої програми буде таким.

Ми не маємо тут жодних текстур, буферів чи груп прив’язки, але ми точно маємо пайплайн, вершинний шейдер, фрагментний шейдер і render pass descriptor, який вказує нашому шейдеру на текстуру нашого полотна.

А ось і результат.

Важливо наголосити, що такі функції як setPipeline та draw тільки додають команди в буфер команд. Вони насправді не виконують їх. Ці команди виконуються тільки тоді, коли ми надсилаємо цей буфер на чергу виконання нашого пристрою.

WebGPU дістає кожних 3 вершини, які повертаються з вершинного шейдера і використовує їх для растеризації трикутника. Він робить це з допомогою визначення пікселів, які перебувають в середині цього трикутника. Далі він викликає фрагментний шейдер для визначення кольору кожного з цих пікселів.

Уявіть собі текстуру розміром 15х11 пікселів, в яку ми хочемо відрендерити наш трикутник. Це виглядатиме ось так:

Отож, ми побачили невеликий робочий приклад WebGPU програми. Зрозуміло, що жорстке кодування трикутника в середині шейдера не є дуже гнучким підходом. Ми повинні мати певні способи передачі даних в шейдер і ми обов’язково розглянемо їх в наступних статтях. Ключові моменти з коду вище, які варто підкреслити:

• WebGPU лише виконує шейдери. Код, який робитиме щось корисне, це ваша відповідальність.
• Шейдери описуються в шейдерному модулі, а потім перетворюються на пайплайн.
• WebGPU здатен малювати трикутники.
• WebGPU малює в текстурах (в нашому випадку це текстура нашого вебполотна).
• WebGPU кодує команди, а потім надсилає їх на виконання.

Запускаємо обчислення на графічному процесорі

Давайте створимо базовий приклад обчислень на графічному процесорі.

Ми почнемо з того ж самого коду, який отримує WebGPU пристрій.

async function main() {
  const adapter = await navigator.gpu?.requestAdapter();
  const device = await adapter?.requestDevice();
  if (!device) {
    fail('цей код потребує браузера, який підтримує WebGPU');
    return;
  }

Далі ми створюємо шейдерний модуль.

  const module = device.createShaderModule({
    label: 'doubling compute module',
    code: `
      @group(0) @binding(0) var<storage, read_write> data: array<f32>;

      @compute @workgroup_size(1) fn computeSomething(
        @builtin(global_invocation_id) id: vec3u
      ) {
        let i = id.x;
        data[i] = data[i] * 2.0;
      }
    `,
  });

Спершу ми оголошуємо змінну data типу storage, яка повинна мати можливість читання та запису. First, we declare a variable called data of type storage that we want to be able to both read from and write to.

      @group(0) @binding(0) var<storage, read_write> data: array<f32>;

Ми оголошуємо її тип як array<f32>, що означає масив 32-бітних чисел з рухомою комою. Ми вказуємо, що цей масив буде в локації прив’язки під індексом 0 (binding(0)), яка буде в групі прив’язок під індексом 0 (@group(0)).

Далі ми оголошуємо функцію computeSomething з атрибутом @compute, який робить її обчислювальний шейдер.

      @compute @workgroup_size(1) fn computeSomething(
        @builtin(global_invocation_id) id: vec3u
      ) {
        ...

Обчислювальний шейдер вимагає від нас оголосити також значення розміру робочої групи (workgroup size), значення якої ми пояснимо пізніше. Наразі ми встановимо це значення в 1 з допомогою атрибута @workgroup_size(1). Ми оголошуємо один параметр функції id з типом vec3u. vec3u це три беззнакових 32-бітних числа. Як і в нашому вершинному шейдері вище, це номер ітерації. Різниця в тому, що в обчислювальному шейдері номер ітерації є тривимірним (має 3 значення). Ми вказуємо на те, що id отримає своє значення з вбудованої змінної global_invocation_id.

Ви можете приблизно вважати, що обчислювальні шейдери працюють якось так. Це дуже спрощений приклад, але поки цього достатньо.

// pseudo code
function dispatchWorkgroups(width, height, depth) {
  for (z = 0; z < depth; ++z) {
    for (y = 0; y < height; ++y) {
      for (x = 0; x < width; ++x) {
        const workgroup_id = {x, y, z};
        dispatchWorkgroup(workgroup_id)
      }
    }
  }
}

function dispatchWorkgroup(workgroup_id) {
  // from @workgroup_size in WGSL
  const workgroup_size = shaderCode.workgroup_size;
  const {x: width, y: height, z: depth} = workgroup_size;
  for (z = 0; z < depth; ++z) {
    for (y = 0; y < height; ++y) {
      for (x = 0; x < width; ++x) {
        const local_invocation_id = {x, y, z};
        const global_invocation_id =
            workgroup_id * workgroup_size + local_invocation_id;
        computeShader(global_invocation_id)
      }
    }
  }
}

Оскільки ми встановили @workgroup_size(1), то цей псевдокод перетворюється на:

// pseudo code
function dispatchWorkgroups(width, height, depth) {
  for (z = 0; z < depth; ++z) {
    for (y = 0; y < height; ++y) {
      for (x = 0; x < width; ++x) {
        const workgroup_id = {x, y, z};
        dispatchWorkgroup(workgroup_id)
      }
    }
  }
}

function dispatchWorkgroup(workgroup_id) {
  const global_invocation_id = workgroup_id;
  computeShader(global_invocation_id)
}

Нарешті ми використаємо поле x параметру id для індексації масиву data і помножимо кожне число на 2.

        let i = id.x;
        data[i] = data[i] * 2.0;

Вище i є лише першим із 3 чисел ітерації.

Тепер, коли ми створили шейдер, нам потрібно створити пайплайн.

  const pipeline = device.createComputePipeline({
    label: 'doubling compute pipeline',
    layout: 'auto',
    compute: {
      module,
      entryPoint: 'computeSomething',
    },
  });

Тут вказуємо лише на те, що ми використаємо етап compute з шейдерного модуля module, який ми щойно створили, і оскільки там є тільки одна точка запуску з атрибутом @compute WebGPU буде знати, що саме її потрібно викликати. В поле layout встановлюємо значення 'auto', що вказує WebGPU самому визначити цей це значення з шейдерів. ^[5]

Далі нам потрібні якісь дані.

  const input = new Float32Array([1, 3, 5]);

Ці дані існують лише в JavaScript. Щоб WebGPU використав їх, нам потрібно створити буфер, який існуватиме на графічному процесорі, і скопіювати ті дані в цей буфер.

  // створюємо буфер на графічному процесорі
  const workBuffer = device.createBuffer({
    label: 'work buffer',
    size: input.byteLength,
    usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST,
  });
  // копіюємо наші дані в цей буфер
  device.queue.writeBuffer(workBuffer, 0, input);

Ми викликаємо device.createBuffer, щоб створити буфер. size це розмір буфера в байтах. В цьому випадку він буде 12, оскільки розмір масиву Float32Array з трьох елементів буде дорівнювати 12 байтам. Якщо ви не знайомі з Float32Array та типізованими масивами, то гляньте цю статтю.

Кожен створений WebGPU буфер має визначити значення поля usage. Існує багато прапорців, які ми можемо передати туди, але не всі з них можна використовувати разом. Тут ми вказуємо, що хочемо використовувати як storage передаючи прапорець GPUBufferUsage.STORAGE. Це робить його сумісним з var<storage,...> у шейдері. Далі, ми хочемо мати змогу копіювати дані в цей буфер, тому ми передаємо прапорець GPUBufferUsage.COPY_DST. І нарешті, ми хочемо мати змогу копіювати дані з цього буфера, тому ми передаємо прапорець GPUBufferUsage.COPY_SRC.

Зверніть увагу на те, що ви не можете напряму читати вміст WebGPU буфера з допомогою JavaScript. Натомість ви повинні “відобразити (map)” їх, що є просто іншим способом отримати доступ до буфера з допомогою WebGPU, через те, що цей буфер може бути зайнятим, або існувати тільки на графічному процесорі.

WebGPU буфер, який можна відобразити в JavaScript, неможливо використати для чогось іншого. Іншими словами, ми не можемо відобразити буфер, який ми створили вище і навіть якщо ми спробуємо додати прапорець для відображення, то ми отримаємо помилку про те, що цей буфер не сумісний з типом використання STORAGE.

Тож, для того щоб отримати результати обчислень, нам потрібен інший буфер. Після запуску цих обчислень, ми скопіюємо цей наш буфер в новий результуючий буфер і передамо йому прапорець, який дозволить відображати його.

  // створюємо буфер на графічному процесорі, щоб отримати копію результату
  const resultBuffer = device.createBuffer({
    label: 'result buffer',
    size: input.byteLength,
    usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST
  });

MAP_READ означає, що ми хочемо мати можливість відобразити цей масив для читання даних з нього.

Для того, щоб наш шейдер дізнався про буфер з яким він повинен буде працювати, ми маємо створити групу прив’язки (bindGroup).

  // налаштовуємо групу прив’язки для того,
  // щоб вказати шейдеру на буфер, з яким потрібно буде працювати
  const bindGroup = device.createBindGroup({
    label: 'bindGroup for work buffer',
    layout: pipeline.getBindGroupLayout(0),
    entries: [
      { binding: 0, resource: { buffer: workBuffer } },
    ],
  });

Ми отримуємо схему (layout) для цієї групи прив’язки з нашого пайплайну. Далі ми встановлюємо значення поля entries. Цифра 0 в pipeline.getBindGroupLayout(0) відповідає значенню @group(0) в нашому шейдері. Значення {binding: 0 ... в полі entries відповідає значенню @group(0) @binding(0) в шейдері.

Тепер ми можемо почати кодувати команди.

  // кодуємо команди для наших розрахунків
  const encoder = device.createCommandEncoder({
    label: 'doubling encoder',
  });
  const pass = encoder.beginComputePass({
    label: 'doubling compute pass',
  });
  pass.setPipeline(pipeline);
  pass.setBindGroup(0, bindGroup);
  pass.dispatchWorkgroups(input.length);
  pass.end();

Тут ми створюємо кодек команд. Запускаємо обчислювальний pass. Встановлюємо пайплайн, а потім групу прив’язок. Цифра 0 в pass.setBindGroup(0, bindGroup) відповідає значенню @group(0) в шейдері. Далі ми викликаємо метод dispatchWorkgroups, якому в цьому випадку ми передаємо значення input.length, що дорівнює числу 3 і вказує WebGPU запустити обчислювальний шейдер 3 рази. Після цього ми зупиняємо наш pass.

Таким можна уявити діаграму стану WebGPU після виконання методу dispatchWorkgroups.

Після закінчення обчислень ми просимо WebGPU скопіювати дані буфера workBuffer в буфер resultBuffer.

  // кодуємо команду для копіювання результатів в resultBuffer
  encoder.copyBufferToBuffer(workBuffer, 0, resultBuffer, 0, resultBuffer.size);

Тепер ми можемо завершити (finish) наш кодек, отримати буфер команд та надіслати цей буфер на виконання.

  // завершуємо кодування і відправляємо команди на виконання
  const commandBuffer = encoder.finish();
  device.queue.submit([commandBuffer]);

Далі ми відображаємо результуючий буфер і отримуємо копію його даних.

  // Read the results
  await resultBuffer.mapAsync(GPUMapMode.READ);
  const result = new Float32Array(resultBuffer.getMappedRange());

  console.log('input', input);
  console.log('result', result);

  resultBuffer.unmap();

Для відображення буфера ми викликаємо метод mapAsync і чекаємо (await) на його завершення. Після цього, ми можемо викликати resultBuffer.getMappedRange(), який поверне ArrayBuffer з вмістом усього буфера. Ми передаємо це значення в конструктор типізованого масиву Float32Array і після цього маємо змогу переглядати ці значення. Одна важлива деталь - ArrayBuffer, який повертається з методу getMappedRange, дійсний тільки до того моменту поки ми не викличемо метод unmap. Після того, як ми викликали unmap, значення довжини масиву ArrayBuffer буде встановлене в 0 і його вміст перестане бути доступний.

Запустивши увесь цей код ми побачимо наші результати - усі числа перемножені на 2.

По пояснимо, як по справжньому користуватись обчислювальними шейдерами в іншій статті. Наразі, сподіваюсь, ви отримали розуміння того, що робить WebGPU. УСЕ ІНШЕ ЦЕ ВАША ВІДПОВІДАЛЬНІСТЬ! Думайте про WebGPU, як про іншу мову програмування. Він надає нам декілька базових функцій та залишає простір для вашої креативності.

Те, що вершинні, фрагментні та обчислювальні шейдери запускаються на графічному процесорі, робить програмування на WebGPU дуже особливим. Графічний процесор може мати більше 10000 підпроцесорів, які потенційно можуть виконати більше 10000 обчислень паралельно, що, ймовірно, на 3 або більше порядків величин більше, ніж ваш центральний процесор може виконати паралельно.

Проста зміна розміру полотна

Перед тим, як ми перейдемо до інших статей, варто повернутись до прикладу з малюванням трикутника і додати базову підтримку зміни розміру полотна. Зміна розміру полотна це тема, яка може мати багато тонкощів тож існує ціла окрема стаття про це. Наразі ж давайте додамо базову підтримку цього.

Спочатку, ми додамо трішки CSS, щоб заповнити усю сторінку нашим полотном.

<style>
html, body {
  margin: 0;       /* remove the default margin          */
  height: 100%;    /* make the html,body fill the page   */
}
canvas {
  display: block;  /* make the canvas act like a block   */
  width: 100%;     /* make the canvas fill its container */
  height: 100%;
}
</style>

Цей CSS код сам по собі розтягне наше полотно на весь екран, але не змінить його розширення, тож ви можете помітити, що краї трикутник нагадують маленькі блоки. Особливо, якщо збільшити приклад, натиснувши на кнопку повного екрану.

Тег <canvas> тут має розширення 300x150 пікселів. Ми б хотіли підігнати розширення полотна до розміру йог відображення. Один хороший спосіб - використати ResizeObserver. Ви створюєте ResizeObserver і передаєте йому функцію, яку він має запустити кожного разу, коли об’єкт, за яким ви встановили спостереження, змінює свій розмір. Ми можете самостійно вказати, за яким елементом цей об’єкт повинен вести спостереження.

    ...
-    render();

+    const observer = new ResizeObserver(entries => {
+      for (const entry of entries) {
+        const canvas = entry.target;
+        const width = entry.contentBoxSize[0].inlineSize;
+        const height = entry.contentBoxSize[0].blockSize;
+        canvas.width = Math.max(1, Math.min(width, device.limits.maxTextureDimension2D));
+        canvas.height = Math.max(1, Math.min(height, device.limits.maxTextureDimension2D));
+        // re-render
+        render();
+      }
+    });
+    observer.observe(canvas);

В коді вище, ми проходимось по масиву entries, але тут має бути тільки один елемент, оскільки ми стежимо тільки за полотном. Нам потрібно обмежити розмір полотна до найбільшого значення, яке може підтримати наш пристрій. В іншому випадку WebGPU почне генерувати помилки про те, що ми намагаємось створити завелику текстуру. Також, нам потрібно запевнитись, що це значення не буде нулем, бо знову отримаємо помилки. Перегляньте цю довшу статтю для розуміння деталей.

Ми викликаємо метод render, щоб заново намалювати трикутник з новим розширенням. Ми також видалили старий виклик render тому, що він більше не потрібен. ResizeObserver запустить свій колбек щонайменше один раз, щоб сповістити про розмір елементу в момент, коли за ним почали слідкувати.

Текстура з новим розміром створюється в момент, коли ми викликаємо context.getCurrentTexture() в методі render, тож більше тут немає що змінювати.

В наступних статтях, ми розглянемо різні методи передання даних в шейдери.

• міжетапні змінні
• юніформи
• буфери для зберігання
• вершинні буфери
• текстури
• константи

Потім ми розглянемо основи WGSL.

В такому порядку ми підемо від найпростішого до найскладнішого. Міжетапні змінні не потребують додаткових зовнішнього налаштування для пояснення. Ми можемо побачити, як їх використовувати, лише з допомогою змін до WGSL коду, який ми використовували вище. Юніформи фактично є глобальними змінними і використовуються у всіх 3 видах шейдерів (вершинні, фрагментні та обчислювальні). Перехід від юніформних буферів до буферів зберігання досить тривіальний. Вершинні буфери використовуються лише у вершинних шейдерах. Вони складніші, оскільки вимагають опису макета даних у WebGPU. Текстури є найскладнішими, оскільки вони мають безліч типів і варіантів.

Я трохи хвилююся, що спочатку ці статті будуть нудними. Не соромтеся перестрибувати з однієї на іншу, якщо цього хочете. Просто пам’ятайте, якщо ви чогось не розумієте, вам, ймовірно, потрібно прочитати або переглянути основи розглянуті в цій статті. Опанувавши основи, ми почнемо розглядати справжні техніки роботи з WebGPU.

Ще одна річ. Усі приклади програм можна редагувати на цій же веб-сторінці. Крім того, їх можна легко експортувати в jsfiddle, codepen і навіть stackoverflow. Просто натисніть «Експорт».

async function start() {
  if (!navigator.gpu) {
    fail('this browser does not support WebGPU');
    return;
  }

  const adapter = await navigator.gpu.requestAdapter();
  if (!adapter) {
    fail('this browser supports webgpu but it appears disabled');
    return;
  }

  const device = await adapter?.requestDevice();
  device.lost.then((info) => {
    console.error(`WebGPU device was lost: ${info.message}`);

    // 'reason' will be 'destroyed' if we intentionally destroy the device.
    if (info.reason !== 'destroyed') {
      // try again
      start();
    }
  });
  
  main(device);
}
start();

function main(device) {
  ... do webgpu ...
}