Pierwsze kroki z GPU Compute w internecie

W tym poście omawiamy eksperymentalny interfejs WebGPU API na przykładach i pomagamy zaczniesz wykonywać obliczenia równoległe dotyczące danych z wykorzystaniem GPU.

François Beaufort

Tło

Jak zapewne wiecie, procesor graficzny (GPU) to elektroniczny podsystemu w komputerze, który pierwotnie służył do przetwarzania danych grafiki. Jednak w ciągu ostatnich 10 lat przekształciła się w kierunku większej elastyczności, architektura pozwalająca programistom wdrażać wiele rodzajów algorytmów, nie tylko grafiki 3D, wykorzystując unikalną architekturę GPU. Te możliwości określa się jako GPU, a użycie GPU jako współprocesor do ogólnego przeznaczenia naukowego jest nazywany ogólnym przeznaczeniem programowania GPU (GPGPU).

W ostatnim okresie bojów systemów uczących się nastąpił gwałtowny wzrost liczby GPU, ponieważ splotowe sieci neuronowe i inne modele mogą wykorzystać do wydajniejszego działania na procesorach GPU. Dzięki obecnej platformie internetowej ponieważ nie ma możliwości wykorzystania procesorów graficznych, opracowana przez W3C „GPU dla sieci” Grupa społeczności projektuje interfejs API pod kątem udostępniania nowoczesnych interfejsów API GPU, urządzeń. Ten interfejs API nazywa się WebGPU.

WebGPU to niskopoziomowy interfejs API, taki jak WebGL. To bardzo intensywne i szczegółowe, zobaczysz. Ale to nie problem. Interesuje nas skuteczność reklam.

W tym artykule skupię się na GPU w ramach WebGPU. szczerość, wierzę tylko, że mogę zacząć grać na swoim własnych. Zajmę się szczegółowo renderowaniem WebGPU (canvas, texture, itp.) w kolejnych artykułach.

Dostęp do GPU

W WebGPU dostęp do GPU jest prosty. Dzwonię pod navigator.gpu.requestAdapter() zwraca obietnicę JavaScript, która zostanie asynchronicznie rozwiązana za pomocą GPU przejściówkę. Potraktuj ten adapter jako kartę graficzną. Można ją zintegrować (na tym samym układzie scalonym co procesor) lub dyskretny (zwykle jest to karta PCIe o większej wydajności, ale zużywa więcej energii).

Gdy będziesz już mieć adapter GPU, wywołaj funkcję adapter.requestDevice(), aby uzyskać obietnicę który będzie rozwiązywany na procesorze GPU, które będzie służyć do wykonywania obliczeń.

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

Obie funkcje przyjmują opcje, które pozwalają dokładnie określić rodzaj i urządzenia (rozszerzenia, limity). W przypadku atrybutu dla uproszczenia użyjemy opcji domyślnych opisanych w tym artykule.

Pamięć bufora zapisu

Zobaczmy, jak użyć JavaScriptu do zapisywania danych w pamięci przez GPU. Ten nie jest prosty ze względu na używany we współczesnym internecie model piaskownicy, przeglądarki.

Przykład poniżej pokazuje, jak zapisać 4 bajty w celu buforowania dostępnej pamięci z procesora graficznego. Wywołuje on funkcję device.createBuffer(), która przyjmuje rozmiar i jego wykorzystaniu. Mimo że flaga użycia GPUBufferUsage.MAP_WRITE jest nie jest wymagane w przypadku tego konkretnego wywołania, wyraźnie zaznaczmy, że chcemy napisać do tego bufora. Powoduje to utworzenie obiektu bufora GPU zmapowanego podczas tworzenia dzięki mappedAtCreation ma wartość Prawda. Wtedy powiązany bufor nieprzetworzonych danych binarnych może można pobierać, wywołując metodę bufora GPU getMappedRange().

Zapisywanie bajtów jest znajome, jeśli znasz już funkcję ArrayBuffer. 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 mapowany jest bufor GPU, co oznacza, że należy on do procesora. jest dostępna w trybie odczytu/zapisu z JavaScriptu. Aby GPU mógł uzyskać do niego dostęp, musi być niezmapowany. To tak proste, jak wywołanie gpuBuffer.unmap().

Koncepcja map lub niezmapowanych jest potrzebna, by zapobiec wyścigom w warunkach, w których GPU i korzystać z pamięci procesora w tym samym czasie.

Odczyt pamięci bufora

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

Ponieważ zapisujemy w pierwszym buforze GPU, chcemy go skopiować do drugiej Bufor GPU, wymagana jest nowa flaga wykorzystania GPUBufferUsage.COPY_SRC. Druga Bufor GPU jest tym razem tworzony w stanie niezmapowanym z użyciem: device.createBuffer() Flaga użycia tej karty to GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ, ponieważ będzie ona używana jako miejsce docelowe pierwszego GPU. buforować i odczytywać w JavaScripcie 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
});

