Questo post esplora l'API WebGPU sperimentale tramite esempi e ti aiuta a iniziare a eseguire calcoli in parallelo con i dati utilizzando la GPU.
Sfondo
Come forse già saprai, la GPU (Graphic Processing Unit) è un sottosistema elettronico all'interno di un computer, originariamente specializzato per l'elaborazione delle immagini. Tuttavia, negli ultimi 10 anni si è evoluto verso un'architettura più flessibile che consente agli sviluppatori di implementare molti tipi di algoritmi, non solo per il rendering di grafica 3D, sfruttando al contempo l'architettura unica della GPU. Queste funzionalità sono chiamate computing GPU e l'utilizzo di una GPU come coprocessore per il calcolo scientifico per uso generico è chiamato programmazione GPU per uso generico (GPGPU).
Il calcolo GPU ha contribuito in modo significativo al recente boom del machine learning, poiché le reti neurali convoluzionali e altri modelli possono sfruttare l'architettura per funzionare in modo più efficiente sulle GPU. Poiché l'attuale piattaforma web non dispone di funzionalità di calcolo GPU, il gruppo della community "GPU per il web" del W3C sta progettando un'API per esporre le API GPU moderne disponibili sulla maggior parte dei dispositivi attuali. Questa API si chiama WebGPU.
WebGPU è un'API a basso livello, come WebGL. È molto potente e piuttosto dettagliato, come vedrai. Ma non importa. Quello che cerchiamo è il rendimento.
In questo articolo mi concentrerò sulla parte di calcolo GPU di WebGPU e, a essere sincero, sto solo grattando la superficie, in modo che tu possa iniziare a giocare autonomamente. In articoli futuri approfondirò il rendering WebGPU (canvas, texture e così via).
Accedere alla GPU
Accedere alla GPU è facile in WebGPU. La chiamata a navigator.gpu.requestAdapter()
restituisce una promessa JavaScript che verrà risolta in modo asincrono con un adattatore GPU. Puoi considerare questo adattatore come la scheda grafica. Può essere integrata (sullo stesso chip della CPU) o discreta (di solito una scheda PCIe più performante, ma che consuma più energia).
Una volta ottenuto l'adattatore GPU, chiama adapter.requestDevice()
per ottenere una promessa che verrà risolta con un dispositivo GPU che utilizzerai per eseguire alcuni calcoli GPU.
const adapter = await navigator.gpu.requestAdapter();
if (!adapter) { return; }
const device = await adapter.requestDevice();
Entrambe le funzioni accettano opzioni che ti consentono di specificare il tipo di adattatore (preferenza di alimentazione) e di dispositivo (estensioni, limiti) che preferisci. Per semplicità, in questo articolo utilizzeremo le opzioni predefinite.
Memoria buffer di scrittura
Vediamo come utilizzare JavaScript per scrivere dati nella memoria della GPU. Questa procedura non è semplice a causa del modello di sandboxing utilizzato nei browser web moderni.
L'esempio seguente mostra come scrivere quattro byte nella memoria buffer accessibile dalla GPU. Chiama device.createBuffer()
, che prende le dimensioni
del buffer e il suo utilizzo. Anche se il flag di utilizzo GPUBufferUsage.MAP_WRITE
non è obbligatorio per questa chiamata specifica, dichiariamo esplicitamente che vogliamo scrivere in questo buffer. Il risultato è un oggetto buffer della GPU mappato al momento della creazione grazie amappedAtCreation
impostato su true. Il buffer di dati binari non elaborati associato può essere recuperato chiamando il metodo del buffer della GPU getMappedRange()
.
È possibile scrivere byte se hai già giocato con ArrayBuffer
. Utilizza un TypedArray
e copia i valori al suo interno.
// 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]);
A questo punto, il buffer GPU è mappato, ovvero è di proprietà della CPU ed è accessibile in lettura/scrittura da JavaScript. Affinché la GPU possa accedervi, deve
essere rimossa la mappatura, semplice come chiamare gpuBuffer.unmap()
.
Il concetto di mappato/non mappato è necessario per evitare condizioni di gara in cui GPU e CPU accedono contemporaneamente alla memoria.
Memoria buffer di lettura
Ora vediamo come copiare un buffer della GPU in un altro buffer della GPU e leggerlo di nuovo.
Poiché stiamo scrivendo nel primo buffer GPU e vogliamo copiarlo in un secondo
buffer GPU, è richiesto un nuovo flag di utilizzo GPUBufferUsage.COPY_SRC
. Il secondo buffer GPU viene creato questa volta in uno stato non mappato con device.createBuffer()
. Il relativo flag di utilizzo è GPUBufferUsage.COPY_DST |
GPUBufferUsage.MAP_READ
, in quanto verrà utilizzato come destinazione del primo buffer GPU e letto in JavaScript una volta eseguiti i comandi di copia della 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
});
Poiché la GPU è un coprocessore indipendente, tutti i comandi GPU vengono eseguiti
in modo asincrono. Ecco perché esiste un elenco di comandi GPU creati e inviati in batch quando necessario. In WebGPU, l'encoder dei comandi GPU restituito da
device.createCommandEncoder()
è l'oggetto JavaScript che genera un batch di
comandi "in buffer" che verranno inviati alla GPU in un secondo momento. I metodi su
GPUBuffer
, invece, sono "non in buffer", il che significa che vengono eseguiti in modo atomico
al momento in cui vengono chiamati.
Una volta ottenuto il codificatore dei comandi GPU, chiama copyEncoder.copyBufferToBuffer()
come mostrato di seguito per aggiungere questo comando alla coda dei comandi ed eseguirlo in un secondo momento.
Infine, completa i comandi di codifica chiamando copyEncoder.finish()
e inviali alla coda dei comandi del dispositivo GPU. La coda è responsabile della gestione dei contenuti inviati tramite device.queue.submit()
con i comandi GPU come argomenti.
In questo modo verranno eseguiti in modo atomico tutti i comandi memorizzati nell'array in ordine.
// 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]);
A questo punto, i comandi della coda della GPU sono stati inviati, ma non necessariamente eseguiti.
Per leggere il secondo buffer della GPU, chiama gpuReadBuffer.mapAsync()
con
GPUMapMode.READ
. Restituisce una promessa che si risolverà quando il buffer GPU viene mappato. Dopodiché, recupera l'intervallo mappato con gpuReadBuffer.getMappedRange()
che contiene gli stessi valori del primo buffer della GPU una volta eseguiti tutti i comandi GPU in coda.
// Read buffer.
await gpuReadBuffer.mapAsync(GPUMapMode.READ);
const copyArrayBuffer = gpuReadBuffer.getMappedRange();
console.log(new Uint8Array(copyArrayBuffer));
Puoi provare questo esempio.
In breve, ecco cosa devi ricordare in merito alle operazioni di memoria buffer:
- La mappatura dei buffer GPU deve essere annullata per poter essere utilizzati nell'invio della coda dei dispositivi.
- Una volta mappati, i buffer della GPU possono essere letti e scritti in JavaScript.
- I buffer della GPU vengono mappati quando vengono chiamati
mapAsync()
ecreateBuffer()
conmappedAtCreation
impostato su true.
Programmazione degli shader
I programmi in esecuzione sulla GPU che eseguono solo calcoli (e non disegnano triangoli) sono chiamati shader di calcolo. Vengono eseguiti in parallelo da centinaia di core GPU (più piccoli dei core della CPU) che operano insieme per elaborare i dati. I relativi input e output sono buffer in WebGPU.
Per illustrare l'utilizzo degli shader di calcolo in WebGPU, faremo pratica con la moltiplicazione di matrici, un algoritmo comune nel machine learning illustrato di seguito.
In breve, ecco cosa faremo:
- Crea tre buffer della GPU (due per le matrici da moltiplicare e uno per la matrice del risultato)
- Descrivi input e output per lo shader di calcolo
- Compila il codice dello shader di calcolo
- Configura una pipeline di computing
- Invia in batch i comandi codificati alla GPU
- Leggi il buffer della GPU della matrice di risultati
Creazione di buffer GPU
Per semplicità, le matrici saranno rappresentate come un elenco di numeri in virgola mobile. Il primo elemento è il numero di righe, il secondo è il numero di colonne e il resto sono i numeri effettivi della matrice.
I tre buffer della GPU sono buffer di archiviazione perché dobbiamo archiviare e recuperare i dati nel compute shader. Questo spiega perché i flag di utilizzo del buffer GPU includono
GPUBufferUsage.STORAGE
per tutti. Il flag di utilizzo della matrice di risultati ha anche GPUBufferUsage.COPY_SRC
perché verrà copiato in un altro buffer per la lettura una volta eseguiti tutti i comandi della coda della 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
});
Associa layout gruppo e associa gruppo
I concetti di layout del gruppo di binding e gruppo di binding sono specifici di WebGPU. Un layout di gruppo di associazioni definisce l'interfaccia di input/output prevista da uno shaker, mentre un gruppo di associazione rappresenta i dati effettivi di input/output per uno shaker.
Nell'esempio seguente, il layout del gruppo di binding prevede due buffer di archiviazione di sola lettura nelle associazioni di voci numerate 0
, 1
e un buffer di archiviazione in 2
per lo shader di calcolo.
Il gruppo di binding, invece, definito per questo layout del gruppo di binding, associa
i buffer della GPU alle voci: gpuBufferFirstMatrix
alla associazione 0
, gpuBufferSecondMatrix
alla associazione 1
e resultMatrixBuffer
alla
associazione 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
}
}
]
});
Codice dello shader di calcolo
Il codice dello streamr di computing per moltiplicare le matrici è scritto in WGSL, il linguaggio
WebGPU Shader Language, che è traducibile banalmente in SPIR-V. Senza
entrare nei dettagli, dovresti trovare sotto i tre buffer di archiviazione identificati
con var<storage>
. Il programma utilizzerà firstMatrix
e secondMatrix
come input e resultMatrix
come output.
Tieni presente che ogni buffer di archiviazione ha una decorazione binding
utilizzata che corrisponde allo stesso indice definito nei layout dei gruppi di binding e nei gruppi di binding dichiarati sopra.
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;
}
`
});
Configurazione della pipeline
La pipeline di calcolo è l'oggetto che descrive effettivamente l'operazione di calcolo che eseguiremo. Crealo chiamando device.createComputePipeline()
.
Richiede due argomenti: il layout del gruppo di binding che abbiamo creato in precedenza e una fase di calcolo che definisce il punto di ingresso del nostro shader di calcolo (la funzione WGSL main
) e il modulo shader di calcolo effettivo creato con device.createShaderModule()
.
const computePipeline = device.createComputePipeline({
layout: device.createPipelineLayout({
bindGroupLayouts: [bindGroupLayout]
}),
compute: {
module: shaderModule,
entryPoint: "main"
}
});
Invio dei comandi
Dopo aver creato un gruppo di binding con i tre buffer della GPU e una pipeline di calcolo con un layout del gruppo di binding, è il momento di utilizzarli.
Iniziamo con un codificatore di pass di calcolo programmabile con
commandEncoder.beginComputePass()
. Lo utilizzeremo per codificare i comandi GPU
che eseguiranno la moltiplicazione della matrice. Imposta la pipeline con
passEncoder.setPipeline(computePipeline)
e il relativo gruppo di associazione su un indice 0 con
passEncoder.setBindGroup(0, bindGroup)
. L'indice 0 corrisponde alla decorazione group(0)
nel codice WGSL.
Ora parliamo di come questo compute shader verrà eseguito sulla GPU. Il nostro obiettivo è eseguire questo programma in parallelo per ogni cella della matrice di risultati, passo dopo passo. Ad esempio, per una matrice di risultati di dimensioni 16 x 32, per codificare il comando di esecuzione su un @workgroup_size(8, 8)
, chiameremo passEncoder.dispatchWorkgroups(2, 4)
o passEncoder.dispatchWorkgroups(16 / 8, 32 / 8)
.
Il primo argomento "x" è la prima dimensione, il secondo "y" è la seconda dimensione e l'ultimo"z" è la terza dimensione, che per impostazione predefinita è 1 perché non ce ne serve qui.
Nel mondo del computing GPU, la codifica di un comando per l'esecuzione di una funzione kernel su un set di dati è chiamata invio.
La dimensione della griglia del gruppo di lavoro per il nostro shader di calcolo è (8, 8)
nel codice WGSL. Per questo motivo, "x" e "y", che rappresentano rispettivamente il numero di righe della prima matrice e il numero di colonne della seconda matrice, verranno divisi per 8. A questo punto, possiamo inviare una chiamata di computing con
passEncoder.dispatchWorkgroups(firstMatrix[0] / 8, secondMatrix[1] / 8)
. Il
numero di griglie di gruppi di lavoro da eseguire sono gli argomenti dispatchWorkgroups()
.
Come mostrato nel disegno sopra, ogni shader avrà accesso a un oggetto builtin(global_invocation_id)
univoco che verrà utilizzato per sapere quale cella della matrice del risultato calcolare.
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();
Per terminare il codificatore pass di calcolo, chiama passEncoder.end()
. Poi, crea un buffer GPU da utilizzare come destinazione per copiare il buffer della matrice del risultato con copyBufferToBuffer
. Infine, completa i comandi di codifica con copyEncoder.finish()
e inviali alla coda dei dispositivi GPU chiamando device.queue.submit()
con i comandi 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]);
Lettura matrice dei risultati
Leggere la matrice dei risultati è facile come chiamare gpuReadBuffer.mapAsync()
con
GPUMapMode.READ
e attendere la risoluzione della promessa restituita, che indica
che il buffer della GPU è ora mappato. A questo punto, è possibile ottenere l'intervallo mappato con gpuReadBuffer.getMappedRange()
.
Nel nostro codice, il risultato registrato nella console JavaScript di DevTools è "2, 2, 50, 60, 114, 140".
// Read buffer.
await gpuReadBuffer.mapAsync(GPUMapMode.READ);
const arrayBuffer = gpuReadBuffer.getMappedRange();
console.log(new Float32Array(arrayBuffer));
Complimenti! Ce l'hai fatta. Puoi giocare con il Sample.
Un ultimo trucco
Un modo per semplificare la lettura del codice è utilizzare il pratico metodo getBindGroupLayout
della pipeline di calcolo per dedurre il layout del gruppo di binding dal modulo shader. Questo trucco elimina la necessità di creare un layout di gruppo di associazione personalizzato e di specificare un layout di pipeline nella pipeline di computing, come mostrato di seguito.
È disponibile un'illustrazione di getBindGroupLayout
per il Sample precedente.
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: [
Risultati sul rendimento
In che modo eseguire la moltiplicazione delle matrici su una GPU è confrontata con l'esecuzione su una CPU? Per scoprirlo, ho scritto il programma appena descritto per una CPU. Come puoi vedere nel grafico seguente, l'utilizzo della potenza completa della GPU sembra una scelta ovvia quando le dimensioni delle matrici sono maggiori di 256 x 256.
Questo articolo è stato solo l'inizio del mio viaggio di esplorazione di WebGPU. A breve saranno disponibili altri articoli con approfondimenti su GPU Compute e su come funziona il rendering (canvas, texture, sampler) in WebGPU.