Pierwsze kroki z GPU Compute w internecie

W tym artykule za pomocą przykładów omawiamy eksperymentalny interfejs WebGPU API i pomagamy Ci zacząć wykonywać obliczenia równoległe danych za pomocą karty graficznej.

François Beaufort

Tło

Jak pewnie wiesz, procesor graficzny (GPU) to elektroniczny subsistema komputera, który pierwotnie był wyspecjalizowany do przetwarzania grafiki. Jednak w ciągu ostatnich 10 lat architektura stała się bardziej elastyczna, dzięki czemu programiści mogą wdrażać wiele typów algorytmów, nie tylko renderować grafikę 3D, korzystając z unikalnej architektury GPU. Te możliwości określa się jako GPU, a korzystanie z GPU jako współprocesora w ogólnych obliczeniach naukowych nazywane jest programowaniem ogólnego przeznaczenia GPU (GPGPU).

Procesory GPU przyczyniły się do ostatniego boomu w dziedzinie uczenia maszynowego, ponieważ konwolucyjne sieci neuronowe i inne modele mogą korzystać z architektury, aby działać wydajniej na procesorach GPU. Obecna platforma internetowa nie ma możliwości obliczeń GPU, dlatego grupa społeczności „GPU for the Web” w W3C opracowuje interfejs API, który udostępnia nowoczesne interfejsy API GPU dostępne na większości obecnych urządzeń. Ten interfejs API nazywa się WebGPU.

WebGPU to niskopoziomowy interfejs API, podobnie jak WebGL. Jak widać, jest ono bardzo skuteczne i dokładne. To nie problem. Interesuje nas skuteczność.

W tym artykule skupię się na części WebGPU dotyczącej przetwarzania przez GPU. Szczerze mówiąc, tylko troszeczkę zagłębię się w tą tematykę, abyś mógł/mogła zacząć samodzielnie eksperymentować. W przyszłych artykułach zajmę się bardziej szczegółowo tematem renderowania WebGPU (płótna, tekstur itp.).

Dostęp do GPU

W WebGPU dostęp do GPU jest łatwy. Wywołanie funkcji navigator.gpu.requestAdapter() zwraca obietnicę JavaScript, która zostanie rozwiązana asynchronicznie za pomocą adaptera GPU. Możesz traktować ten adapter jak kartę graficzną. Może być zintegrowana (na tym samym układzie scalonym co procesor) lub dedykowana (zazwyczaj karta PCIe, która jest wydajniejsza, ale zużywa więcej energii).

Gdy masz już adapter GPU, wywołaj adapter.requestDevice(), aby uzyskać obietnicę, która zostanie rozwiązana z urządzeniem GPU, którego użyjesz do wykonania obliczeń GPU.

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

Obie funkcje mają opcje, które umożliwiają określenie rodzaju adaptera (preferencje dotyczące mocy) i urządzenia (rozszerzenia, limity). Dla uproszczenia użyjemy opcji domyślnych w tym artykule.

Bufor pamięci zapisu

Zobaczmy, jak za pomocą JavaScript zapisać dane do pamięci dla GPU. Ten proces nie jest prosty ze względu na model piaskownicy używany w nowoczesnych przeglądarkach internetowych.

Przykład poniżej pokazuje, jak zapisać 4 bajty w celu buforowania pamięci dostępnej przez GPU. Wywołuje on funkcję device.createBuffer(), która pobiera rozmiar bufora i jego użycie. Mimo że w przypadku tego konkretnego wywołania flaga użycia GPUBufferUsage.MAP_WRITE nie jest wymagana, wyraźnie określmy, że chcemy zapisać dane w tym buforze. W efekcie obiekt bufora procesora GPU jest mapowany podczas tworzenia dzięki ustawieniu parametru mappedAtCreation na wartość true. Następnie powiązany bufor danych binarnych w postaci binarnej można pobrać, wywołując metodę bufora GPU getMappedRange().

Pisanie bajtów jest znajome, jeśli już pracowałeś/pracowałaś z 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]);

Na tym etapie bufor GPU jest mapowany, co oznacza, że należy on do procesora i jest dostępny w trybie odczytu i zapisu z poziomu JavaScriptu. Aby GPU mógł uzyskać do niego dostęp, musi być odmapowany. Można to zrobić, wywołując funkcję gpuBuffer.unmap().

Pojęcie mapowania/niemapowania jest potrzebne, aby zapobiec warunkom wyścigu, w których GPU i procesor uzyskują dostęp do pamięci w tym samym czasie.

Odczyt pamięci bufora

