Pierwsze kroki z GPU Compute w internecie

W tym poście omawiamy eksperymentalny interfejs WebGPU API na przykładach i pomagamy Ci rozpocząć wykonywanie obliczeń równoległych na danych za pomocą procesora graficznego.

François Beaufort

Tło

Jak zapewne wiesz, procesor graficzny (GPU) to elektroniczny podsystem komputera, który pierwotnie był przeznaczony do przetwarzania grafiki. Jednak w ciągu ostatnich 10 lat ewoluowała w kierunku bardziej elastycznej architektury, która umożliwia deweloperom wdrażanie wielu typów algorytmów, a nie tylko renderowanie grafiki 3D, przy jednoczesnym wykorzystaniu unikalnej architektury procesora graficznego. Te możliwości są określane jako obliczenia na GPU, a używanie GPU jako koprocesora do obliczeń naukowych ogólnego przeznaczenia jest nazywane programowaniem na GPU ogólnego przeznaczenia (GPGPU).

Obliczenia na GPU w znacznym stopniu przyczyniły się do niedawnego boomu w dziedzinie uczenia maszynowego, ponieważ konwolucyjne sieci neuronowe i inne modele mogą wykorzystywać tę architekturę do wydajniejszego działania na GPU. Obecna platforma internetowa nie ma możliwości obliczeniowych GPU, dlatego grupa społecznościowa W3C „GPU for the Web” projektuje interfejs API, który udostępni nowoczesne interfejsy API GPU dostępne na większości obecnych urządzeń. Ten interfejs API nazywa się WebGPU.

WebGPU to interfejs API niskiego poziomu, podobnie jak WebGL. Jest bardzo rozbudowany i zawiera dużo informacji, jak zobaczysz. Ale to nic. Zależy nam na wydajności.

W tym artykule skupię się na części WebGPU związanej z obliczeniami na GPU. Szczerze mówiąc, tylko zarysowuję temat, aby umożliwić Ci samodzielne eksperymentowanie. W kolejnych artykułach zajmę się szczegółowo renderowaniem WebGPU (canvas, tekstura itp.).

Dostęp do GPU

Dostęp do GPU w WebGPU jest łatwy. Wywołanie navigator.gpu.requestAdapter() zwraca obietnicę JavaScript, która asynchronicznie rozwiązuje problem z adapterem GPU. Możesz traktować ten adapter jako kartę graficzną. Może być zintegrowana (na tym samym chipie co procesor) lub dyskretna (zwykle karta PCIe, która jest wydajniejsza, ale zużywa więcej energii).

Gdy uzyskasz adapter GPU, wywołaj funkcję adapter.requestDevice(), aby otrzymać obietnicę, która zostanie spełniona przez urządzenie GPU, którego użyjesz do obliczeń na GPU.

const adapter = await navigator.gpu.requestAdapter();
if (!adapter) { return; }
const device = await adapter.requestDevice();

Obie funkcje przyjmują opcje, które pozwalają określić rodzaj adaptera (preferencje dotyczące zasilania) i urządzenia (rozszerzenia, limity). W tym artykule dla uproszczenia użyjemy opcji domyślnych.

Bufor zapisu

Zobaczmy, jak za pomocą JavaScriptu zapisywać dane w pamięci procesora graficznego. Ten proces nie jest prosty ze względu na model piaskownicy używany w nowoczesnych przeglądarkach internetowych.

Z przykładu poniżej dowiesz się, jak zapisać 4 bajty w pamięci bufora dostępnej z GPU. Wywołuje funkcję device.createBuffer(), która przyjmuje rozmiar bufora i jego wykorzystanie. Chociaż flaga użycia GPUBufferUsage.MAP_WRITE nie jest wymagana w przypadku tego konkretnego wywołania, wyraźnie zaznaczmy, że chcemy zapisać dane w tym buforze. W wyniku tego powstaje obiekt bufora GPU zmapowany podczas tworzenia dzięki ustawieniu wartościmappedAtCreation na true. Powiązany bufor surowych danych binarnych można następnie pobrać, wywołując metodę bufora GPU getMappedRange().

