In diesem Beitrag wird die experimentelle WebGPU API anhand von Beispielen erläutert. beginnen Sie mit datenparallelen Berechnungen mit der GPU.
Hintergrund
Wie Sie vielleicht schon wissen, ist der Grafikprozessor (Graphic Processing Unit, GPU) ein elektronischer Subsystem eines Computers, der ursprünglich auf die Verarbeitung spezialisiert war, Grafiken. In den letzten 10 Jahren hat sie sich jedoch zu einer flexibleren Architektur, mit der Entwickler viele Arten von Algorithmen implementieren können, 3D-Grafiken rendern und dabei die einzigartige Architektur der GPU Diese Funktionen werden als GPU-Computing bezeichnet. Dabei wird eine GPU Der Koprozessor für das wissenschaftliche Rechnen für allgemeine Zwecke wird als GPU-Programmierung (GPGPU)
GPU-Computing hat wesentlich zum aktuellen Boom des maschinellen Lernens da Convolution Neural Networks und andere Modelle für eine effizientere Ausführung auf GPUs. Mit der aktuellen Webplattform die mangelnde GPU-Computing-Funktionalität, die "GPU für das Web" des W3C Community-Gruppe entwickelt eine API, um die modernen GPU-APIs bereitzustellen, die auf den meisten aktuellen Geräten. Diese API wird als WebGPU bezeichnet.
WebGPU ist ein Low-Level-API wie WebGL. Sie ist sehr wirkungsvoll und sehr ausführlich, die Sie sehen werden. Das ist auch vollkommen in Ordnung. Wir achten auf die Leistung.
In diesem Artikel werde ich mich auf den GPU-Computing-Teil von WebGPU konzentrieren und Ehrlich gesagt, ich kratzer nur an der Oberfläche, damit ihr gleich auf eurem gehören. Ich werde tiefer in das WebGPU-Rendering (Canvas, Textur, künftige Artikel.
Auf GPU zugreifen
Der Zugriff auf die GPU ist über WebGPU einfach. navigator.gpu.requestAdapter() wird angerufen
gibt ein JavaScript-Promise zurück, das asynchron mit einer GPU aufgelöst wird
Adapter. Stellen Sie sich diesen Adapter als Grafikkarte vor. Sie können sie entweder
(auf demselben Chip wie die CPU) oder diskret (normalerweise eine PCIe-Karte,
ist aber leistungsfähiger, verbraucht jedoch mehr Energie).
Sobald Sie den GPU-Adapter haben, rufen Sie adapter.requestDevice() auf, um ein Promise 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 bieten Optionen, mit denen Sie genau angeben können, Netzteil (Stromeinstellung) und das gewünschte Gerät (Erweiterungen, Limits) ein. Für die 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. Dieses Der Prozess ist aufgrund des Sandbox-Modells, das im modernen Web verwendet wird, nicht ganz einfach. Browser.
Das folgende Beispiel zeigt, wie Sie vier Byte in den Pufferspeicher schreiben, auf den zugegriffen werden kann.
von der GPU. Sie ruft device.createBuffer() auf, wobei die Größe des
und dessen Nutzung. Auch wenn das Nutzungs-Flag GPUBufferUsage.MAP_WRITE
für diesen speziellen Aufruf nicht erforderlich ist,
sollten wir ausdrücklich darauf hinweisen,
in diesen Puffer ein. Dies führt zu einem GPU-Zwischenspeicherobjekt, das bei der Erstellung
mappedAtCreation auf „true“ gesetzt. Dann kann der zugehörige Zwischenspeicher der binären Rohdaten
durch Aufrufen der GPU-Puffermethode getMappedRange() abgerufen werden.
Das Schreiben von Byte ist vertraut, wenn Sie bereits mit ArrayBuffer gespielt haben. eine
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.
kann über
Lesen/Schreiben aus JavaScript aufgerufen werden. Damit die GPU darauf zugreifen kann,
muss die Zuordnung aufgehoben werden. Dazu müssen Sie lediglich gpuBuffer.unmap() aufrufen.
Das Konzept der Zuordnung/Nichtzuordnung ist erforderlich, um Race-Bedingungen zu vermeiden, und CPU-Zugriff auf den Arbeitsspeicher.
Zwischenspeicher 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-Zwischenspeicher schreiben,
GPU-Zwischenspeicher, neues Nutzungs-Flag GPUBufferUsage.COPY_SRC ist erforderlich. Die zweite
Der GPU-Zwischenspeicher wird diesmal in einem nicht zugeordneten Status mit
device.createBuffer() Das Nutzungs-Flag ist GPUBufferUsage.COPY_DST |
GPUBufferUsage.MAP_READ, da es als Ziel der ersten GPU verwendet wird
nach Ausführung der GPU-Kopierbefehle in JavaScript zu puffern und zu lesen.
// 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 Koprozessor ist, werden alle GPU-Befehle ausgeführt
asynchron programmiert. Aus diesem Grund wird eine Liste
mit GPU-Befehlen erstellt und
bei Bedarf Batches ausführen. In WebGPU hat der
GPU-Befehlsencoder, der von
device.createCommandEncoder() ist das JavaScript-Objekt, das einen Batch von
"gepuffert" die irgendwann an die GPU gesendet werden. Die Methoden auf
GPUBuffer hingegen sind „nicht gepuffert“, d. h. sie werden atomar ausgeführt
zum Zeitpunkt ihres Anrufs.
Wenn du den GPU-Befehlsencoder hast, rufe copyEncoder.copyBufferToBuffer() auf.
wie unten gezeigt, um diesen Befehl zur späteren Ausführung zur Befehlswarteschlange hinzuzufügen.
Schließe zum Abschluss die Codierungsbefehle ab, indem du copyEncoder.finish() aufrufst und
an die Befehlswarteschlange des GPU-Geräts. Die Warteschlange ist für die Verarbeitung
Einreichungen über device.queue.submit() mit den GPU-Befehlen als Argumente.
Dadurch werden alle im Array gespeicherten Befehle in kleinstmöglicher Reihenfolge 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]);
Bisher wurden GPU-Warteschlangenbefehle gesendet, aber nicht unbedingt ausgeführt.
Rufen Sie zum Lesen des zweiten GPU-Zwischenspeichers gpuReadBuffer.mapAsync() mit
GPUMapMode.READ. Sie gibt ein Promise zurück, das aufgelöst wird, wenn der GPU-Zwischenspeicher
zugeordnet. Dann rufen Sie den zugeordneten Bereich mit gpuReadBuffer.getMappedRange() ab,
enthält dieselben Werte wie der erste GPU-Zwischenspeicher, sobald alle GPU-Befehle in der Warteschlange
ausgeführt wurden.
// Read buffer.
await gpuReadBuffer.mapAsync(GPUMapMode.READ);
const copyArrayBuffer = gpuReadBuffer.getMappedRange();
console.log(new Uint8Array(copyArrayBuffer));
Sie können dieses Beispiel ausprobieren.
Kurz gesagt: Was Sie in Bezug auf Zwischenspeichervorgänge beachten müssen:
- Die Zuordnung von GPU-Zwischenspeichern muss aufgehoben werden, damit sie bei der Einreichung von Gerätewarteschlangen verwendet werden können.
- Bei der Zuordnung können GPU-Zwischenspeicher in JavaScript gelesen und geschrieben werden.
- GPU-Zwischenspeicher werden zugeordnet, wenn
mapAsync()undcreateBuffer()mitmappedAtCreation, die auf „true“ gesetzt sind, werden aufgerufen.
Shader-Programmierung
Programme, die auf der GPU ausgeführt werden und nur Berechnungen (und keine Dreiecken) werden als Compute-Shader bezeichnet. Sie werden von Hunderten Nutzern parallel ausgeführt. an GPU-Kernen (die kleiner als CPU-Kerne sind), die zusammen zur Verarbeitung Daten. Ihre Ein- und Ausgabe sind Puffer in einer WebGPU.
Um die Verwendung von Compute-Shadern in WebGPU zu veranschaulichen, werden wir mit der Matrixstruktur Multiplikation, ein gängiger Algorithmus beim maschinellen Lernen, der unten dargestellt ist.
<ph type="x-smartling-placeholder">
Hier ist unsere Vorgehensweise:
- Erstellen Sie drei GPU-Zwischenspeicher (zwei für die Matrizen zum Multiplizieren und einen für die Ergebnismatrix)
- Eingabe und Ausgabe für den Compute-Shader beschreiben
- Compute Shader-Code kompilieren
- Computing-Pipeline einrichten
- Codierte Befehle im Batch an die GPU senden
- Ergebnismatrix-GPU-Zwischenspeicher lesen
GPU-Puffererstellung
Matrizen werden der Einfachheit halber als Liste Punktzahlen. Das erste Element ist die Anzahl der Zeilen, das zweite Element Spaltenanzahl und der Rest sind die tatsächlichen Zahlen der Matrix.
<ph type="x-smartling-placeholder">
Die drei GPU-Zwischenspeicher sind Speicherpuffer, da wir Daten in
den Compute-Shader. Dies erklärt, warum die Flags für die Nutzung von GPU-Puffern Folgendes umfassen:
GPUBufferUsage.STORAGE für alle. Das Flag für die Ergebnismatrix-Verwendung enthält ebenfalls
GPUBufferUsage.COPY_SRC, da es in einen anderen Zwischenspeicher kopiert wird für
nachdem 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
});
Layout der Gruppe binden und Gruppe binden
Die Konzepte für das Layout und die Bindungsgruppe der Bindung gelten nur für WebGPU. Bindung Das Gruppenlayout definiert die Eingabe-/Ausgabeschnittstelle, die von einem Shader erwartet wird, während ein Bindungsgruppe stellt die tatsächlichen Eingabe-/Ausgabedaten für einen Shader dar.
Im folgenden Beispiel erwartet das Layout der Bindungsgruppe zwei schreibgeschützte Speicherzwischenspeicher,
nummerierte Eintragbindungen 0, 1 und ein Speicherpuffer unter 2 für den Compute-Shader.
Die Bindungsgruppe hingegen, die für dieses Layout der Bindungsgruppe definiert ist, verbindet
GPU puffert die Einträge: gpuBufferFirstMatrix für die Bindung 0
gpuBufferSecondMatrix zur Bindung 1 und resultMatrixBuffer zur Bindung
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 zum Multiplizieren von Matrizen ist in WGSL geschrieben:
WebGPU Shader Language, das sich einfach in SPIR-V übersetzen lässt. Ohne
Unter den drei angegebenen Speicherpuffern finden Sie
mit var<storage>. Das Programm verwendet firstMatrix und secondMatrix als
Eingaben und resultMatrix als Ausgabe.
Beachten Sie, dass für jeden Speicherzwischenspeicher ein binding-Element verwendet wird, das
denselben Index, der in den Bindungsgruppenlayouts und den oben deklarierten Bindungsgruppen definiert ist.
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 Computing-Pipeline ist das Objekt, das den Computing-Vorgang beschreibt.
die wir ausführen werden. Rufen Sie zum Erstellen device.createComputePipeline() auf.
Dafür sind zwei Argumente erforderlich: das zuvor erstellte Layout der Bindungsgruppe und ein
Phase, in der der Einstiegspunkt unseres Compute-Shaders definiert wird (die WGSL-Funktion main)
und dem tatsächlichen Compute-Shader-Modul, das mit device.createShaderModule() erstellt wurde.
const computePipeline = device.createComputePipeline({
layout: device.createPipelineLayout({
bindGroupLayouts: [bindGroupLayout]
}),
compute: {
module: shaderModule,
entryPoint: "main"
}
});
Einreichung von Befehlen
Nach der Instanziierung einer Bindungsgruppe mit unseren drei GPU-Zwischenspeichern und einem mit einem Bindungsgruppenlayout verwenden, ist es an der Zeit, sie zu verwenden.
Starten wir einen programmierbaren
Compute Pass-Encoder mit
commandEncoder.beginComputePass() Wir codieren damit GPU-Befehle
die die Matrixmultiplikation durchführt. Pipeline festlegen mit
passEncoder.setPipeline(computePipeline) und seine Bindungsgruppe bei Index 0 mit
passEncoder.setBindGroup(0, bindGroup). Der Index 0 entspricht dem
group(0) im WGSL-Code.
Sprechen wir nun darüber, wie dieser Compute-Shader auf der GPU ausgeführt wird. Unsere
dieses Programm für jede Zelle der Ergebnismatrix
gleichzeitig auszuführen,
Schritt für Schritt. Für eine Ergebnismatrix der Größe 16 x 32, um beispielsweise
Im Ausführungsbefehl würden wir auf einem @workgroup_size(8, 8)
passEncoder.dispatchWorkgroups(2, 4) oder passEncoder.dispatchWorkgroups(16 / 8, 32 / 8).
Das erste Argument „x“ ist die erste Dimension, die zweite „y“. ist die zweite Dimension,
und das letzte „z“ ist die dritte Dimension, die standardmäßig den Wert 1 hat, da sie hier nicht benötigt wird.
In der GPU-Computing-Welt wird das Codieren eines Befehls zum Ausführen einer Kernel-Funktion für einen Datensatz als Weiterleitungsfunktion bezeichnet.
Die Größe des Arbeitsgruppenrasters für unseren Compute-Shader beträgt in unserer WGSL (8, 8)
Code. Aus diesem Grund wird „x“ und „y“ die jeweils der Anzahl der Zeilen
werden die erste Matrix und die Anzahl der Spalten der zweiten Matrix
bis zum 8. Damit können wir jetzt einen Compute-Aufruf mit
passEncoder.dispatchWorkgroups(firstMatrix[0] / 8, secondMatrix[1] / 8) Die
Anzahl der Arbeitsgruppenraster, die ausgeführt werden sollen, sind die Argumente dispatchWorkgroups().
Wie in der obigen Zeichnung zu sehen, hat jeder Shader Zugriff auf einen eindeutigen
builtin(global_invocation_id)-Objekt, das als Ergebnis verwendet wird
Matrixzelle zu berechnen.
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
GPU-Zwischenspeicher, der als Ziel zum Kopieren des Ergebnismatrix-Zwischenspeichers verwendet werden soll
copyBufferToBuffer Beenden Sie die Codierungsbefehle mit
copyEncoder.finish() und senden Sie diese an die GPU-Gerätewarteschlange, indem Sie Folgendes aufrufen:
device.queue.submit() durch die GPU-Befehle.
// 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 so einfach wie das Aufrufen von gpuReadBuffer.mapAsync() mit
GPUMapMode.READ und warten, bis das zurückgegebene Versprechen aufgelöst wird. Dies weist darauf hin,
ist der GPU-Zwischenspeicher zugeordnet. An dieser Stelle ist es möglich, die Kartendaten
Bereich mit gpuReadBuffer.getMappedRange().
In unserem Code wird in der JavaScript-Konsole der Entwicklertools das 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, ist die Verwendung der praktischen
Methode getBindGroupLayout der Compute-Pipeline zum Ableiten der Bindungsgruppe
aus dem Shader-Modul. Mit diesem Trick ist es nicht mehr erforderlich,
Layout für benutzerdefiniertes Bindungsgruppen und Angeben eines Pipelinelayouts in Ihrem Computing
wie Sie unten sehen können.
Eine Abbildung von getBindGroupLayout für das vorherige Beispiel ist verfügbar.
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: [
Leistungsergebnisse
Wie verhält sich die Matrixmultiplikation auf einer GPU im Vergleich zur Ausführung auf einer GPU? CPU? Um dies herauszufinden, habe ich das gerade beschriebene Programm für eine CPU geschrieben. Und wie Sie Wie Sie im Diagramm unten sehen, erscheint es naheliegend, die volle GPU-Leistung zu nutzen. wenn die Größe der Matrizen größer als 256 mal 256 ist.
<ph type="x-smartling-placeholder">
Dieser Artikel war erst der Anfang meiner Reise zum Erlernen von WebGPU. Mehr erwartet Artikel zu GPU-Computing und zum Rendern (Canvas, Textur, Sampler) funktioniert mit WebGPU.