Questo post esplora l'API WebGPU sperimentale attraverso esempi e ti aiuta a iniziare a eseguire calcoli paralleli sui dati utilizzando la GPU.
Pubblicato il: 28 agosto 2019, ultimo aggiornamento: 12 agosto 2025
Sfondo
Come forse già saprai, la GPU (Graphics Processing Unit) è un sottosistema elettronico all'interno di un computer originariamente specializzato nell'elaborazione della grafica. Tuttavia, negli ultimi 10 anni si è evoluta verso un'architettura più flessibile che consente agli sviluppatori di implementare molti tipi di algoritmi, non solo il rendering di grafica 3D, sfruttando al contempo l'architettura unica della GPU. Queste funzionalità sono chiamate GPU Compute e l'utilizzo di una GPU come coprocessore per il calcolo scientifico per uso generico è chiamato programmazione GPU per uso generico (GPGPU).
GPU Compute ha contribuito in modo significativo al recente boom del machine learning, in quanto le reti neurali convoluzionali e altri modelli possono sfruttare l'architettura per essere eseguiti in modo più efficiente sulle GPU. Poiché l'attuale piattaforma web non dispone di funzionalità di calcolo GPU, il gruppo della community "GPU for the Web" del W3C sta progettando un'API per esporre le moderne API GPU disponibili sulla maggior parte dei dispositivi attuali. Questa API è chiamata WebGPU.
WebGPU è un'API di basso livello, come WebGL. È molto potente e piuttosto dettagliato, come vedrai. Ma non importa. Ciò che cerchiamo è il rendimento.
In questo articolo mi concentrerò sulla parte di calcolo della GPU di WebGPU e, a dire il vero, mi limiterò a grattare la superficie, in modo che tu possa iniziare a giocare da solo. Approfondirò l'argomento e tratterò il rendering WebGPU (canvas, texture, ecc.) nei prossimi articoli.
Accedere alla GPU
Accedere alla GPU è facile in WebGPU. La chiamata di navigator.gpu.requestAdapter() restituisce una promessa JavaScript che verrà risolta in modo asincrono con una scheda GPU. Considera 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 della 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 dispositivo (estensioni, limiti) che desideri. Per semplicità, in questo articolo utilizzeremo le opzioni predefinite.
Memoria buffer di scrittura
Vediamo come utilizzare JavaScript per scrivere dati nella memoria della GPU. Questo processo non è semplice a causa del modello di sandboxing utilizzato nei moderni browser web.
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, specifichiamo che vogliamo scrivere in questo buffer. Il risultato è un oggetto buffer GPU mappato al momento della creazione grazie a mappedAtCreation impostato su true. Il buffer di dati binari non elaborati associato può essere recuperato chiamando il metodo del buffer GPU getMappedRange().
La scrittura di byte è familiare se hai già utilizzato ArrayBuffer; usa 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 della GPU è mappato, il che significa che è di proprietà della CPU ed è accessibile in lettura/scrittura da JavaScript. Affinché la GPU possa accedervi, deve essere annullata la mappatura, operazione semplice come chiamare gpuBuffer.unmap().
Il concetto di mappato/non mappato è necessario per evitare condizioni di competizione in cui la GPU e la CPU accedono alla memoria contemporaneamente.
Memoria buffer di lettura
Ora vediamo come copiare un buffer GPU in un altro buffer GPU e leggerlo.
Poiché stiamo scrivendo nel primo buffer della GPU e vogliamo copiarlo in un secondo buffer della GPU, è necessario un nuovo flag di utilizzo GPUBufferUsage.COPY_SRC. Il secondo buffer GPU viene creato in uno stato non mappato questa volta con device.createBuffer(). Il relativo flag di utilizzo è GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ, in quanto verrà utilizzato come destinazione del primo buffer della 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. Per questo motivo, viene creato e inviato un elenco di comandi GPU in batch quando necessario. In WebGPU, il codificatore di comandi GPU restituito da device.createCommandEncoder()è l'oggetto JavaScript che crea un batch di comandi "memorizzati nel buffer" che verranno inviati alla GPU a un certo punto. I metodi su GPUBuffer, invece, sono "senza buffer", il che significa che vengono eseguiti in modo atomico nel momento in cui vengono chiamati.
Una volta ottenuto il codificatore di comandi della GPU, chiama copyEncoder.copyBufferToBuffer() come mostrato di seguito per aggiungere questo comando alla coda di comandi per l'esecuzione successiva. Infine, termina i comandi di codifica chiamando copyEncoder.finish() e invia questi alla coda di 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, tutti i comandi memorizzati nell'array verranno eseguiti in modo atomico e in ordine.
// Encode commands for copying buffer to buffer. const copyEncoder = device.createCommandEncoder(); copyEncoder.copyBufferToBuffer(gpuWriteBuffer, gpuReadBuffer); // 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 della GPU, chiama gpuReadBuffer.mapAsync() con GPUMapMode.READ. Restituisce una promessa che verrà risolta quando il buffer della GPU viene mappato. Poi ottieni l'intervallo mappato con gpuReadBuffer.getMappedRange() che contiene gli stessi valori del primo buffer 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 del buffer:
- I buffer GPU devono essere annullati per essere utilizzati nell'invio della coda del dispositivo.
- 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()conmappedAtCreationimpostato 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 CPU) che operano insieme per elaborare i dati. L'input e l'output sono buffer in WebGPU.
Per illustrare l'utilizzo degli shader di calcolo in WebGPU, giocheremo con la moltiplicazione di matrici, un algoritmo comune nel machine learning illustrato di seguito.
In breve, ecco cosa faremo:
- Crea tre buffer GPU (due per le matrici da moltiplicare e uno per la matrice dei risultati)
- Descrivi l'input e l'output per lo shader di calcolo
- Compila il codice dello shader di calcolo
- Configura una pipeline di calcolo
- Invia in batch i comandi codificati alla GPU
- Leggi il buffer della GPU della matrice dei risultati
Creazione di buffer GPU
Per semplicità, le matrici verranno 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 GPU sono buffer di archiviazione perché dobbiamo archiviare e recuperare i dati nello shader di calcolo. Questo spiega perché i flag di utilizzo del buffer GPU includono GPUBufferUsage.STORAGE per tutti. Il flag di utilizzo della matrice dei 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 }); Layout del gruppo di binding e gruppo di binding
I concetti di layout del gruppo di binding e gruppo di binding sono specifici di WebGPU. Un layout del gruppo di binding definisce l'interfaccia di input/output prevista da uno shader, mentre un gruppo di binding rappresenta i dati di input/output effettivi per uno shader.
Nell'esempio seguente, il layout del gruppo di binding prevede due buffer di archiviazione di sola lettura in corrispondenza dei binding di voci numerati 0, 1 e un buffer di archiviazione in corrispondenza di 2 per lo shader di calcolo. Il gruppo di binding, invece, definito per questo layout, associa i buffer della GPU alle voci: gpuBufferFirstMatrix al binding 0, gpuBufferSecondMatrix al binding 1 e resultMatrixBuffer al binding 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: gpuBufferFirstMatrix }, { binding: 1, resource: gpuBufferSecondMatrix }, { binding: 2, resource: resultMatrixBuffer } ] }); Codice dello shader di calcolo
Il codice dello shader di calcolo per la moltiplicazione delle matrici è scritto in WGSL, il WebGPU Shader Language, che è facilmente traducibile in SPIR-V. Senza entrare nei dettagli, di seguito troverai 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 il numero device.createComputePipeline(). Accetta due argomenti: il layout del gruppo di binding creato in precedenza e una fase di calcolo che definisce il punto di ingresso dello shader di calcolo (la funzione main WGSL) e il modulo shader di calcolo effettivo creato con device.createShaderModule().
const computePipeline = device.createComputePipeline({ layout: device.createPipelineLayout({ bindGroupLayouts: [bindGroupLayout] }), compute: { module: shaderModule } }); Invio dei comandi
Dopo aver creato un gruppo di binding con i nostri tre buffer 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 della GPU che eseguiranno la moltiplicazione delle matrici. Imposta la pipeline con passEncoder.setPipeline(computePipeline) e il relativo gruppo di binding 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 shader di calcolo sulla GPU. Il nostro obiettivo è eseguire questo programma in parallelo per ogni cella della matrice dei risultati, passo dopo passo. Per una matrice di risultati di dimensioni 16 x 32, ad esempio, per codificare il comando di esecuzione, su un @workgroup_size(8, 8), chiameremmo 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 ci serve qui. Nel mondo del calcolo GPU, la codifica di un comando per eseguire una funzione kernel su un insieme di dati viene chiamata dispatching.
La dimensione della griglia del gruppo di lavoro per il nostro shader di calcolo è (8, 8) nel nostro codice WGSL. Per questo motivo, "x" e "y", che sono 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 calcolo con passEncoder.dispatchWorkgroups(firstMatrix[0] / 8, secondMatrix[1] / 8). Il numero di griglie del gruppo 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 dei 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 il codificatore di tessere di calcolo, chiama passEncoder.end(). Poi, crea un buffer GPU da utilizzare come destinazione per copiare il buffer della matrice dei risultati con copyBufferToBuffer. Infine, termina i comandi di codifica con copyEncoder.finish() e inviali alla coda del dispositivo 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, gpuReadBuffer); // Submit GPU commands. const gpuCommands = commandEncoder.finish(); device.queue.submit([gpuCommands]); Leggere la matrice dei risultati
Leggere la matrice dei risultati è semplice come chiamare gpuReadBuffer.mapAsync() con GPUMapMode.READ e attendere che la promessa di ritorno si risolva, il 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 provare il campione.
Un ultimo trucco
Un modo per rendere il codice più facile da leggere è 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 binding personalizzato e specificare un layout della pipeline nella pipeline di calcolo, come puoi vedere di seguito.
Un'illustrazione di getBindGroupLayout per il campione 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 al rendimento
Quindi, come si confronta l'esecuzione della moltiplicazione di matrici su una GPU con l'esecuzione su una CPU? Per scoprirlo, ho scritto il programma appena descritto per una CPU. Come puoi vedere nel grafico sottostante, utilizzare tutta la potenza della GPU sembra una scelta ovvia quando le dimensioni delle matrici sono superiori a 256 x 256.
Questo articolo è stato solo l'inizio del mio viaggio alla scoperta di WebGPU. A breve saranno disponibili altri articoli con approfondimenti sul calcolo della GPU e sul funzionamento del rendering (canvas, texture, campionatore) in WebGPU.