Jeśli znasz już ArrayBuffer, zapisywanie bajtów nie będzie dla Ciebie problemem. Użyj TypedArray i skopiuj do niego wartości.

// Get a GPU buffer in a mapped state and an arrayBuffer for writing.
const gpuBuffer = device.createBuffer({
  mappedAtCreation: true,
  size: 4,
  usage: GPUBufferUsage.MAP_WRITE
});
const arrayBuffer = gpuBuffer.getMappedRange();

// Write bytes to buffer.
new Uint8Array(arrayBuffer).set([0, 1, 2, 3]);

W tym momencie bufor GPU jest mapowany, co oznacza, że należy do procesora i jest dostępny do odczytu i zapisu z JavaScriptu. Aby GPU mógł uzyskać do niego dostęp, musi zostać odmapowany. Wystarczy wywołać funkcję gpuBuffer.unmap().

Koncepcja mapowania/niemapowania jest potrzebna, aby zapobiec sytuacjom wyścigu, w których GPU i CPU uzyskują dostęp do pamięci w tym samym czasie.

Odczyt bufora

Teraz zobaczmy, jak skopiować bufor GPU do innego bufora GPU i odczytać go.

Ponieważ zapisujemy dane w pierwszym buforze GPU i chcemy je skopiować do drugiego bufora GPU, wymagana jest nowa flaga użycia GPUBufferUsage.COPY_SRC. Drugi bufor GPU jest tym razem tworzony w stanie niezamapowanym z wartością device.createBuffer(). Jego flaga użycia to GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ, ponieważ będzie używany jako miejsce docelowe pierwszego bufora GPU i będzie odczytywany w JavaScript po wykonaniu poleceń kopiowania GPU.

// Get a GPU buffer in a mapped state and an arrayBuffer for writing.
const gpuWriteBuffer = device.createBuffer({
  mappedAtCreation: true,
  size: 4,
  usage: GPUBufferUsage.MAP_WRITE | GPUBufferUsage.COPY_SRC
});
const arrayBuffer = gpuWriteBuffer.getMappedRange();

// Write bytes to buffer.
new Uint8Array(arrayBuffer).set([0, 1, 2, 3]);

// Unmap buffer so that it can be used later for copy.
gpuWriteBuffer.unmap();

// Get a GPU buffer for reading in an unmapped state.
const gpuReadBuffer = device.createBuffer({
  size: 4,
  usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ
});

GPU jest niezależnym koprocesorem, więc wszystkie polecenia GPU są wykonywane asynchronicznie. Dlatego tworzona jest lista poleceń GPU, które w razie potrzeby są wysyłane partiami. W WebGPU koder poleceń GPU zwracany przezdevice.createCommandEncoder()to obiekt JavaScript, który tworzy partię „buforowanych” poleceń, które zostaną w pewnym momencie wysłane do GPU. Metody w GPUBuffer są natomiast „niebuforowane”, co oznacza, że są wykonywane w sposób niepodzielny w momencie wywołania.

Gdy uzyskasz koder poleceń procesora graficznego, wywołaj copyEncoder.copyBufferToBuffer(), jak pokazano poniżej, aby dodać to polecenie do kolejki poleceń do późniejszego wykonania. Na koniec zakończ kodowanie poleceń, wywołując copyEncoder.finish(), i prześlij je do kolejki poleceń urządzenia GPU. Kolejka odpowiada za obsługę zgłoszeń przesyłanych za pomocą polecenia device.queue.submit() z poleceniami GPU jako argumentami. Spowoduje to atomowe wykonanie wszystkich poleceń przechowywanych w tablicy w odpowiedniej kolejności.

// Encode commands for copying buffer to buffer.
const copyEncoder = device.createCommandEncoder();
copyEncoder.copyBufferToBuffer(
  gpuWriteBuffer /* source buffer */,
  0 /* source offset */,
  gpuReadBuffer /* destination buffer */,
  0 /* destination offset */,
  4 /* size */
);

