W tym artykule za pomocą przykładów omawiamy eksperymentalny interfejs WebGPU API i pomagamy Ci rozpocząć wykonywanie obliczeń równoległych z wykorzystaniem GPU.
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 ewoluowała w kierunku bardziej elastycznej architektury, która pozwala deweloperom stosować wiele typów algorytmów, a nie tylko do renderowania grafiki 3D, wykorzystując przy tym unikalną architekturę GPU. Te możliwości są określane jako obliczanie na GPU, a korzystanie z GPU jako procesora pomocniczego do ogólnego przetwarzania naukowego nazywa się programowaniem GPU ogólnego przeznaczenia (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 interfejs API niskiego poziomu, podobnie jak WebGL. Jest bardzo wydajny i bardzo szczegółowy, jak się przekonasz. Ale 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ń na 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). Ze względu na prostotę w tym artykule użyjemy opcji domyślnych.
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 stosowany w nowoczesnych przeglądarkach internetowych.
Przykład poniżej pokazuje, jak zapisać 4 bajty do pamięci buforowej dostępnej z 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]);
W tym momencie bufor GPU jest mapowany, co oznacza, że jest on własnością procesora i dostępny do odczytu/zapisu z JavaScript. 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 buforowej
Zobaczmy teraz, jak skopiować bufor GPU do innego bufora GPU i go odczytać.
Ponieważ zapisujemy dane w pierwszym buforze GPU i chcemy je skopiować do drugiego bufora GPU, potrzebna jest nowa flaga użycia 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() z wartością 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 najważniejsze informacje o operacjach na buforze:
- 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()icreateBuffer()z wartością true w polumappedAtCreation.
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 (które są mniejsze niż rdzenie procesora) i działają razem, aby przetwarzać dane. Ich dane wejściowe i wyjściowe są buforami w WebGPU.
Aby zilustrować użycie shaderów obliczeniowych w WebGPU, pobawimy się mnożeniem macierzy, czyli częstym algorytmem uczenia maszynowego, który jest pokazany poniżej.
Krótko mówiąc, zrobimy to w ten sposób:
- Utwórz 3 bufory GPU (2 dla macierzy do mnożenia i 1 dla macierzy wynikowej).
- Opisz dane wejściowe i wyjściowe dla shadera obliczeniowego
- Kompilowanie kodu shadera obliczeniowego
- Konfigurowanie potoku obliczeniowego
- Przesyłanie zakodowanych poleceń do GPU wsadowo
- 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. Z tego powodu flagi wykorzystania bufora GPU zawierają flagę GPUBufferUsage.STORAGE dla wszystkich. 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
});
Bind group layout i bind group
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.
Zdefiniowana w tym układzie grupa wiązania skojarzy bufory GPU z tymi wpisami: gpuBufferFirstMatrix z wiązaniem 0, gpuBufferSecondMatrix z wiązaniem 1 oraz resultMatrixBuffer z wią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 (funkcję 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 uruchomieniu grupy wiązania z 3 buforami GPU i przetwarzania z użyciem architektury procesora graficznego z grupą wiązania nadszedł 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 przepustowość za pomocą parametru passEncoder.setPipeline(computePipeline) i grupę wiązania o indeksie 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. Na przykład w przypadku macierzy wyników o wymiarach 16 x 32, aby zakodować polecenie wykonania na @workgroup_size(8, 8), użyjemy funkcji 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 dotyczące przetwarzania danych z adresu 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 kopiowania za pomocą copyBufferToBuffer bufora macierzy wyników. 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.
Dostępny jest obrazek getBindGroupLayout z poprzedniego przykładu.
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 omówiony powyżej na potrzeby procesora 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ł był dopiero początkiem mojej podróży w badaniu 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.