In diesem Beitrag wird die experimentelle WebGPU API anhand von Beispielen erläutert und Sie erhalten einen Einstieg in die Durchführung von datenparrallel ausgeführten Berechnungen mit der GPU.
Hintergrund
Wie Sie vielleicht bereits wissen, ist die Graphic Processing Unit (GPU) ein elektronisches Subsystem in einem Computer, das ursprünglich auf die Verarbeitung von Grafiken spezialisiert war. In den letzten zehn Jahren hat sie sich jedoch zu einer flexibleren Architektur weiterentwickelt, mit der Entwickler viele Arten von Algorithmen implementieren können, nicht nur 3D-Grafiken rendern, und dabei die einzigartige Architektur der GPU nutzen. Diese Funktionen werden als GPU-Computing bezeichnet. Die Verwendung einer GPU als Coprozessor für allgemeine wissenschaftliche Berechnungen wird als GPGPU-Programmierung (General-Purpose GPU) bezeichnet.
GPU-Computing hat wesentlich zum aktuellen Boom des maschinellen Lernens beigetragen, da Convolution Neural Networks und andere Modelle die Architektur nutzen können, um effizienter auf GPUs ausgeführt zu werden. Da die aktuelle Webplattform keine GPU-Computing-Funktionen bietet, entwickelt die Community-Gruppe „GPU for the Web“ des W3C eine API, um die modernen GPU-APIs bereitzustellen, die auf den meisten aktuellen Geräten verfügbar sind. Diese API heißt WebGPU.
WebGPU ist ein Low-Level-API wie WebGL. Sie ist sehr wirkungsvoll und sehr ausführlich. Das ist auch vollkommen in Ordnung. Wir suchen nach Leistung.
In diesem Artikel konzentriere ich mich auf den GPU-Computing-Teil von WebGPU und gehe ehrlich gesagt nur an der Oberfläche entlang, damit Sie selbst loslegen können. In zukünftigen Artikeln werde ich das WebGPU-Rendering (Canvas, Textur usw.) genauer behandeln.
Auf die GPU zugreifen
Der Zugriff auf die GPU ist in WebGPU ganz einfach. Wenn du navigator.gpu.requestAdapter() aufrufst, wird ein JavaScript-Promise zurückgegeben, das asynchron mit einem GPU-Adapter aufgelöst wird. Stellen Sie sich diesen Adapter als Grafikkarte vor. Sie kann entweder integriert (auf demselben Chip wie die CPU) oder diskret (in der Regel eine PCIe-Karte, die leistungsstärker ist, aber mehr Strom verbraucht) sein.
Nachdem Sie den GPU-Adapter haben, rufen Sie adapter.requestDevice() auf, um ein Versprechen zu erhalten, das mit einem GPU-Gerät aufgelöst wird, das Sie für GPU-Berechnungen verwenden.
const adapter = await navigator.gpu.requestAdapter();
if (!adapter) { return; }
const device = await adapter.requestDevice();
Beide Funktionen haben Optionen, mit denen Sie die Art des Adapters (Energiespareinstellungen) und des Geräts (Erweiterungen, Limits) angeben können. Der Einfachheit halber verwenden wir in diesem Artikel die Standardoptionen.
Zwischenspeicher schreiben
Sehen wir uns an, wie Sie mit JavaScript Daten für die GPU in den Arbeitsspeicher schreiben. Dieser Vorgang ist aufgrund des Sandboxing-Modells, das in modernen Webbrowsern verwendet wird, nicht ganz einfach.
Im folgenden Beispiel wird gezeigt, wie vier Byte in den GPU-zugänglichen Pufferspeicher geschrieben werden. Es ruft device.createBuffer() auf, das die Größe des Puffers und seine Nutzung berücksichtigt. Auch wenn das Nutzungsflag GPUBufferUsage.MAP_WRITE für diesen bestimmten Aufruf nicht erforderlich ist, geben wir an, dass wir in diesen Puffer schreiben möchten. Da mappedAtCreation auf „wahr“ gesetzt ist, wird beim Erstellen ein GPU-Pufferobjekt zugeordnet. Anschließend kann der zugehörige Roh-Binärdatenpuffer durch Aufrufen der GPU-Puffermethode getMappedRange() abgerufen werden.
Wenn Sie schon einmal mit ArrayBuffer gearbeitet haben, ist das Schreiben von Bytes kein Problem. Verwenden Sie einfach einen TypedArray und kopieren Sie die Werte hinein.
// 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]);
An dieser Stelle ist der GPU-Zwischenspeicher zugeordnet, d. h. er gehört der CPU und ist mit Lese-/Schreibvorgängen aus JavaScript zugänglich. Damit die GPU darauf zugreifen kann, muss die Zuordnung aufgehoben werden. Das geht ganz einfach mit gpuBuffer.unmap().
Das Konzept „zugewiesen/nicht zugewiesen“ ist erforderlich, um Wettlaufsituationen zu vermeiden, bei denen GPU und CPU gleichzeitig auf den Arbeitsspeicher zugreifen.
Pufferspeicher lesen
Sehen wir uns nun an, wie wir einen GPU-Zwischenspeicher in einen anderen GPU-Zwischenspeicher kopieren und auslesen können.
Da wir in den ersten GPU-Puffer schreiben und ihn in einen zweiten GPU-Puffer kopieren möchten, ist ein neues Nutzungsflag GPUBufferUsage.COPY_SRC erforderlich. Der zweite GPU-Puffer wird diesmal in einem nicht zugeordneten Zustand mit device.createBuffer() erstellt. Das Nutzungs-Flag ist GPUBufferUsage.COPY_DST |
GPUBufferUsage.MAP_READ, da es als Ziel des ersten GPU-Zwischenspeichers verwendet und in JavaScript gelesen wird, nachdem GPU-Kopierbefehle ausgeführt wurden.
// 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
});
Da die GPU ein unabhängiger Coprozessor ist, werden alle GPU-Befehle asynchron ausgeführt. Aus diesem Grund wird eine Liste von GPU-Befehlen erstellt und bei Bedarf in Batches gesendet. In WebGPU ist der von device.createCommandEncoder() zurückgegebene GPU-Befehls-Encoder das JavaScript-Objekt, das einen Batch von „gepufferten“ Befehlen erstellt, die zu einem bestimmten Zeitpunkt an die GPU gesendet werden. Die Methoden in GPUBuffer hingegen sind „nicht gepuffert“, was bedeutet, dass sie zum Zeitpunkt des Aufrufs atomar ausgeführt werden.
Sobald Sie den GPU-Befehlsencoder haben, rufen Sie copyEncoder.copyBufferToBuffer() wie unten gezeigt auf, um diesen Befehl zur späteren Ausführung zur Befehlswarteschlange hinzuzufügen.
Schließen Sie die Codierungsbefehle ab, indem Sie copyEncoder.finish() aufrufen und sie an die GPU-Gerätewarteschlange senden. Die Warteschlange ist für die Verarbeitung von über device.queue.submit() eingereichten Einträgen mit den GPU-Befehlen als Argumenten verantwortlich.
Dadurch werden alle im Array gespeicherten Befehle nacheinander ausgeführt.
// 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]);
Zu diesem Zeitpunkt wurden GPU-Warteschlangenbefehle gesendet, aber nicht unbedingt ausgeführt.
Wenn Sie den zweiten GPU-Puffer lesen möchten, rufen Sie gpuReadBuffer.mapAsync() mit GPUMapMode.READ auf. Es gibt ein Versprechen zurück, das aufgelöst wird, wenn der GPU-Puffer zugeordnet ist. Rufen Sie dann den zugeordneten Bereich mit gpuReadBuffer.getMappedRange() ab, der dieselben Werte wie der erste GPU-Puffer enthält, sobald alle angeforderten GPU-Befehle ausgeführt wurden.
// Read buffer.
await gpuReadBuffer.mapAsync(GPUMapMode.READ);
const copyArrayBuffer = gpuReadBuffer.getMappedRange();
console.log(new Uint8Array(copyArrayBuffer));
Probieren Sie dieses Beispiel aus.
Kurz gesagt: Was Sie in Bezug auf Zwischenspeichervorgänge beachten müssen:
- GPU-Buffer müssen entkoppelt werden, damit sie in der Gerätewarteschlange verwendet werden können.
- Nach der Zuordnung können GPU-Buffer in JavaScript gelesen und geschrieben werden.
- GPU-Buffer werden zugeordnet, wenn
mapAsync()undcreateBuffer()mitmappedAtCreationauf „true“ gesetzt aufgerufen werden.
Shader-Programmierung
Programme, die auf der GPU ausgeführt werden und nur Berechnungen ausführen (und keine Dreiecke zeichnen), werden als Compute Shader bezeichnet. Sie werden parallel von Hunderten von GPU-Kernen (die kleiner als CPU-Kerne sind) ausgeführt, die gemeinsam Daten verarbeiten. Ihre Eingabe und Ausgabe sind Buffers in WebGPU.
Um die Verwendung von Rechen-Shadern in WebGPU zu veranschaulichen, werden wir mit der Matrixmultiplikation spielen, einem gängigen Algorithmus im maschinellen Lernen, wie unten dargestellt.
Kurz gesagt, gehen wir so vor:
- Erstellen Sie drei GPU-Zwischenspeicher (zwei für die zu multiplizierenden Matrizen und einen für die Ergebnismatrix).
- Eingabe und Ausgabe für den Compute-Shader beschreiben
- Compute Shader-Code kompilieren
- Compute-Pipeline einrichten
- Die codierten Befehle im Batch an die GPU senden
- GPU-Puffer der Ergebnismatrix lesen
Erstellen von GPU-Buffers
Der Einfachheit halber werden Matrizen als Liste von Gleitkommazahlen dargestellt. Das erste Element ist die Anzahl der Zeilen, das zweite Element die Anzahl der Spalten und der Rest sind die tatsächlichen Zahlen der Matrix.
Die drei GPU-Buffer sind Speicherbuffer, da wir Daten im Compute-Shader speichern und abrufen müssen. Das erklärt, warum die Flags für die GPU-Puffernutzung für alle GPUBufferUsage.STORAGE enthalten. Das Flag für die Verwendung der Ergebnismatrix enthält ebenfalls GPUBufferUsage.COPY_SRC, da es zum Lesen in einen anderen Puffer kopiert wird, sobald alle GPU-Warteschlangenbefehle ausgeführt wurden.
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 und bind group
Die Konzepte „Bindegruppenlayout“ und „Bindegruppe“ sind spezifisch für WebGPU. Ein Bindungsgruppen-Layout definiert die von einem Shader erwartete Eingabe-/Ausgabeschnittstelle, während eine Bindungsgruppe die tatsächlichen Eingabe-/Ausgabedaten für einen Shader darstellt.
Im folgenden Beispiel erwartet das Bindungsgruppen-Layout zwei schreibgeschützte Speicherbuffer an den nummerierten Eintragsbindungen 0 und 1 sowie einen Speicherbuffer an 2 für den Compute-Shader.
Die Bindungsgruppe hingegen, die für dieses Bindungsgruppenlayout definiert ist, ordnet den Einträgen GPU-Buffer zu: gpuBufferFirstMatrix der Bindung 0, gpuBufferSecondMatrix der Bindung 1 und resultMatrixBuffer der Bindung 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
}
}
]
});
Shader-Code berechnen
Der Compute-Shader-Code für die Matrixmultiplikation ist in WGSL (WebGPU Shader Language) geschrieben, einer Sprache, die sich problemlos in SPIR-V übersetzen lässt. Ohne ins Detail zu gehen, finden Sie unten die drei Speicherbuffer, die mit var<storage> gekennzeichnet sind. Das Programm verwendet firstMatrix und secondMatrix als Eingaben und resultMatrix als Ausgabe.
Jeder Speicherpuffer hat eine binding-Dekoration, die dem Index entspricht, der in den Bindungsgruppenlayouts und Bindungsgruppen oben definiert wurde.
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;
}
`
});
Pipelineeinrichtung
Die Compute-Pipeline ist das Objekt, das den zu führenden Berechnungsvorgang beschreibt. Rufen Sie device.createComputePipeline() auf, um sie zu erstellen.
Es nimmt zwei Argumente entgegen: das Bindungsgruppen-Layout, das wir zuvor erstellt haben, und eine Compute-Phase, die den Einstiegspunkt unseres Compute-Shaders (die main-WGSL-Funktion) und das eigentliche Compute-Shader-Modul definiert, das mit device.createShaderModule() erstellt wurde.
const computePipeline = device.createComputePipeline({
layout: device.createPipelineLayout({
bindGroupLayouts: [bindGroupLayout]
}),
compute: {
module: shaderModule,
entryPoint: "main"
}
});
Befehle senden
Nachdem wir eine Bindungsgruppe mit unseren drei GPU-Buffern und einer Compute-Pipeline mit einem Bindungsgruppenlayout instanziiert haben, ist es an der Zeit, sie zu verwenden.
Starten wir einen programmierbaren Encoder für Fahrkarten/Tickets mit commandEncoder.beginComputePass(). Damit codieren wir GPU-Befehle, die die Matrixmultiplikation ausführen. Legen Sie die Pipeline mit passEncoder.setPipeline(computePipeline) und die Bindungsgruppe mit dem Index 0 mit passEncoder.setBindGroup(0, bindGroup) fest. Der Index 0 entspricht der Verzierung group(0) im WGSL-Code.
Sehen wir uns nun an, wie dieser Compute-Shader auf der GPU ausgeführt wird. Unser Ziel ist es, dieses Programm Schritt für Schritt parallel für jede Zelle der Ergebnismatrix auszuführen. Für eine Ergebnismatrix mit einer Größe von 16 × 32 würden wir beispielsweise passEncoder.dispatchWorkgroups(2, 4) oder passEncoder.dispatchWorkgroups(16 / 8, 32 / 8) aufrufen, um den Ausführungsbefehl auf einem @workgroup_size(8, 8) zu codieren.
Das erste Argument „x“ ist die erste Dimension, das zweite „y“ die zweite Dimension und das letzte „z“ die dritte Dimension, die standardmäßig auf 1 festgelegt ist, da wir sie hier nicht benötigen.
In der GPU-Rechenwelt wird das Codieren eines Befehls zum Ausführen einer Kernelfunktion für einen Datensatz als Dispatcher bezeichnet.
Die Größe des Arbeitsgruppen-Rasters für unseren Compute-Shader ist (8, 8) in unserem WGSL-Code. Aus diesem Grund werden "x" und "y", die jeweils die Anzahl der Zeilen der ersten Matrix und die Anzahl der Spalten der zweiten Matrix sind, durch 8 geteilt. Damit können wir jetzt einen Compute-Aufruf mit passEncoder.dispatchWorkgroups(firstMatrix[0] / 8, secondMatrix[1] / 8) auslösen. Die Anzahl der Arbeitsgruppenraster, die ausgeführt werden sollen, entspricht den Argumenten dispatchWorkgroups().
Wie in der Abbildung oben zu sehen, hat jeder Shader Zugriff auf ein eindeutiges builtin(global_invocation_id)-Objekt, anhand dessen ermittelt wird, welche Ergebnismatrixzelle berechnet werden soll.
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();
Rufe passEncoder.end() auf, um den Encoder für die Berechnung von Karten/Tickets zu beenden. Erstellen Sie dann einen GPU-Puffer, der als Ziel verwendet wird, um den Ergebnismatrixpuffer mit copyBufferToBuffer zu kopieren. Schließen Sie die Codierungsbefehle mit copyEncoder.finish() ab und reichen Sie sie über device.queue.submit() mit den GPU-Befehlen in der GPU-Gerätewarteschlange ein.
// 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]);
Ergebnismatrix lesen
Das Lesen der Ergebnismatrix ist ganz einfach: Rufe gpuReadBuffer.mapAsync() mit GPUMapMode.READ auf und warte, bis das zurückgegebene Promise gelöst wird, was bedeutet, dass der GPU-Puffer jetzt zugeordnet ist. An dieser Stelle ist es möglich, den zugeordneten Bereich mit gpuReadBuffer.getMappedRange() abzurufen.
In unserem Code ist das in der JavaScript-Konsole der Entwicklertools protokollierte Ergebnis „2, 2, 50, 60, 114, 140“.
// Read buffer.
await gpuReadBuffer.mapAsync(GPUMapMode.READ);
const arrayBuffer = gpuReadBuffer.getMappedRange();
console.log(new Float32Array(arrayBuffer));
Glückwunsch! Sie haben es geschafft Sie können mit dem Beispiel experimentieren.
Ein letzter Trick
Eine Möglichkeit, Ihren Code leichter lesbar zu machen, besteht darin, die praktische Methode getBindGroupLayout der Compute-Pipeline zu verwenden, um das Layout der Bindungsgruppe aus dem Shader-Modul abzuleiten. Dadurch müssen Sie kein benutzerdefiniertes Bindungsgruppenlayout erstellen und kein Pipeline-Layout in Ihrer Compute-Pipeline angeben, wie unten dargestellt.
Eine Abbildung von getBindGroupLayout für das vorherige Beispiel finden Sie hier.
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: [
Leistungsbewertung
Wie unterscheidet sich die Ausführung der Matrixmultiplikation auf einer GPU von der Ausführung auf einer CPU? Um das herauszufinden, habe ich das eben beschriebene Programm für eine CPU geschrieben. Wie Sie in der folgenden Grafik sehen, ist die volle Leistung der GPU eine naheliegende Wahl, wenn die Größe der Matrizen größer als 256 × 256 ist.
Dieser Artikel war nur der Anfang meiner Erkundung der WebGPU. Demnächst folgen weitere Artikel mit detaillierten Informationen zu GPU-Computing und zur Funktionsweise von Rendering (Canvas, Textur, Sampler) in WebGPU.