Zobaczmy teraz, jak skopiować bufor GPU do innego bufora GPU i go odczytać.

Piszemy w pierwszym buforze GPU i chcemy skopiować go do drugiego, więc wymagana jest nowa flaga wykorzystania GPUBufferUsage.COPY_SRC. Drugi bufor GPU jest tym razem tworzony w stanie niezamapowanym za pomocą funkcji device.createBuffer(). Jego flaga użycia to GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ, ponieważ będzie on używany jako miejsce docelowe pierwszego bufora GPU i czytany w JavaScript po wykonaniu poleceń kopiowania przez 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
});

Ponieważ GPU jest niezależnym koprocesorem, wszystkie polecenia GPU są wykonywane asynchronicznie. Dlatego w razie potrzeby jest tworzona i wysyłana w partiach lista poleceń GPU. W WebGPU koder poleceń GPU zwracany przez device.createCommandEncoder()to obiekt JavaScript, który tworzy partię „buforowanych” poleceń, które w danym momencie zostaną wysłane do GPU. Z drugiej strony metody w GPUBuffer są „niebuforowane”, co oznacza, że są wykonywane w czasie wywołania.

Gdy masz już koder poleceń GPU, 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 funkcję copyEncoder.finish(), i prześlij je do kolejki poleceń urządzenia GPU. Kolejka odpowiada za obsługę przesyłania za pomocą polecenia device.queue.submit() z poleceniami GPU jako argumentami. Spowoduje to wykonanie w kolejności wszystkich poleceń zapisanych w tablicy.

// 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 zostały wykonane. Aby odczytać drugi bufor GPU, wywołaj gpuReadBuffer.mapAsync() za pomocą GPUMapMode.READ. Zwraca obietnicę, która zostanie spełniona, gdy bufor GPU zostanie zamapowany. Następnie pobierz zmapowany zakres za pomocą gpuReadBuffer.getMappedRange(), który zawiera te same wartości co pierwszy bufor GPU po wykonaniu wszystkich poleceń GPU w kole.

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

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

Oto, co należy pamiętać o operacjach na pamięci buforowej:

• Bufory GPU muszą być odmapowane, aby można było ich używać w przypadku przesyłania do kolejki urządzenia.
• Po zmapowaniu bufory GPU mogą być odczytywane i zapisywane w JavaScript.
• Bufory procesora graficznego są mapowane, gdy wywołane są funkcje mapAsync() i createBuffer() z wartością true w polu mappedAtCreation.

Programowanie shaderów

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 (mniejszych niż rdzenie procesora) współpracujących ze sobą w celu przetwarzania danych. Ich dane wejściowe i wyjściowe są buforami w WebGPU.

Aby zilustrować zastosowanie cieniowania obliczeniowego w WebGPU, zajmiemy się mnożeniem macierzy, co jest często spotykanym algorytmem uczącym się opisanym poniżej.

Krótko mówiąc, zrobimy to w ten sposób:

1. Utwórz 3 bufory GPU (2 dla macierzy do mnożenia i 1 dla macierzy wynikowej).
2. Opisz dane wejściowe i wyjściowe dla cieniowania obliczeniowego
3. Kompilowanie kodu shadera obliczeniowego
4. Konfigurowanie potoku obliczeniowego
5. Przesyłanie zakodowanych poleceń do GPU w grupach
6. Odczytywanie zasobu 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 element to liczba kolumn, a reszta to rzeczywiste liczby macierzy.

Trzy bufory GPU to bufory pamięci, ponieważ musimy przechowywać i pobierać dane w shaderze obliczeniowym. To wyjaśnia, dlaczego flagi wykorzystania bufora GPU zawierają w ich przypadku wartość GPUBufferUsage.STORAGE. Flaga użycia macierzy wyników ma też wartość GPUBufferUsage.COPY_SRC, ponieważ zostanie skopiowana do innego bufora 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ń

Pojęcia układu grupy wiązania i grupy wiązania są specyficzne dla WebGPU. Układ grupy wiązania definiuje interfejs wejścia/wyjścia oczekiwany przez shader, a grupa wiązania reprezentuje rzeczywiste dane wejścia/wyjścia dla shadera.