// Submit copy commands.
const copyCommands = copyEncoder.finish();
device.queue.submit([copyCommands]);

Na tym etapie polecenia kolejki GPU zostały wysłane, ale niekoniecznie wykonane. Aby odczytać drugi bufor GPU, wywołaj funkcję gpuReadBuffer.mapAsync() z parametrem GPUMapMode.READ. Zwraca obietnicę, która zostanie spełniona, gdy bufor GPU zostanie zmapowany. Następnie uzyskaj zamapowany zakres za pomocą funkcji gpuReadBuffer.getMappedRange(), która zawiera te same wartości co pierwszy bufor GPU po wykonaniu wszystkich poleceń GPU w kolejce.

// Read buffer.
await gpuReadBuffer.mapAsync(GPUMapMode.READ);
const copyArrayBuffer = gpuReadBuffer.getMappedRange();
console.log(new Uint8Array(copyArrayBuffer));

Możesz wypróbować ten przykład.

Podsumowując, o operacjach na pamięci buforowej musisz pamiętać o tych kwestiach:

• Bufory GPU muszą zostać odmapowane, aby można było ich używać podczas przesyłania do kolejki urządzeń.
• Po zmapowaniu bufory GPU można odczytywać i zapisywać w JavaScript.
• Bufory procesora graficznego są mapowane, gdy wywoływane są funkcje mapAsync() i createBuffer() z parametrem mappedAtCreation ustawionym na wartość true.

Programowanie cieniowania

Programy działające na procesorze graficznym, które wykonują tylko obliczenia (i nie rysują trójkątów), nazywane są shaderami obliczeniowymi. Są one wykonywane równolegle przez setki rdzeni GPU (które są mniejsze niż rdzenie CPU), które działają razem, aby przetwarzać dane. Dane wejściowe i wyjściowe to bufory w WebGPU.

Aby zilustrować użycie programów cieniujących obliczenia w WebGPU, zajmiemy się mnożeniem macierzy, czyli popularnym algorytmem w uczeniu maszynowym, który przedstawiono poniżej.

Krótko mówiąc, zrobimy to:

1. Utwórz 3 bufory GPU (2 na macierze do pomnożenia i 1 na macierz wynikową).
2. Opisz dane wejściowe i wyjściowe shadera obliczeniowego.
3. Skompiluj kod shadera obliczeniowego.
4. Konfigurowanie potoku obliczeniowego
5. Przesyłanie do procesora graficznego zakodowanych poleceń w ramach zadania wsadowego
6. Odczytywanie bufora GPU macierzy wyników

Tworzenie buforów GPU

Dla uproszczenia macierze będą reprezentowane jako lista liczb zmiennoprzecinkowych. Pierwszy element to liczba wierszy, drugi to liczba kolumn, a pozostałe to rzeczywiste liczby macierzy.

Trzy bufory GPU to bufory pamięci masowej, ponieważ musimy przechowywać i pobierać dane w cieniowaniu obliczeniowym. Wyjaśnia to, dlaczego flagi wykorzystania bufora GPU zawierają GPUBufferUsage.STORAGE. Flaga użycia macierzy wyników ma też wartość GPUBufferUsage.COPY_SRC, ponieważ po wykonaniu wszystkich poleceń kolejki GPU zostanie ona skopiowana do innego bufora w celu odczytu.

const adapter = await navigator.gpu.requestAdapter();
if (!adapter) { return; }
const device = await adapter.requestDevice();


// First Matrix

const firstMatrix = new Float32Array([
  2 /* rows */, 4 /* columns */,
  1, 2, 3, 4,
  5, 6, 7, 8
]);

const gpuBufferFirstMatrix = device.createBuffer({
  mappedAtCreation: true,
  size: firstMatrix.byteLength,
  usage: GPUBufferUsage.STORAGE,
});
const arrayBufferFirstMatrix = gpuBufferFirstMatrix.getMappedRange();
new Float32Array(arrayBufferFirstMatrix).set(firstMatrix);
gpuBufferFirstMatrix.unmap();


