Inizia a utilizzare GPU Compute sul web

Questo post esplora l'API WebGPU sperimentale attraverso esempi e aiuta si inizia a eseguire calcoli paralleli ai dati utilizzando la GPU.

François Beaufort
François Beaufort
GitHub

Sfondo

Come forse già saprai, la GPU (Graphic Processing Unit) è un'unità un sottosistema di un computer originariamente specializzato per l'elaborazione le immagini. Tuttavia, negli ultimi 10 anni, si è evoluto verso una maggiore flessibilità che consente agli sviluppatori di implementare molti tipi di algoritmi, non solo il rendering di grafica 3D, sfruttando al contempo l'architettura unica GPU. Queste funzionalità sono chiamate computing GPU e utilizzano una GPU come coprocessore per il calcolo scientifico per uso generico è chiamato programmazione GPU (GPGPU).

GPU Compute ha contribuito in modo significativo al recente boom del machine learning, poiché le reti neurali di convoluzione e altri modelli possono sfruttare per l'esecuzione in modo più efficiente sulle GPU. Con l'attuale piattaforma web non dispone di capacità di calcolo GPU, la GPU "GPU for the Web" del W3C Gruppo della community è la progettazione di un'API per esporre le moderne API GPU disponibili dispositivi attuali. Questa API è denominata WebGPU.

WebGPU è un'API di basso livello, come WebGL. È molto potente e piuttosto dettagliato, che vedrai. Ma va bene così. Quello che vogliamo è il rendimento.

In questo articolo mi soffermerò sulla parte di GPU Compute di WebGPU. onestamente, sto solo grattando la superficie, così potrai iniziare a giocare sui tuoi personali. Approfondiremo l'argomento e parlerò del rendering di WebGPU (canvas, texture, e così via) nei prossimi articoli.

Accedi alla GPU

L'accesso alla GPU è facile in WebGPU. Chiamata a navigator.gpu.requestAdapter() restituisce una promessa JavaScript che si risolverà in modo asincrono con una GPU dell'adattatore. Considera questo adattatore come la scheda grafica. Può essere integrato (sullo stesso chip della CPU) o discreto (di solito una scheda PCIe che prestazioni migliori, ma consumano più energia).

Una volta che hai l'adattatore GPU, chiama adapter.requestDevice() per ottenere una promessa che si risolverà con un dispositivo GPU che userai 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 consentono di specificare l'alimentatore (preferenza alimentazione) e il dispositivo (estensioni, limiti) che desideri. Per Per semplicità, utilizzeremo le opzioni predefinite riportate in questo articolo.

Memoria buffer di scrittura

Vediamo come usare JavaScript per scrivere dati nella memoria della GPU. Questo non è semplice per via del modello di sandbox utilizzato nel web moderno browser.

L'esempio seguente mostra come scrivere 4 byte nella memoria di 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 sono richiesti per questa specifica chiamata, vogliamo chiarire che vogliamo scrivere a questo buffer. Risulta in un oggetto buffer GPU mappato al momento della creazione grazie mappedAtCreation impostato su true. Il buffer di dati binari non elaborati associato può quindi recuperabile chiamando il metodo di buffer GPU getMappedRange().

È possibile scrivere byte se hai già giocato con 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 GPU è mappato, ovvero è di proprietà della CPU, è accessibile in lettura/scrittura da JavaScript. Per consentire alla GPU di accedervi, deve essere rimosso, il che è semplice come chiamare gpuBuffer.unmap().

Il concetto di mappato/non mappato è necessario per impedire le gare in cui le GPU e alla memoria di accesso alla CPU.

Memoria buffer di lettura

Vediamo ora come copiare un buffer GPU in un altro buffer GPU e rileggere il buffer.

