W tym poście omawiamy eksperymentalny interfejs API WebGPU z użyciem przykładów i ułatwia rozpoczęcie pracy z wykonywaniem obliczeń równoległych danych za pomocą GPU.
Wprowadzenie
Jak zapewne wiesz, procesor graficzny (GPU) to podsystem elektroniczny w komputerze, który pierwotnie był wyspecjalizowany do przetwarzania grafiki. Jednak w ciągu ostatnich 10 lat system rozwinął się w kierunku bardziej elastycznej architektury, dzięki czemu deweloperzy mogą wykorzystywać unikalną architekturę GPU, nie tylko renderowanie grafiki 3D, ale także implementowanie wielu rodzajów algorytmów. Funkcje te nazywane są procesorami graficznymi (GPU), a użycie GPU jako koprocesora do ogólnego przeznaczenia obliczeń naukowych jest nazywane programowaniem GPU (GPUU) ogólnego przeznaczenia.
Obliczenia w GPU przyczyniły się w znacznym stopniu do ostatniego boomu związanego z systemami uczącymi się, ponieważ splotowe sieci neuronowe i inne modele mogą korzystać z tej architektury, aby wydajniej działać na procesorach graficznych. Obecnie platforma Web Platform nie ma możliwości procesora graficznego (GPU), dlatego grupa społeczności W3C „GPU for the Web” pracuje nad projektem API, który udostępnia nowoczesne interfejsy API korzystające z GPU, które są dostępne na większości współczesnych urządzeń. Nazwa tego interfejsu API to WebGPU.
WebGPU to interfejs API niskiego poziomu, taki jak WebGL. Jak zobaczycie, jest to bardzo mocne i szczegółowe. Ale nie martw się. Liczy się skuteczność.
W tym artykule skupię się na GPU w WebGPU, a przy okazji – tylko z głowy, aby umożliwić Ci rozpoczęcie własnej gry. W następnych artykułach zagłębię się w temat renderowania opartego na WebGPU (odbitki na płótnie, tekstury itd.).
Dostęp do GPU
Dostęp do GPU jest łatwy w WebGPU. Wywołanie navigator.gpu.requestAdapter() zwraca obietnicę JavaScriptu, która zostanie asynchronicznie rozpoznana za pomocą adaptera GPU. Potraktuj ten adapter jak kartę graficzną. Może być zintegrowana (z tym samym układem scalonym co procesor) lub dyskretna (zwykle jest to karta PCIe, która jest bardziej wydajna, ale zużywa więcej mocy).
Po zainstalowaniu adaptera GPU wywołaj adapter.requestDevice(), aby uzyskać obietnicę rozwiązania problemu z urządzeniem GPU, którego użyjesz do obliczeń GPU.
const adapter = await navigator.gpu.requestAdapter();
if (!adapter) { return; }
const device = await adapter.requestDevice();
Obie funkcje korzystają z opcji, które pozwalają wybrać konkretny rodzaj zasilacza (preferencje dotyczące mocy) i urządzenia (rozszerzenia, limity). Dla uproszczenia w tym artykule użyjemy opcji domyślnych.
Zapisuj w pamięci bufora
Zobaczmy, jak za pomocą JavaScriptu zapisywać dane w pamięci GPU. Proces ten nie jest łatwy ze względu na model piaskownicy stosowany w nowoczesnych przeglądarkach.
Przykład poniżej pokazuje, jak zapisać 4 bajty w celu buforowania pamięci dostępnej przez GPU. Wywołuje on device.createBuffer(), który bierze pod uwagę rozmiar bufora i jego wykorzystanie. Mimo że w przypadku tego wywołania flaga użytkowania GPUBufferUsage.MAP_WRITE nie jest wymagana, wyraźnie zaznaczmy, że chcemy zapisać ją w tym buforze. Powoduje to powstanie obiektu bufora GPU zmapowanego podczas tworzenia, bo mappedAtCreation ma wartość Prawda. Następnie powiązany bufor nieprzetworzonych danych binarnych można pobrać, wywołując metodę bufora GPU getMappedRange().
Zapisywanie bajtów jest znane, 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]);
Na tym etapie bufor GPU jest zmapowany, co oznacza, że należy on do CPU i jest dostępny w zakresie odczytu/zapisu z poziomu JavaScriptu. Aby procesor graficzny miał dostęp do niego, musi być niezmapowany. To tak proste jak wywołanie gpuBuffer.unmap().
Koncepcja mapowania/niezmapowania jest potrzebna, aby uniknąć wyścigów, gdy GPU i CPU mają dostęp do pamięci w tym samym czasie.
Odczytywanie pamięci bufora
Zobaczmy teraz, jak skopiować bufor GPU do innego bufora GPU i go odczytać.
Ponieważ zapisujemy dane w pierwszym buforze GPU i chcemy go skopiować do drugiego bufora GPU, wymagana jest nowa flaga wykorzystania GPUBufferUsage.COPY_SRC. Drugi bufor GPU jest tym razem tworzony w niezmapowanym stanie za pomocą device.createBuffer(). Jego flaga wykorzystania to GPUBufferUsage.COPY_DST |
GPUBufferUsage.MAP_READ, ponieważ będzie używana jako miejsce docelowe pierwszego bufora GPU i będzie odczytywana 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
});
Procesor graficzny jest niezależnym współprocesorem, dlatego wszystkie polecenia GPU są wykonywane asynchronicznie. Dlatego istnieje lista poleceń dotyczących GPU. W WebGPU koder poleceń GPU zwracany przez device.createCommandEncoder() to obiekt JavaScript, który tworzy grupę poleceń „buffered” (buforowanych), które będą w pewnym momencie wysyłane do GPU. Z kolei metody w GPUBuffer są „niebuforowane”, co oznacza, że są wykonywane atomowo w momencie ich wywołania.
Gdy masz koder poleceń GPU, wywołaj copyEncoder.copyBufferToBuffer() w sposób pokazany poniżej, aby dodać to polecenie do kolejki poleceń i wykonać je później.
Na koniec zakończ kodowanie poleceń, wywołując copyEncoder.finish() i przesyłając je do kolejki poleceń urządzenia z GPU. Kolejka odpowiada za obsługę zgłoszeń przesłanych przez device.queue.submit() z użyciem poleceń GPU jako argumentów.
Spowoduje to atomowe wykonywanie wszystkich poleceń zapisanych 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]);
W tym momencie polecenia kolejki GPU zostały wysłane, ale niekoniecznie wykonane.
Aby odczytać drugi bufor GPU, wywołaj gpuReadBuffer.mapAsync() za pomocą polecenia GPUMapMode.READ. Zwraca obietnicę, która zniknie po zmapowaniu bufora GPU. Następnie pobierz zmapowany zakres z funkcją gpuReadBuffer.getMappedRange(), który 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 fragment.
Krótko mówiąc, o operacjach związanych z pamięcią buforującą należy pamiętać o tych kwestiach:
- Aby można było korzystać z buforów GPU podczas przesyłania kolejki urządzenia, trzeba usunąć ich mapowanie.
- Po zmapowaniu buforów GPU można odczytywać i zapisywać bufory w JavaScripcie.
- Bufory GPU są mapowane, gdy wywoływane są wywołania funkcji
mapAsync()icreateBuffer()z wartościąmappedAtCreationi wartością prawda.
Programowanie Shader
Programy działające na GPU, które wykonują tylko obliczenia (i nie rysują trójkątów), to tak zwane programy do cieniowania. Są one wykonywane równolegle przez setki rdzeni GPU (mniejszych niż rdzenie procesora), które działają razem, aby przetwarzać dane. Ich dane wejściowe i wyjściowe są buforami w WebGPU.
Aby zilustrować wykorzystanie cieniowania obliczeniowego w WebGPU, zajmiemy się mnożeniem macierzy, czyli popularnym algorytmem systemów uczących się przedstawionym poniżej.
Oto, co będziemy robić:
- Utwórz 3 bufory GPU (dwa na potrzeby mnożenia macierzy i drugi na macierz wyników)
- Opisz dane wejściowe i wyjściowe dla cieniowania Compute
- Kompilowanie kodu do cieniowania w Compute
- Konfigurowanie potoku obliczeniowego
- Zbiorcze przesyłanie zakodowanych poleceń do GPU
- Odczytywanie bufora GPU macierzy wyników
Tworzenie buforów GPU
Dla uproszczenia macierze będą przedstawiane jako lista liczb zmiennoprzecinkowych. Pierwszy element to liczba wierszy, drugi element to liczba kolumn, a reszta to rzeczywiste liczby w tablicy.
Te 3 bufory GPU pełnią funkcję buforów pamięci, ponieważ musimy przechowywać i pobierać dane w obrębie Compute Engine. To wyjaśnia, dlaczego flagi wykorzystania bufora GPU w przypadku wszystkich tych zdarzeń zawierają parametr GPUBufferUsage.STORAGE. Flaga wykorzystania macierzy wyników zawiera też właściwość GPUBufferUsage.COPY_SRC, ponieważ jest kopiowana do innego bufora w celu 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
});
Układ grupy powiązań i grupa powiązań
Koncepcje układu grupy powiązań i grupy powiązań są specyficzne dla interfejsu WebGPU. Układ grupy powiązań określa interfejs wejścia/wyjścia oczekiwany przez cieniowanie, a grupa powiązań przedstawia rzeczywiste dane wejściowe/wyjściowe dla cieniowania.
W przykładzie poniżej układ grupy powiązań wymaga 2 buforów pamięci tylko do odczytu w ponumerowanych powiązaniach wpisów 0 i 1 oraz bufora pamięci masowej 2 na potrzeby programu do cieniowania Compute.
Z kolei grupa powiązań, zdefiniowana dla tego układu grupy powiązań, wiąże bufory GPU z wpisami: gpuBufferFirstMatrix z wiązaniem 0, gpuBufferSecondMatrix z wiązaniem 1 i 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 do cieniowania w Compute
Kod narzędzia do cieniowania, który służy do mnożenia macierzy, jest napisany w języku WGSL (WebGPU Shader Language), który można łatwo przetłumaczyć na język SPIR-V. Poniżej znajdują się 3 bufory pamięci opisane przez var<storage>. Program będzie używać firstMatrix i secondMatrix jako danych wejściowych i resultMatrix jako danych wyjściowych.
Pamiętaj, że każdy bufor pamięci ma ustawioną ozdobę binding, która odpowiada temu samemu indeksowi zdefiniowanemu 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 opisujący operacje obliczeniowe, które będziemy wykonywać. Aby go utworzyć, zadzwoń pod numer device.createComputePipeline().
Wykorzystuje 2 argumenty: utworzony wcześniej układ grupy powiązań i etap obliczeniowy definiujący punkt wejścia naszego cienia danych (funkcja WGSL main) oraz rzeczywisty moduł cieniowania obliczeniowy 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 3 naszymi buforami GPU i potokiem obliczeniowym z układem grupy powiązań czas ich użyć.
Uruchommy programowalny koder Compute Pass z funkcją commandEncoder.beginComputePass(). Będziemy go używać do zakodowania poleceń GPU,
które będą wykonywać mnożenie macierzy. Ustaw potok z passEncoder.setPipeline(computePipeline) i jego grupą powiązań na indeks 0 z passEncoder.setBindGroup(0, bindGroup). Indeks 0 odpowiada opakowaniu elementu group(0) w kodzie WGSL.
Porozmawiajmy teraz o tym, jak ten program do cieniowania będzie działać w GPU. Naszym celem jest wykonanie tego programu równolegle dla każdej komórki macierzy wyników, krok po kroku. W przypadku macierzy o rozmiarach 16 x 32, aby zakodować polecenie wykonania, w tabeli @workgroup_size(8, 8) wywoływaćmy funkcję 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 ma domyślną wartość 1, ponieważ w tym miejscu go nie potrzebujemy.
W świecie zasobów obliczeniowych w GPU kodowanie polecenia uruchomienia funkcji jądra na zbiorze danych jest nazywane dyspozycją.
Rozmiar siatki grup roboczych dla naszego programu do cieniowania wynosi (8, 8) w naszym kodzie WGSL. Z tego względu wartości „x” i „y”, które są odpowiednio liczbą wierszy pierwszej macierzy, a liczbą kolumn drugiej macierzy, zostaną podzielone przez 8. Możemy teraz wysyłać wywołanie obliczeniowe przez passEncoder.dispatchWorkgroups(firstMatrix[0] / 8, secondMatrix[1] / 8). Liczba siatki grup roboczych do uruchomienia to argumenty dispatchWorkgroups().
Jak widać na ilustracji powyżej, każdy program do cieniowania będzie miał dostęp do unikalnego obiektu builtin(global_invocation_id), który będzie określać, 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 wyłączyć koder Compute Pass, wywołaj polecenie passEncoder.end(). Następnie utwórz bufor procesora graficznego, który będzie używany jako miejsce docelowe do skopiowania bufora macierzy wyników za pomocą copyBufferToBuffer. Na koniec zakończ kodowanie poleceń copyEncoder.finish() i prześlij je do kolejki na urządzeniu 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]);
Macierz wyników odczytu
Aby odczytać macierz wyników, wystarczy wywołać funkcję gpuReadBuffer.mapAsync() za pomocą metody GPUMapMode.READ i poczekać na rozwiązanie powracającej obietnicy, co oznacza, że bufor GPU jest teraz zmapowany. Zmapowany zakres możesz teraz pobrać za pomocą funkcji gpuReadBuffer.getMappedRange().
W naszym kodzie wynik zapisany w konsoli JavaScript Narzędzi 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 zagrać z samplem.
Ostatnia sztuczka
Jednym ze sposobów na ułatwienie odczytu kodu jest użycie poręcznej metody getBindGroupLayout potoku obliczeniowego do wywnioskowania układu grupy powiązań z modułu cieniowania. Ta sztuczka eliminuje potrzebę tworzenia niestandardowego układu grup powiązań i określania układu potoku w potoku obliczeniowym, 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 więc działa mnożenie macierzy na GPU w porównaniu do działania na procesorze? Aby to sprawdzić, napisałem program właśnie dla procesora. Jak widać na wykresie poniżej, wykorzystanie pełnej mocy GPU wydaje się oczywistym wyborem, gdy rozmiar matryc przekracza 256 x 256.
Ten artykuł to dopiero początek mojej przygody z poznawaniem WebGPU. Wkrótce pojawi się więcej artykułów z bardziej szczegółowymi informacjami o GPU Compute i o tym, jak działa renderowanie (płótno, tekstura, próbkowanie) w WebGPU.