W przykładzie poniżej układ grupy bind oczekuje 2 buforów pamięci tylko do odczytu w numerowanych elementach bindowania 0 i 1 oraz bufora pamięci w miejscu 2 dla shadera obliczeniowego. Z drugiej strony grupa powiązań, zdefiniowana dla tego układu grupy powiązań, wiąże bufory GPU z wpisami: gpuBufferFirstMatrix z powiązaniem 0, gpuBufferSecondMatrix z powiązaniem 1, a resultMatrixBuffer z powiązaniem 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 języku shaderów WebGPU, który można łatwo przetłumaczyć na SPIR-V. Nie wdając się w szczegóły, poniżej znajdziesz 3 bufory pamięci identyfikowane za pomocą var<storage>. Program będzie używać danych wejściowych firstMatrix i secondMatrix oraz danych wyjściowych resultMatrix.

Pamiętaj, że każdy bufor pamięci ma ozdobę binding, która odpowiada temu samemu indeksowi zdefiniowanemu w deklarowanych powyżej grupach bind i ich układach.

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órą zamierzamy wykonać. Utwórz je, dzwoniąc pod numer device.createComputePipeline(). Przyjmuje 2 argumenty: wcześniej utworzony układ grupy wiązania oraz etap przetwarzania definiujący punkt wejścia naszego shadera obliczeniowego (funkcja WGSL main) i sam moduł shadera obliczeniowego utworzony za pomocą funkcji 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 i potokiem obliczeniowym z układem grupy powiązań przyszedł czas na ich użycie.

Uruchom koder z przesyłką obliczeniową programowalną za pomocą commandEncoder.beginComputePass(). Użyjemy go do zakodowania poleceń GPU, które wykonają mnożenie macierzy. Ustaw potok passEncoder.setPipeline(computePipeline) i jego grupę powiązań na poziomie 0 za pomocą parametru passEncoder.setBindGroup(0, bindGroup). Indeks 0 odpowiada dekoracji group(0) w kodzie WGSL.

Porozmawiajmy teraz o tym, jak ten shader obliczeniowy będzie działać na karcie graficznej. Naszym celem jest równoległe wykonywanie tego programu w przypadku każdej komórki macierzy wyników, krok po kroku. W przypadku macierzy wyników o rozmiarach 16 x 32 na przykład do zakodowania polecenia wykonania na @workgroup_size(8, 8) użyjemy 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 procesorach graficznych kodowanie polecenia do wykonania funkcji jądra na zbiorze danych nazywa się wysyłaniem.

Rozmiar siatki grup roboczych dla naszego shadera obliczeniowego to (8, 8) w kodzie WGSL. W związku z tym wartości „x” i „y”, czyli odpowiednio liczba wierszy pierwszej macierzy i liczba kolumn drugiej macierzy, zostaną podzielone przez 8. Teraz możemy wysłać do Ciebie zapytanie o pomoc dotyczące 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 shader będzie mieć dostęp do unikalnego obiektu builtin(global_invocation_id), który posłuży do określenia, 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ć kodowanie, wywołaj funkcję passEncoder.end(). Następnie utwórz bufor GPU, który będzie używany jako miejsce docelowe do skopiowania bufora macierzy wyników za pomocą polecenia copyBufferToBuffer. Na koniec zakończ kodowanie poleceń za pomocą copyEncoder.finish() i prześlij je do kolejki urządzenia GPU, wywołując 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]);

Czytanie macierzy wyników

Odczytywanie macierzy wyników jest tak proste, jak wywołanie funkcji gpuReadBuffer.mapAsync() z parametrem GPUMapMode.READ i oczekiwanie na rozwiązanie obietnicy zwracanej przez tę funkcję, co oznacza, że bufor GPU jest już zamapowany. W tym momencie można uzyskać zmapowany zakres za pomocą funkcji gpuReadBuffer.getMappedRange().

W naszym kodzie wynik zapisany w konsoli JavaScriptu DevTools 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 zapoznać się z próbką.

Ostatni trik

Jednym ze sposobów ułatwienia odczytania kodu jest użycie przydatnej metody getBindGroupLayout w pipeline’ie obliczeniowym, aby wywnioskować układ grupy bindowania z modułu shadera. Dzięki temu nie musisz tworzyć niestandardowego układu grupy wiązania ani określać układu potoku w potoku obliczeniowym, jak widać poniżej.

getBindGroupLayout w przypadku poprzedniego przykładu 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

Jakie są różnice między mnożeniem macierzy na procesorze graficznym a na procesorze CPU? Aby to sprawdzić, napisałem program opisany właśnie dla CPU. 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 tylko początek mojej przygody z poznawaniem WebGPU. Wkrótce opublikujemy więcej artykułów, w których znajdziesz więcej informacji o obliczeniach na procesorze graficznym oraz o tym, jak działa renderowanie (płótno, tekstura, próbnik) w WebGPU.