Procesor graficzny jest niezależnym współprocesorem, więc wszystkie jego polecenia są wykonywane asynchronicznie. Dlatego tworzy się i wysyła listę poleceń GPU wsadów. W WebGPU koder poleceń GPU zwracany przez funkcję device.createCommandEncoder()to obiekt JavaScript, który tworzy wsad „buforowany” które będą w pewnym momencie wysyłane do GPU. Metody włączone Z kolei GPUBuffer są „niebuforowane”, co oznacza, że działają atomowo. w momencie, gdy są wzywani.

Gdy masz już koder poleceń GPU, wywołaj copyEncoder.copyBufferToBuffer() jak pokazano poniżej, aby dodać to polecenie do kolejki poleceń i wykonać je później. Na koniec dokończ polecenia kodowania, wywołując copyEncoder.finish() i przesyłając do kolejki poleceń urządzenia GPU. Kolejka odpowiada za obsługę przesłane w device.queue.submit() z użyciem poleceń GPU jako argumentów. Powoduje to atomowe wykonywanie wszystkich poleceń zapisanych w tablicy w określonej 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]);

W tym momencie polecenia kolejki GPU zostały wysłane, ale niekoniecznie są one wykonywane. Aby odczytać drugi bufor GPU, wywołaj funkcję gpuReadBuffer.mapAsync() z GPUMapMode.READ Zwraca obietnicę, która zostanie zrealizowana, gdy bufor GPU będzie zmapowane. Następnie pobierz zmapowany zakres dla argumentu gpuReadBuffer.getMappedRange(), który zawiera te same wartości, co pierwszy bufor GPU po wszystkich poleceniach GPU w kolejce zostały wykonane.

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

Możesz wypróbować ten fragment.

W skrócie musisz pamiętać o operacjach buforowania pamięci:

• Aby można było używać ich podczas przesyłania kolejki urządzenia, bufory GPU muszą być niezmapowane.
• Po zmapowaniu bufory GPU można odczytywać i zapisywać w JavaScripcie.
• Bufory GPU są mapowane, gdy mapAsync() i createBuffer() z atrybutem Metoda mappedAtCreation z wartością Prawda jest wywoływana.

Programowanie oparte na cieniowaniu

Programy działające z procesorem graficznym, które wykonują tylko obliczenia (i nie rysują). trójkąty) są nazywane programami do cieniowania obliczeń. Są one wykonywane równolegle przez setki rdzeni GPU (mniejszych od rdzeni procesora) działających razem w celu procesora i skalowalnych danych. Ich dane wejściowe i wyjściowe są buforami w WebGPU.

Aby pokazać zastosowanie cieniowania obliczeniowego w WebGPU, zagramy w matrycy, mnożenia, czyli typowego algorytmu w systemach uczących się, który pokazano poniżej.

W skrócie:

1. Utwórz trzy bufory GPU (dwa bufory dla macierzy mnożenia i jeden dla macierzy macierz wyników)
2. Opisz dane wejściowe i wyjściowe dla cieniowania obliczeniowego
3. Kompilowanie kodu cieniowania Compute
4. Konfigurowanie potoku obliczeniowego
5. Zbiorcze przesyłanie zakodowanych poleceń do GPU
6. Odczyt bufora GPU macierzy wyników

Tworzenie buforów GPU

Dla uproszczenia macierze będą przedstawione jako lista pływająca. liczby punktów. Pierwszym z nich jest liczba wierszy, , a reszta to rzeczywiste liczby macierzy.

Trzy bufory GPU są buforami pamięci masowej, ponieważ musimy przechowywać i pobierać dane w do cieniowania obliczeniowego. To wyjaśnia, dlaczego flagi wykorzystania bufora GPU zawierają GPUBufferUsage.STORAGE. Flaga użycia macierzy wyników również zawiera GPUBufferUsage.COPY_SRC, ponieważ zostanie skopiowany do innego bufora przez do odczytu po wykonaniu wszystkich poleceń kolejki GPU.

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
});

Powiąż układ grupy i grupę powiązań

Koncepcje układu grupy powiązań i grupy powiązań są specyficzne dla WebGPU. Powiązanie układ grupy definiuje interfejs wejścia/wyjścia oczekiwany przez program do cieniowania, grupa powiązań reprezentuje rzeczywiste dane wejściowe/wyjściowe dla cieniowania.

W przykładzie poniżej układ grupy powiązań oczekuje 2 buforów pamięci masowej tylko do odczytu w wiązania numerowanych wpisów 0, 1 oraz bufor pamięci masowej na poziomie 2 dla cieniowania obliczeniowego. Z drugiej strony grupa powiązań, zdefiniowana dla tego układu grupy powiązań, wiąże GPU buforuje do wpisów: gpuBufferFirstMatrix do powiązania 0, gpuBufferSecondMatrix do powiązania 1 i resultMatrixBuffer z powiązanie 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 cieniowania Compute

