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 wissen, ist die GPU (Graphics Processing Unit) ein elektronisches Subsystem in einem Computer, das ursprünglich für die Verarbeitung von Grafiken entwickelt wurde. In den letzten zehn Jahren hat sich die GPU jedoch zu einer flexibleren Architektur entwickelt, mit der Entwickler nicht nur 3D-Grafiken rendern, sondern auch viele Arten von Algorithmen implementieren können, während sie gleichzeitig 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 jüngsten Boom des maschinellen Lernens beigetragen, da Convolutional 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 eine Low-Level-API wie WebGL. Sie ist sehr leistungsstark und recht umfangreich, wie Sie sehen werden. 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.
Schreibpufferspeicher
Sehen wir uns an, wie Sie mit JavaScript Daten in den Arbeitsspeicher der GPU 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, wollen wir explizit darauf hinweisen, 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 diesem Punkt ist der GPU-Puffer zugeordnet, d. h., er gehört der CPU und kann von JavaScript aus gelesen und geschrieben werden. 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 Sie einen GPU-Puffer in einen anderen GPU-Puffer kopieren und wieder lesen.
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 mit device.createBuffer() in einem nicht zugeordneten Zustand erstellt. Das Flag für die Verwendung ist GPUBufferUsage.COPY_DST |
GPUBufferUsage.MAP_READ, da es als Ziel des ersten GPU-Buffers verwendet und in JavaScript gelesen wird, sobald GPU-Kopienbefehle 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 auf GPUBuffer sind dagegen „nicht gepuffert“, d. h., sie werden zum Zeitpunkt des Aufrufs atomar ausgeführt.
Rufen Sie nach dem Erstellen des GPU-Befehlscode-Encoders copyEncoder.copyBufferToBuffer() wie unten gezeigt auf, um diesen Befehl der Befehlswarteschlange zur späteren Ausführung 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.
Im Folgenden finden Sie eine kurze Zusammenfassung der wichtigsten Punkte zu Bufferspeicher-Vorgängen:
- GPU-Buffer müssen entkoppelt werden, um in der Gerätewarteschlange verwendet zu werden.
- Nach der Zuordnung können GPU-Buffer in JavaScript gelesen und geschrieben werden.
- GPU-Buffer werden zugeordnet, wenn
mapAsync()undcreateBuffer()mitmappedAtCreationauf „true“ 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 Compute Shadern in WebGPU zu veranschaulichen, spielen wir mit der Matrixmultiplikation, einem gängigen Algorithmus im maschinellen Lernen, der unten dargestellt ist.
Kurz gesagt, gehen wir so vor:
- Drei GPU-Buffer erstellen (zwei für die zu multiplizierenden Matrizen und einer 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
});
Bindegruppenlayout und Bindegruppe
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
}
}
]
});
Compute-Shader-Code
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;
}
`
});
Pipeline einrichten
Die Compute-Pipeline ist das Objekt, das den zu führenden Berechnungsvorgang beschreibt. Erstellen Sie sie durch Aufrufen von device.createComputePipeline().
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.
Beginnen wir mit einem programmierbaren Encoder für Fahrkarten/Tickets für öffentliche Verkehrsmittel 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 passEncoder.setBindGroup(0, bindGroup) am Index 0 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“ gesetzt ist, da wir sie hier nicht benötigen.
In der GPU-Rechenwelt wird das Codieren eines Befehls zum Ausführen einer Kernelfunktion auf einem Datensatz als Dispatcher bezeichnet.
Die Größe des Arbeitsgruppen-Rasters für unseren Compute-Shader ist (8, 8) in unserem WGSL-Code. Daher werden „x“ und „y“, die Anzahl der Zeilen der ersten Matrix bzw. der Spalten der zweiten Matrix, durch 8 geteilt. Damit können wir jetzt einen Compute-Aufruf mit passEncoder.dispatchWorkgroups(firstMatrix[0] / 8, secondMatrix[1] / 8) senden. Die Anzahl der zu erstellenden Arbeitsgruppenraster wird durch die dispatchWorkgroups()-Argumente festgelegt.
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();
Wenn Sie den Encoder für den persönlichen Ausweis beenden möchten, drücken Sie die Taste passEncoder.end(). 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 können Sie den zugeordneten Bereich mit gpuReadBuffer.getMappedRange() abrufen.
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 Du kannst mit dem Sample herumspielen.
Ein letzter Trick
Mit der praktischen Methode getBindGroupLayout der Compute-Pipeline können Sie das Bindegruppen-Layout aus dem Shadermodul ableiten, um Ihren Code leichter lesbar zu machen. 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 schneidet die Ausführung der Matrixmultiplikation auf einer GPU im Vergleich zur Ausführung auf einer CPU ab? 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.