// Second Matrix

const secondMatrix = new Float32Array([
  4 /* rows */, 2 /* columns */,
  1, 2,
  3, 4,
  5, 6,
  7, 8
]);

const gpuBufferSecondMatrix = device.createBuffer({
  mappedAtCreation: true,
  size: secondMatrix.byteLength,
  usage: GPUBufferUsage.STORAGE,
});
const arrayBufferSecondMatrix = gpuBufferSecondMatrix.getMappedRange();
new Float32Array(arrayBufferSecondMatrix).set(secondMatrix);
gpuBufferSecondMatrix.unmap();


// Result Matrix

const resultMatrixBufferSize = Float32Array.BYTES_PER_ELEMENT * (2 + firstMatrix[0] * secondMatrix[1]);
const resultMatrixBuffer = device.createBuffer({
  size: resultMatrixBufferSize,
  usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC
});

Układ grupy powiązań i grupa powiązań

Pojęcia układu grupy powiązań i grupy powiązań są specyficzne dla WebGPU. Układ grupy powiązań określa interfejs wejścia/wyjścia oczekiwany przez shader, a grupa powiązań reprezentuje rzeczywiste dane wejścia/wyjścia dla shadera.

W przykładzie poniżej układ grupy powiązań oczekuje dwóch buforów pamięci tylko do odczytu w powiązaniach wpisów o numerach 0 i 1 oraz bufora pamięci w powiązaniu wpisu o numerze 2 dla shadera obliczeniowego. Z kolei grupa wiązań zdefiniowana dla tego układu grupy wiązań przypisuje bufory GPU do wpisów: gpuBufferFirstMatrix do wiązania 0, gpuBufferSecondMatrix do wiązania 1 i resultMatrixBuffer do wiązania 2.

const bindGroupLayout = device.createBindGroupLayout({
  entries: [
    {
      binding: 0,
      visibility: GPUShaderStage.COMPUTE,
      buffer: {
        type: "read-only-storage"
      }
    },
    {
      binding: 1,
      visibility: GPUShaderStage.COMPUTE,
      buffer: {
        type: "read-only-storage"
      }
    },
    {
      binding: 2,
      visibility: GPUShaderStage.COMPUTE,
      buffer: {
        type: "storage"
      }
    }
  ]
});

const bindGroup = device.createBindGroup({
  layout: bindGroupLayout,
  entries: [
    {
      binding: 0,
      resource: {
        buffer: gpuBufferFirstMatrix
      }
    },
    {
      binding: 1,
      resource: {
        buffer: gpuBufferSecondMatrix
      }
    },
    {
      binding: 2,
      resource: {
        buffer: resultMatrixBuffer
      }
    }
  ]
});

Kod shadera obliczeniowego

Kod shadera obliczeniowego do mnożenia macierzy jest napisany w języku WGSL, czyli WebGPU Shader Language, który można łatwo przetłumaczyć na SPIR-V. Poniżej znajdziesz 3 bufory pamięci oznaczone symbolem var<storage>. Program będzie używać firstMatrix i secondMatrix jako danych wejściowych, a resultMatrix jako danych wyjściowych.

Zwróć uwagę, że każdy bufor pamięci ma dekorację binding, która odpowiada temu samemu indeksowi zdefiniowanemu w układach grup wiązań i grupach wiązań zadeklarowanych powyżej.