Kod programu do cieniowania obliczeniowego służący do mnożenia macierzy jest napisany w języku WGSL, Język WebGPU Shader, który można łatwo przetłumaczyć jako SPIR-V. Bez poniżej trzech buforów pamięci masowej dzięki var<storage>. Program będzie używać firstMatrix i secondMatrix jako: i resultMatrix jako dane wyjściowe.

Pamiętaj, że każdy bufor pamięci ma ustawioną dekorację binding odpowiadającą ten sam indeks zdefiniowany w układach grup powiązań i grupach powią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 faktycznie opisuje operację obliczeniową które przedstawimy. Utwórz ją, dzwoniąc pod numer device.createComputePipeline(). Przyjmuje 2 argumenty: utworzony wcześniej układ grupy powiązań oraz funkcję obliczeniową etap określający punkt wejścia naszego programu do cieniowania obliczeń (funkcja WGSL main) oraz rzeczywisty moduł cieniowania 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 instancji grupy powiązań z naszymi 3 buforami GPU oraz z układem grupy powiązań, czas ich użyć.

Zacznijmy od programowalnego kodera Compute Pass, commandEncoder.beginComputePass() Wykorzystamy go do kodowania poleceń GPU który wykona mnożenie macierzy. Ustaw swój potok za pomocą passEncoder.setPipeline(computePipeline) i jego grupa powiązań w indeksie 0 z passEncoder.setBindGroup(0, bindGroup) Indeks 0 odpowiada Dekoracja group(0) w kodzie WGSL.

Teraz omówmy, jak ten program do cieniowania będzie działać z układem graficznym. Nasze celem jest uruchomienie tego programu równolegle dla każdej komórki macierzy wyników, krok po kroku. Dla macierzy wyników o rozmiarach 16 x 32 na przykład do zakodowania polecenia wykonania, na @workgroup_size(8, 8) nazywaliśmy się passEncoder.dispatchWorkgroups(2, 4) lub passEncoder.dispatchWorkgroups(16 / 8, 32 / 8). Pierwszy argument „x”. oznacza pierwszy wymiar, drugi „y” jest drugi wymiar, a najnowszy „z” to trzeci wymiar, który ma domyślną wartość 1, bo nie jest on tu potrzebny. W świecie obliczeniowym GPU kodowanie polecenia wykonującego funkcję jądra na zbiorze danych nazywa się dyspozycją.

Rozmiar siatki grup roboczych dla naszego mechanizmu cieniowania obliczeniowego wynosi (8, 8) w WGSL w kodzie. Z tego powodu „x” i „y” , które są odpowiednio liczbą wierszy zostanie podzielona pierwsza macierz i liczba kolumn drugiej macierzy przed 8. Możemy więc wysłać włączenie obliczeniowe passEncoder.dispatchWorkgroups(firstMatrix[0] / 8, secondMatrix[1] / 8) liczba siatek grup roboczych do uruchomienia to argumenty dispatchWorkgroups().

Jak widać na rysunku powyżej, każdy program do cieniowania ma dostęp do unikalnego builtin(global_invocation_id) obiekt, który będzie używany do określania, który wynik komórki macierzy do obliczeń.

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 Compute Pass, wywołaj passEncoder.end(). Następnie utwórz Bufor GPU używany jako miejsce docelowe do skopiowania bufora macierzy wyników z copyBufferToBuffer Na koniec dokończ polecenia kodowania, używając copyEncoder.finish() i prześlij je do kolejki urządzenia GPU, wywołując device.queue.submit() za pomocą poleceń 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 tablicy wyników jest tak proste, jak wywołanie funkcji gpuReadBuffer.mapAsync() przy użyciu GPUMapMode.READ i oczekiwanie na rozwiązanie obietnicy zwrotu, co wskazuje, bufor GPU jest teraz mapowany. Na tym etapie można pobrać zakres z funkcją gpuReadBuffer.getMappedRange().

Wynik zarejestrowany w konsoli JavaScript Narzędzi deweloperskich w naszym kodzie 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 wykorzystać sampel.

Ostatnia sztuczka

Jednym ze sposobów na zwiększenie czytelności kodu jest użycie poręcznej Metoda getBindGroupLayout potoku obliczeniowego w celu wywnioskowania grupy powiązań z modułu cieniowania. Dzięki temu nie musisz tworzyć niestandardowy układ grupy powiązań i określenie układu potoku w obliczeniach Jak widać poniżej.

Ilustracja przedstawiająca getBindGroupLayout dla 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 działa mnożenie macierzy na GPU w porównaniu z działaniem CPU? Aby to sprawdzić, napisałem program opisany właśnie dla CPU. Jak możesz, co widać na wykresie poniżej, wykorzystanie pełnych możliwości GPU wydaje się oczywistym wyborem. gdy rozmiar macierzy jest większy niż 256 x 256.

Ten artykuł to tylko początek mojej przygody z poznawaniem WebGPU. Oczekuj więcej artykułów z bardziej szczegółowymi informacjami o GPU Compute i o tym, jak renderowanie (canvas, texture, sampler) działa w WebGPU.