Inizia a utilizzare GPU Compute sul web

Questo post esplora l'API WebGPU sperimentale tramite esempi e ti aiuta a iniziare a eseguire calcoli paralleli dei dati utilizzando la GPU.

François Beaufort

Sfondo

Come forse già saprai, la GPU (Graphics Processing Unit) è un sottosistema elettronico all'interno di un computer originariamente specializzato per l'elaborazione della grafica. Tuttavia, negli ultimi 10 anni si è evoluto 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 calcolo GPU e l'utilizzo di una GPU come coprocessore per il calcolo scientifico generico è chiamato programmazione GPU (GPGPU) generica.

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. 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 (nello 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 relativo 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().

Se hai già utilizzato ArrayBuffer, la scrittura di byte ti sarà familiare. 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 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 smappata, il che è semplicissimo chiamando 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 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 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 crea un batch di comandi "in buffer" che verranno inviati alla GPU a un certo punto. I metodi su GPUBuffer, invece, sono "non con buffer", il che significa che vengono eseguiti in modo atomico al momento in cui vengono chiamati.

Una volta ottenuto l'encoder dei comandi GPU, chiama copyEncoder.copyBufferToBuffer() come mostrato di seguito per aggiungere questo comando alla coda dei comandi per l'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 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 verrà risolta quando il buffer della 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:

• I buffer della GPU devono essere rimappati 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() e createBuffer() con mappedAtCreation 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 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 qualche esperimento con la moltiplicazione di matrici, un algoritmo comune nel machine learning illustrato di seguito.

In breve, ecco cosa faremo:

1. Crea tre buffer della GPU (due per le matrici da moltiplicare e uno per la matrice del risultato)
2. Descrivi input e output per lo shader di calcolo
3. Compila il codice dello shader di calcolo
4. Configura una pipeline di calcolo
5. Invia in batch i comandi codificati alla GPU
6. Leggi il buffer della GPU della matrice di risultati

Creazione di buffer GPU

Per semplicità, le matrici verranno rappresentate come un elenco di numeri con 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é gli indicatori di utilizzo del buffer della GPU includonoGPUBufferUsage.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
});

Layout del gruppo di associazione e gruppo di associazione

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 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 shader di calcolo per la moltiplicazione delle matrici è scritto in WGSL, il linguaggio shader WebGPU, che è facilmente traducibile in SPIR-V. Senza entrare nei dettagli, di seguito dovresti trovare 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 di 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.

Avvia un codificatore di pass di calcolo programmabile con commandEncoder.beginComputePass(). Lo utilizzeremo per codificare i comandi GPU che eseguiranno la moltiplicazione delle matrici. Imposta la pipeline con passEncoder.setPipeline(computePipeline) e il 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 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 dell'elaborazione GPU, la codifica di un comando per eseguire una funzione del kernel su un insieme 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. Ora possiamo inviare una chiamata di calcolo 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]);

Leggi la 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 personalizzato del gruppo di associazione e di specificare un layout della pipeline nella pipeline di calcolo, come puoi vedere 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

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 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 che approfondiranno il calcolo GPU e il funzionamento del rendering (canvas, texture, sampler) in WebGPU.