const shaderModule = device.createShaderModule({
  code: `
    struct Matrix {
      size : vec2f,
      numbers: array<f32>,
    }

    @group(0) @binding(0) var<storage, read> firstMatrix : Matrix;
    @group(0) @binding(1) var<storage, read> secondMatrix : Matrix;
    @group(0) @binding(2) var<storage, read_write> resultMatrix : Matrix;

    @compute @workgroup_size(8, 8)
    fn main(@builtin(global_invocation_id) global_id : vec3u) {
      // Guard against out-of-bounds work group sizes
      if (global_id.x >= u32(firstMatrix.size.x) || global_id.y >= u32(secondMatrix.size.y)) {
        return;
      }

      resultMatrix.size = vec2(firstMatrix.size.x, secondMatrix.size.y);

      let resultCell = vec2(global_id.x, global_id.y);
      var result = 0.0;
      for (var i = 0u; i < u32(firstMatrix.size.y); i = i + 1u) {
        let a = i + resultCell.x * u32(firstMatrix.size.y);
        let b = resultCell.y + i * u32(secondMatrix.size.y);
        result = result + firstMatrix.numbers[a] * secondMatrix.numbers[b];
      }

      let index = resultCell.y + resultCell.x * u32(secondMatrix.size.y);
      resultMatrix.numbers[index] = result;
    }
  `
});

Konfiguracja potoku

Potok obliczeniowy to obiekt, który opisuje operację obliczeniową, którą zamierzamy wykonać. Utwórz go, dzwoniąc pod numer device.createComputePipeline(). Przyjmuje 2 argumenty: utworzony wcześniej układ grupy powiązań i etap obliczeniowy określający punkt wejścia do naszego shadera obliczeniowego (funkcja main WGSL) oraz rzeczywisty moduł shadera obliczeniowego utworzony za pomocą device.createShaderModule().

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

Przesyłanie poleceń

Po utworzeniu grupy powiązań z 3 buforami GPU i potoku obliczeniowego z układem grupy powiązań możemy ich użyć.

Zacznijmy od kodera przepustki obliczeniowej z możliwością programowania za pomocą znaku commandEncoder.beginComputePass(). Użyjemy go do kodowania poleceń GPU, które wykonają mnożenie macierzy. Ustaw potok za pomocą passEncoder.setPipeline(computePipeline) i grupę wiązań o indeksie 0 za pomocą passEncoder.setBindGroup(0, bindGroup). Indeks 0 odpowiada dekoracji group(0) w kodzie WGSL.

Porozmawiajmy teraz o tym, jak ten program cieniujący obliczenia będzie działać na procesorze graficznym. Naszym celem jest równoległe wykonywanie tego programu dla każdej komórki macierzy wyników, krok po kroku. Na przykład w przypadku macierzy wyników o rozmiarze 16 x 32, aby zakodować polecenie wykonania na urządzeniu @workgroup_size(8, 8), wywołamy passEncoder.dispatchWorkgroups(2, 4) lub passEncoder.dispatchWorkgroups(16 / 8, 32 / 8). Pierwszy argument „x” to pierwszy wymiar, drugi „y” to drugi wymiar, a ostatni „z” to trzeci wymiar, który domyślnie ma wartość 1, ponieważ nie jest nam tu potrzebny. W świecie obliczeń na GPU kodowanie polecenia wykonania funkcji jądra na zbiorze danych nazywa się wysyłaniem.

Rozmiar siatki grupy roboczej dla naszego shadera obliczeniowego to (8, 8) w naszym kodzie WGSL. Z tego powodu „x” i „y”, które są odpowiednio liczbą wierszy pierwszej macierzy i liczbą kolumn drugiej macierzy, zostaną podzielone przez 8. Możemy teraz wysłać wywołanie obliczeniowe z parametrem passEncoder.dispatchWorkgroups(firstMatrix[0] / 8, secondMatrix[1] / 8). Liczba siatek grup roboczych do uruchomienia to argumenty dispatchWorkgroups().

Jak widać na powyższym rysunku, każdy shader będzie miał dostęp do unikalnego obiektu builtin(global_invocation_id), który będzie używany do określania, którą komórkę macierzy wyników należy obliczyć.

const commandEncoder = device.createCommandEncoder();

const passEncoder = commandEncoder.beginComputePass();
passEncoder.setPipeline(computePipeline);
passEncoder.setBindGroup(0, bindGroup);
const workgroupCountX = Math.ceil(firstMatrix[0] / 8);
const workgroupCountY = Math.ceil(secondMatrix[1] / 8);
passEncoder.dispatchWorkgroups(workgroupCountX, workgroupCountY);
passEncoder.end();