Poiché stiamo scrivendo nel primo buffer GPU e vogliamo copiarlo Buffer GPU, è richiesto un nuovo flag di utilizzo GPUBufferUsage.COPY_SRC. Il secondo Il buffer GPU è stato creato in uno stato non mappato questa volta con device.createBuffer(). Il suo flag di utilizzo è GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ perché verrà usato come destinazione della prima GPU buffer e leggi 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, vengono eseguiti tutti i comandi GPU in modo asincrono. Ecco perché è presente un elenco di comandi GPU creati e inviati i batch quando necessario. In WebGPU, l'encoder dei comandi GPU restituito device.createCommandEncoder()è l'oggetto JavaScript che crea un batch di "con buffer" che a un certo punto verranno inviati alla GPU. I metodi su GPUBuffer, invece, sono "senza buffer", il che significa che vengono eseguiti atomicamente nel 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 per un'esecuzione successiva. Infine, completa i comandi di codifica richiamando copyEncoder.finish() e invia alla coda dei comandi del dispositivo GPU. La coda è responsabile della gestione 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 si risolverà quando il buffer GPU viene mappato. Quindi ottieni l'intervallo mappato con gpuReadBuffer.getMappedRange() contiene gli stessi valori del primo buffer GPU una volta che tutti i comandi GPU sono in coda sono state eseguite.

// Read buffer.
await gpuReadBuffer.mapAsync(GPUMapMode.READ);
const copyArrayBuffer = gpuReadBuffer.getMappedRange();
console.log(new Uint8Array(copyArrayBuffer));

Puoi provare questo esempio.

In breve, ecco cosa bisogna ricordare per quanto riguarda le operazioni di memoria del buffer:

  • La mappatura dei buffer GPU deve essere annullata per poter 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 mapAsync() e createBuffer() con Vengono chiamati i criteri mappedAtCreation impostati su true.

Programmazione Shader

Programmi in esecuzione sulla GPU che eseguono solo calcoli (e non disegnano triangoli) sono chiamati consumatori di computing. vengono eseguite in parallelo di core GPU (che sono più piccoli dei core della CPU) che operano insieme per e i dati di Google Cloud. I rispettivi input e output sono buffer in WebGPU.

Per illustrare l'uso degli shaker di computing in WebGPU, giochiamo con Matrix la moltiplicazione, un algoritmo comune nel machine learning illustrato di seguito.

Diagramma di moltiplicazione matriciale
Diagramma di moltiplicazione delle matrici

In breve, ecco cosa faremo:

  1. Crea tre buffer GPU (due per le matrici da moltiplicare e uno per matrice dei risultati)
  2. Descrivere l'input e l'output per Compute Skillsr
  3. Compila il codice Compute Shar
  4. Configura una pipeline di computing
  5. Invia in batch i comandi codificati alla GPU
  6. Leggere il buffer GPU della matrice dei risultati

Creazione di buffer GPU

Per semplicità, le matrici saranno rappresentate come un elenco di variabili mobili numeri in punti. Il primo elemento è il numero di righe, il secondo elemento numero di colonne, mentre il resto sono i numeri effettivi della matrice.

Semplice rappresentazione 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 buffer di archiviazione perché dobbiamo archiviare e recuperare i dati Compute Dataflow. Questo spiega perché i flag di utilizzo del buffer GPU includono GPUBufferUsage.STORAGE per tutti. Anche il flag di utilizzo della matrice dei risultati GPUBufferUsage.COPY_SRC perché verrà copiata in un altro buffer per dopo l'esecuzione di 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 del gruppo di associazione e del gruppo di associazione sono specifici di WebGPU. Un vincolo il layout dei gruppi definisce l'interfaccia di input/output prevista da uno Shar, mentre Il gruppo di associazione rappresenta i dati di input/output effettivi per uno shaker.

Nell'esempio riportato di seguito, il layout del gruppo di associazione prevede due buffer di archiviazione di sola lettura Le voci numerate associano 0, 1 e un buffer di archiviazione in 2 per il computing shaker. Il gruppo di associazione, invece, definito per questo layout del gruppo di associazione, associa Buffer GPU nelle voci: gpuBufferFirstMatrix all'associazione 0, gpuBufferSecondMatrix all'associazione 1 e resultMatrixBuffer all'associazione 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 Shar Compute

