Inizia a utilizzare GPU Compute sul web

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

François Beaufort
François Beaufort

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() e createBuffer() con mappedAtCreation 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.

Diagramma di moltiplicazione delle matrici
Diagramma di moltiplicazione della matrice

In breve, ecco cosa faremo:

  1. Crea tre buffer GPU (due per moltiplicare le matrici e uno per la matrice dei risultati)
  2. Descrivere input e output per loshar di calcolo
  3. Compila il codice Shaker di elaborazione
  4. Configura una pipeline di computing
  5. Invia in gruppo i comandi codificati alla GPU
  6. 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.

Rappresentazione semplice di una matrice in JavaScript e del suo equivalente in notazione matematica
Rappresentazione semplice di una matrice in JavaScript e del suo equivalente in notazione matematica

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.

Esecuzione in parallelo per ogni cella della matrice dei risultati
Esecuzione in parallelo per ogni cella della matrice dei risultati

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

Risultato della moltiplicazione della matrice
Risultato della moltiplicazione della matrice

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.

Benchmark GPU e CPU
Benchmark GPU e CPU

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.