Aby zakończyć koder przepustki obliczeniowej, wywołaj passEncoder.end(). Następnie utwórz bufor GPU, który będzie służyć jako miejsce docelowe do skopiowania bufora macierzy wyników za pomocą polecenia copyBufferToBuffer. Na koniec zakończ kodowanie poleceń za pomocą znaku copyEncoder.finish() i prześlij je do kolejki urządzenia GPU, wywołując funkcję device.queue.submit() z poleceniami GPU.

// Get a GPU buffer for reading in an unmapped state.
const gpuReadBuffer = device.createBuffer({
  size: resultMatrixBufferSize,
  usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ
});

// Encode commands for copying buffer to buffer.
commandEncoder.copyBufferToBuffer(
  resultMatrixBuffer /* source buffer */,
  0 /* source offset */,
  gpuReadBuffer /* destination buffer */,
  0 /* destination offset */,
  resultMatrixBufferSize /* size */
);

// Submit GPU commands.
const gpuCommands = commandEncoder.finish();
device.queue.submit([gpuCommands]);

Odczytywanie macierzy wyników

Odczytanie macierzy wyników jest tak proste, jak wywołanie gpuReadBuffer.mapAsync() z GPUMapMode.READ i oczekiwanie na rozwiązanie zwracanego obiektu Promise, co oznacza, że bufor GPU jest teraz zmapowany. W tym momencie można uzyskać zamapowany zakres za pomocą funkcji gpuReadBuffer.getMappedRange().

W naszym kodzie wynik zarejestrowany w konsoli JavaScriptu w Narzędziach deweloperskich to „2, 2, 50, 60, 114, 140”.

// Read buffer.
await gpuReadBuffer.mapAsync(GPUMapMode.READ);
const arrayBuffer = gpuReadBuffer.getMappedRange();
console.log(new Float32Array(arrayBuffer));

Gratulacje! Gotowe! Możesz wypróbować próbkę.

Ostatnia sztuczka

Jednym ze sposobów na zwiększenie czytelności kodu jest użycie wygodnej metody getBindGroupLayout potoku obliczeniowego, która pozwala wywnioskować układ grupy powiązań z modułu cieniowania. Dzięki temu nie musisz tworzyć niestandardowego układu grupy powiązań ani określać układu potoku w potoku obliczeniowym, jak widać poniżej.

Ilustracja wartości getBindGroupLayout w przypadku poprzedniej próbki jest dostępna.

 const computePipeline = device.createComputePipeline({
-  layout: device.createPipelineLayout({
-    bindGroupLayouts: [bindGroupLayout]
-  }),
   compute: {

-// Bind group layout and bind group
- const bindGroupLayout = device.createBindGroupLayout({
-   entries: [
-     {
-       binding: 0,
-       visibility: GPUShaderStage.COMPUTE,
-       buffer: {
-         type: "read-only-storage"
-       }
-     },
-     {
-       binding: 1,
-       visibility: GPUShaderStage.COMPUTE,
-       buffer: {
-         type: "read-only-storage"
-       }
-     },
-     {
-       binding: 2,
-       visibility: GPUShaderStage.COMPUTE,
-       buffer: {
-         type: "storage"
-       }
-     }
-   ]
- });
+// Bind group
  const bindGroup = device.createBindGroup({
-  layout: bindGroupLayout,
+  layout: computePipeline.getBindGroupLayout(0 /* index */),
   entries: [

Wyniki dotyczące skuteczności

Jak mnożenie macierzy na GPU wypada w porównaniu z mnożeniem na CPU? Aby się tego dowiedzieć, napisałem opisany wyżej program na procesor. Jak widać na wykresie poniżej, wykorzystanie pełnej mocy GPU wydaje się oczywistym wyborem, gdy rozmiar macierzy jest większy niż 256 x 256.

Ten artykuł to dopiero początek mojej przygody z WebGPU. Wkrótce opublikujemy kolejne artykuły, w których szczegółowo omówimy obliczenia na GPU i działanie renderowania (canvas, tekstura, sampler) w WebGPU.