Questo post esplora l'API WebGPU sperimentale tramite esempi e ti aiuta a iniziare a eseguire i calcoli paralleli ai dati utilizzando la GPU.
Contesto
Come forse già saprai, la GPU (Gric Processing Unit) è un sottosistema elettronico all'interno di un computer originariamente specializzato per l'elaborazione delle immagini. Tuttavia, negli ultimi 10 anni, si è evoluta verso un'architettura più flessibile che consente agli sviluppatori di implementare molti tipi di algoritmi, non solo la visualizzazione di grafica 3D, sfruttando al contempo l'architettura unica della GPU. Queste funzionalità sono note come calcolo GPU, mentre l'utilizzo di una GPU come coprocessore per il calcolo scientifico generico è chiamato programmazione GPU generale (GPGPU).
Il calcolo tramite GPU ha contribuito in modo significativo al recente boom del machine learning, poiché le reti neurali di convoluzione e altri modelli possono sfruttare questa architettura per funzionare in modo più efficiente sulle GPU. Con l'attuale piattaforma web priva di funzionalità di calcolo GPU, il gruppo della community "GPU per il web" di W3C sta progettando un'API per esporre le moderne API GPU disponibili sulla maggior parte dei dispositivi attuali. Questa API si chiama WebGPU.
WebGPU è un'API di basso livello, come WebGL. È molto potente e molto dettagliato, come vedrai. Ma non è un problema. Vogliamo il rendimento.
In questo articolo, mi soffermerò sulla parte di elaborazione della GPU di WebGPU e, ad essere onesti, approfondirò ulteriormente l'argomento, così potrai iniziare a giocare da solo. Approfondirò e tratterò il rendering WebGPU (tela, texture e così via) nei prossimi articoli.
Accedi alla GPU
L'accesso 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. Immagina 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).
Dopo aver 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 offrono opzioni che ti consentono di specificare in modo specifico il tipo di adattatore (preferenza di alimentazione) e il dispositivo (estensioni, limiti) che vuoi. Per semplicità, utilizzeremo le opzioni predefinite in questo articolo.
Scrittura memoria buffer
Vediamo come utilizzare JavaScript per scrivere dati nella memoria per la GPU. Questo processo non è semplice a causa del modello di sandbox utilizzato nei browser web moderni.
L'esempio seguente mostra come scrivere quattro byte per eseguire il buffer della memoria 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, indichiamo esplicitamente di voler scrivere in questo buffer. Genera un oggetto buffer GPU mappato al momento della creazione grazie a mappedAtCreation
impostato su true. Quindi, il buffer di dati binari non elaborati associato può essere recuperato chiamando il metodo del buffer GPU getMappedRange()
.
La scrittura dei byte è familiare 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. Per potervi accedere, la GPU deve essere
non mappata, il che basta chiamare gpuBuffer.unmap()
.
Il concetto di mappato/non mappato è necessario per prevenire condizioni di gara in cui GPU e CPU accedono alla memoria contemporaneamente.
Lettura memoria buffer
Ora vediamo come copiare un buffer GPU in un altro buffer GPU e rileggilo.
Poiché stiamo scrivendo nel primo buffer GPU e vogliamo copiarlo in un secondo
buffer GPU, è necessario 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 suo flag di utilizzo è GPUBufferUsage.COPY_DST |
GPUBufferUsage.MAP_READ
, in quanto verrà utilizzato come destinazione del primo buffer della GPU e verrà letto in JavaScript dopo l'esecuzione dei 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 della GPU vengono eseguiti in modo asincrono. Ecco perché viene creato un elenco di comandi della GPU, che vengono inviati in batch quando necessario. In WebGPU, il codificatore dei comandi GPU restituito da device.createCommandEncoder()
è l'oggetto JavaScript che crea un batch di comandi "bufferizzati" che verranno inviati alla GPU a un certo punto. I metodi su GPUBuffer
, invece, sono "unbuffered", ovvero vengono eseguiti a livello atomico al momento della chiamata.
Una volta ottenuto il codificatore dei comandi GPU, chiama copyEncoder.copyBufferToBuffer()
come mostrato di seguito per aggiungere questo comando alla coda dei comandi per un'esecuzione successiva.
Infine, completa i comandi di codifica chiamando copyEncoder.finish()
e inviali
alla coda dei comandi del dispositivo GPU. La coda è responsabile della gestione degli invii effettuati tramite device.queue.submit()
con i comandi GPU come argomenti.
In questo modo verranno eseguiti a livello atomico tutti i comandi archiviati 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 GPU sono stati inviati, ma non necessariamente eseguiti.
Per leggere il secondo buffer GPU, chiama gpuReadBuffer.mapAsync()
con
GPUMapMode.READ
. Restituisce una promessa che verrà risolta quando il buffer GPU viene mappato. Quindi, ottieni l'intervallo mappato con gpuReadBuffer.getMappedRange()
che contiene gli stessi valori del primo buffer GPU una volta eseguiti tutti i comandi della 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:
- I buffer GPU non devono essere mappati per essere utilizzati nell'invio della coda dei dispositivi.
- Una volta mappati, i buffer GPU possono essere letti e scritti in JavaScript.
- I buffer GPU vengono mappati quando vengono chiamati
mapAsync()
ecreateBuffer()
conmappedAtCreation
impostato su true.
Programmazione Shader
I programmi in esecuzione sulla GPU che eseguono solo calcoli (e non disegnano triangoli) sono denominati mesh di calcolo. Vengono eseguite in parallelo da centinaia di core GPU (più piccoli rispetto ai core della CPU) che operano insieme per elaborare i dati. L'input e l'output sono buffer in WebGPU.
Per illustrare l'uso degli shardr di computing in WebGPU, utilizzeremo la moltiplicazione matriciale, un algoritmo comune nel machine learning illustrato di seguito.
In breve, ecco cosa faremo:
- Crea tre buffer GPU (due per moltiplicare le matrici e uno per la matrice dei risultati)
- Descrivere input e output per loshar di calcolo
- Compila il codice Shaker di elaborazione
- Configura una pipeline di computing
- Invia in gruppo i comandi codificati alla GPU
- Leggi il buffer GPU della matrice dei risultati
Creazione 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 è il numero effettivo della matrice.
I tre buffer GPU sono i buffer di archiviazione, in quanto dobbiamo archiviare e recuperare i dati nello Shadr Compute. Questo spiega perché i flag di utilizzo del buffer GPU includono GPUBufferUsage.STORAGE
per tutti. Il flag di utilizzo della matrice di risultati contiene anche
GPUBufferUsage.COPY_SRC
perché verrà copiato in un altro buffer per la lettura una volta eseguiti tutti i comandi della coda 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 gruppo di associazione e gruppo di associazione sono specifici di WebGPU. Un layout gruppo di associazione definisce l'interfaccia di input/output prevista da uno shabby, mentre un gruppo di associazione rappresenta i dati effettivi di input/output per uno mesh.
Nell'esempio riportato di seguito, il layout del gruppo di associazione prevede due buffer di archiviazione di sola lettura alle associazioni di voci numerate 0
, 1
e un buffer di archiviazione in 2
per lo Shadr Compute.
Il gruppo di associazione, invece, definito per questo layout del gruppo di associazione, associa i
buffer GPU alle voci: gpuBufferFirstMatrix
all'associazione 0
, gpuBufferSecondMatrix
all'associazione 1
e resultMatrixBuffer
all'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
}
}
]
});
Calcola codice shabbyr
Il codice Shader di calcolo per moltiplicare le matrici è scritto in WGSL, il
linguaggio Shader WebGPU, che è banalmente traducibile in SPIR-V. Senza
andare nel dettaglio, dovresti trovare sotto i tre buffer di archiviazione identificati
con var<storage>
. Il programma userà firstMatrix
e secondMatrix
come
input e resultMatrix
come output.
Tieni presente che per ogni buffer di archiviazione viene utilizzata una decorazione binding
che corrisponde allo stesso indice definito nei layout dei gruppi di associazione e nei gruppi di associazione 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 computing è l'oggetto che descrive effettivamente l'operazione
di calcolo che eseguiremo. Creala chiamando il numero device.createComputePipeline()
.
Prevede due argomenti: il layout del gruppo di associazione creato in precedenza e una fase di calcolo che definisce il punto di ingresso del nostro mesh di computing (la funzione WGSL di main
) e l'effettivo modulo di Shadr Compute creato con device.createShaderModule()
.
const computePipeline = device.createComputePipeline({
layout: device.createPipelineLayout({
bindGroupLayouts: [bindGroupLayout]
}),
compute: {
module: shaderModule,
entryPoint: "main"
}
});
Invio dei comandi
Dopo aver creato l'istanza di un gruppo di associazione con i nostri tre buffer GPU e una pipeline di calcolo con un layout di gruppo di associazione, è il momento di utilizzarli.
Avviamo un codificatore pass programmabile per il computing con
commandEncoder.beginComputePass()
. La useremo per codificare i comandi della GPU
che effettueranno la moltiplicazione della matrice. Imposta la pipeline con passEncoder.setPipeline(computePipeline)
e il relativo gruppo di associazione all'indice 0 con passEncoder.setBindGroup(0, bindGroup)
. L'indice 0 corrisponde alla decorazione group(0)
nel codice WGSL.
Ora parliamo di come verrà eseguito questo meshr di computing sulla GPU. Il nostro obiettivo è eseguire questo programma in parallelo per ogni cella della matrice dei risultati, passo dopo passo. Ad esempio, per una matrice di risultati di dimensione 16 x 32, per codificare
il comando di esecuzione, su un elemento @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 il cui valore predefinito è 1, in quanto qui non è necessario.
Nel mondo di computing delle GPU, la codifica di un comando per eseguire una funzione del kernel su un set di dati è chiamata invio.
La dimensione della griglia del gruppo di lavoro per il nostro probe di computing è (8, 8)
nel nostro codice WGSL. Per questo motivo, "x" e "y", che corrispondono rispettivamente al numero di righe della prima matrice e del numero di colonne della seconda matrice, saranno divise per 8. Ora possiamo inviare una chiamata di computing con
passEncoder.dispatchWorkgroups(firstMatrix[0] / 8, secondMatrix[1] / 8)
. Il numero di griglie del gruppo di lavoro da eseguire corrisponde agli argomenti dispatchWorkgroups()
.
Come mostrato nel disegno qui sopra, ogni Shader avrà accesso a un oggetto builtin(global_invocation_id)
univoco che verrà utilizzato per sapere quale cella della matrice di risultati 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 l'encoder Compute Pass, chiama passEncoder.end()
. Quindi, crea un buffer GPU da utilizzare come destinazione per copiare il buffer della matrice dei risultati 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]);
Leggi la matrice dei risultati
Leggere la matrice dei risultati è facile come chiamare gpuReadBuffer.mapAsync()
con
GPUMapMode.READ
e attendere la promessa di ritorno di risoluzione che indica
che il buffer GPU è ora mappato. A questo punto, è possibile ottenere l'intervallo
mappato con gpuReadBuffer.getMappedRange()
.
Nel nostro codice, il risultato registrato nella console JavaScript 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 riprodurre l'anteprima.
Un ultimo trucco
Un modo per semplificare la lettura del codice è utilizzare il pratico
metodo getBindGroupLayout
della pipeline di computing per dedurre il layout del gruppo di associazione
dal modulo shabbyr. Questo trucco elimina la necessità di creare un layout di gruppi di associazione personalizzato e di specificarne uno nella pipeline di computing, come puoi vedere di seguito.
Un'illustrazione di getBindGroupLayout
per l'esempio precedente è disponibile.
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 relativi alle prestazioni
Quindi, come si rapporta l'esecuzione della moltiplicazione della matrice su una GPU all'esecuzione su una CPU? Per scoprirlo, ho scritto il programma appena descritto per una CPU. Come vedi nel grafico qui sotto, usare tutta la potenza della GPU sembra una scelta ovvia quando la dimensione delle matrici è maggiore di 256 x 256.
Questo articolo è solo l'inizio del mio viaggio a esplorare WebGPU. Presto arriveranno altri articoli con ulteriori approfondimenti sul calcolo delle GPU e sul funzionamento del rendering (canvas, texture, sampler) in WebGPU.