Il codice dello shaker di computing per moltiplicare le matrici è scritto in WGSL, WebGPU Shader Language, che è banalmente traducibile in SPIR-V. Senza più in dettaglio, dovresti trovare più in basso i tre buffer di archiviazione identificati con var<storage>. Il programma userà firstMatrix e secondMatrix come e resultMatrix come output.

Tieni presente che ogni buffer di archiviazione ha una decorazione binding utilizzata che corrisponde lo 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 computing che eseguiremo. Per crearlo, chiama device.createComputePipeline(). Prende due argomenti: il layout del gruppo di associazione creato in precedenza e fase che definisce il punto di ingresso del nostro shaker di computing (la funzione WGSL di main) e l'effettivo modulo Compute Skills 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'istanza per un gruppo di associazione con i nostri tre buffer GPU e con un layout di gruppo di associazione, è il momento di usarli.

Avviamo un codificatore di pass di computing programmabile con commandEncoder.beginComputePass(). Lo utilizzeremo per codificare i comandi GPU che eseguirà la moltiplicazione matriciale. Imposta la sua 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 Shader di computing sulla GPU. Le nostre l'obiettivo è eseguire questo programma in parallelo per ogni cella della matrice dei risultati, passo passo. Ad esempio, per una matrice di risultati di dimensione 16 x 32, per codificare il comando di esecuzione, su un @workgroup_size(8, 8), chiamiamo passEncoder.dispatchWorkgroups(2, 4) o passEncoder.dispatchWorkgroups(16 / 8, 32 / 8). Il primo argomento "x" è la prima dimensione, la seconda "y" è la seconda dimensione, e l'ultima "z" è la terza dimensione predefinita di 1, in quanto non è necessaria 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.

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 lo shaker di computing è (8, 8) nel nostro WGSL le API nel tuo codice. Per questo motivo, "x" e "y" che corrispondono rispettivamente al numero di righe la prima matrice e il numero di colonne della seconda matrice verranno divisi per 8. A questo punto possiamo inviare una chiamata di computing passEncoder.dispatchWorkgroups(firstMatrix[0] / 8, secondMatrix[1] / 8). La del numero di griglie del gruppo di lavoro da eseguire sono gli argomenti dispatchWorkgroups().

Come si vede nel disegno qui sopra, ogni Shader avrà accesso a un builtin(global_invocation_id) oggetto che verrà utilizzato per sapere quale risultato cella della matrice da 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 Compute Pass, chiama passEncoder.end(). Quindi, crea un'istanza 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 inviale 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 in attesa della risoluzione della promessa di ritorno, che indica il buffer GPU è mappato. A questo punto, è possibile ottenere la mappatura intervallo con gpuReadBuffer.getMappedRange().

Risultato della moltiplicazione della matrice
Risultato della moltiplicazione della matrice

Nel nostro codice, il risultato collegato alla 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 riprodurre il Sample.

Un ultimo trucco

Un modo per semplificare la lettura del codice consiste nell'utilizzare la pratica Metodo getBindGroupLayout della pipeline di computing per inferire il gruppo di associazione del layout del modulo shar. Questo trucco elimina la necessità di creare layout dei gruppi di associazione personalizzato e specifica un layout della pipeline nelle fasi di computing come mostrato di seguito.

È disponibile un'illustrazione di getBindGroupLayout per l'esempio 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 delle matrici su una GPU e l'esecuzione su una di CPU? Per scoprirlo, ho scritto il programma appena descritto per una CPU. E il più possibile guarda nel grafico qui sotto, utilizzare 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 è stato solo l'inizio del mio viaggio all'esplorazione di WebGPU. Ottieni di più a breve con ulteriori approfondimenti sul computing GPU e sul modo in cui (canvas, texture, sampler) funziona in WebGPU.