Ga aan de slag met GPU Compute op internet

Dit bericht verkent de experimentele WebGPU API aan de hand van voorbeelden en helpt u aan de slag te gaan met het uitvoeren van gegevensparallelle berekeningen met behulp van de GPU.

François Beaufort

Achtergrond

Zoals u wellicht al weet, is de Graphic Processing Unit (GPU) een elektronisch subsysteem binnen een computer dat oorspronkelijk gespecialiseerd was in het verwerken van grafische afbeeldingen. In de afgelopen tien jaar is het echter geëvolueerd naar een flexibelere architectuur waardoor ontwikkelaars vele soorten algoritmen kunnen implementeren, niet alleen 3D-graphics weergeven, terwijl ze profiteren van de unieke architectuur van de GPU. Deze mogelijkheden worden GPU Compute genoemd, en het gebruik van een GPU als coprocessor voor algemeen wetenschappelijk computergebruik wordt General Purpose GPU (GPGPU)-programmering genoemd.

GPU Compute heeft aanzienlijk bijgedragen aan de recente hausse op het gebied van machine learning, omdat convolutie-neurale netwerken en andere modellen kunnen profiteren van de architectuur om efficiënter op GPU's te draaien. Omdat het huidige webplatform geen GPU Compute-mogelijkheden heeft, ontwerpt de W3C's "GPU for the Web" Community Group een API om de moderne GPU API's bloot te leggen die beschikbaar zijn op de meeste huidige apparaten. Deze API heet WebGPU .

WebGPU is een API op laag niveau, zoals WebGL. Het is zeer krachtig en behoorlijk uitgebreid, zoals je zult zien. Maar dat is oké. Wat we zoeken zijn prestaties.

In dit artikel ga ik me concentreren op het GPU Compute-gedeelte van WebGPU en, om eerlijk te zijn, ben ik nog maar aan het begin, zodat je zelf kunt beginnen met spelen. Ik zal dieper duiken en WebGPU-rendering (canvas, textuur, enz.) behandelen in komende artikelen.

Toegang tot de GPU

Toegang tot de GPU is eenvoudig in WebGPU. Het aanroepen van navigator.gpu.requestAdapter() retourneert een JavaScript-belofte die asynchroon wordt opgelost met een GPU-adapter. Beschouw deze adapter als de grafische kaart. Het kan geïntegreerd zijn (op dezelfde chip als de CPU) of discreet (meestal een PCIe-kaart die beter presteert maar meer stroom verbruikt).

Zodra u de GPU-adapter hebt, roept u adapter.requestDevice() aan om een ​​belofte te krijgen die zal worden opgelost met een GPU-apparaat dat u gaat gebruiken om GPU-berekeningen uit te voeren.

const adapter = await navigator.gpu.requestAdapter();
if (!adapter) { return; }
const device = await adapter.requestDevice();

Beide functies maken gebruik van opties waarmee u specifiek kunt aangeven welk type adapter (stroomvoorkeur) en apparaat (extensies, limieten) u wilt. Voor de eenvoud gebruiken we de standaardopties in dit artikel.

Buffergeheugen schrijven

Laten we eens kijken hoe we JavaScript kunnen gebruiken om gegevens naar het geheugen te schrijven voor de GPU. Dit proces is niet eenvoudig vanwege het sandbox-model dat in moderne webbrowsers wordt gebruikt.

In het onderstaande voorbeeld ziet u hoe u vier bytes schrijft naar buffergeheugen dat toegankelijk is vanaf de GPU. Het roept device.createBuffer() aan, die de grootte van de buffer en het gebruik ervan aanneemt. Ook al is de gebruiksvlag GPUBufferUsage.MAP_WRITE niet vereist voor deze specifieke aanroep, laten we duidelijk maken dat we naar deze buffer willen schrijven. Het resulteert in een GPU-bufferobject dat bij het maken in kaart is gebracht dankzij mappedAtCreation ingesteld op true. Vervolgens kan de bijbehorende ruwe binaire gegevensbuffer worden opgehaald door de GPU-buffermethode getMappedRange() aan te roepen.

Het schrijven van bytes is bekend als je al met ArrayBuffer hebt gespeeld; gebruik een TypedArray en kopieer de waarden daarin.

// 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]);

Op dit punt wordt de GPU-buffer in kaart gebracht, wat betekent dat deze eigendom is van de CPU en toegankelijk is via lezen/schrijven vanuit JavaScript. Om ervoor te zorgen dat de GPU er toegang toe heeft, moet deze worden ontkoppeld, wat net zo eenvoudig is als het aanroepen van gpuBuffer.unmap() .

Het concept van mapped/unmapped is nodig om raceomstandigheden te voorkomen waarbij GPU en CPU tegelijkertijd toegang hebben tot het geheugen.

Buffergeheugen lezen

Laten we nu kijken hoe we een GPU-buffer naar een andere GPU-buffer kunnen kopiëren en teruglezen.

Omdat we in de eerste GPU-buffer schrijven en deze naar een tweede GPU-buffer willen kopiëren, is een nieuwe gebruiksvlag GPUBufferUsage.COPY_SRC vereist. De tweede GPU-buffer wordt deze keer in een niet-toegewezen staat gemaakt met device.createBuffer() . De gebruiksvlag is GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ omdat het zal worden gebruikt als de bestemming van de eerste GPU-buffer en in JavaScript zal worden gelezen zodra GPU-kopieeropdrachten zijn uitgevoerd.

// 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
});

Omdat de GPU een onafhankelijke coprocessor is, worden alle GPU-opdrachten asynchroon uitgevoerd. Daarom is er een lijst met GPU-opdrachten opgebouwd en indien nodig in batches verzonden. In WebGPU is de GPU-opdracht-encoder die wordt geretourneerd door device.createCommandEncoder() het JavaScript-object dat een batch "gebufferde" opdrachten bouwt die op een gegeven moment naar de GPU worden verzonden. De methoden op GPUBuffer zijn daarentegen "niet-gebufferd", wat betekent dat ze atomair worden uitgevoerd op het moment dat ze worden aangeroepen.

Zodra u de GPU-opdrachtencoder hebt, roept u copyEncoder.copyBufferToBuffer() aan, zoals hieronder weergegeven, om deze opdracht toe te voegen aan de opdrachtenwachtrij voor latere uitvoering. Voltooi ten slotte de coderingsopdrachten door copyEncoder.finish() aan te roepen en deze naar de opdrachtenwachtrij van het GPU-apparaat te verzenden. De wachtrij is verantwoordelijk voor het afhandelen van inzendingen via device.queue.submit() met de GPU-opdrachten als argumenten. Hierdoor worden alle opdrachten die in de array zijn opgeslagen atomair op volgorde uitgevoerd.

// 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]);

Op dit punt zijn GPU-wachtrijopdrachten verzonden, maar niet noodzakelijkerwijs uitgevoerd. Als u de tweede GPU-buffer wilt lezen, roept u gpuReadBuffer.mapAsync() aan met GPUMapMode.READ . Het retourneert een belofte die zal worden opgelost wanneer de GPU-buffer in kaart wordt gebracht. Haal vervolgens het toegewezen bereik op met gpuReadBuffer.getMappedRange() dat dezelfde waarden bevat als de eerste GPU-buffer zodra alle GPU-opdrachten in de wachtrij zijn uitgevoerd.

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

U kunt dit voorbeeld uitproberen .

Kortom, dit is wat u moet onthouden met betrekking tot buffergeheugenbewerkingen:

• GPU-buffers moeten worden ontkoppeld om te kunnen worden gebruikt bij het indienen van apparaatwachtrijen.
• Wanneer ze in kaart zijn gebracht, kunnen GPU-buffers worden gelezen en geschreven in JavaScript.
• GPU-buffers worden toegewezen wanneer mapAsync() en createBuffer() met mappedAtCreation ingesteld op true worden aangeroepen.

Shader-programmering

Programma's die op de GPU draaien en die alleen berekeningen uitvoeren (en geen driehoeken tekenen), worden compute shaders genoemd. Ze worden parallel uitgevoerd door honderden GPU-kernen (die kleiner zijn dan CPU-kernen) die samenwerken om gegevens te verwerken. Hun invoer en uitvoer zijn buffers in WebGPU.

Om het gebruik van compute shaders in WebGPU te illustreren, spelen we met matrixvermenigvuldiging, een veelgebruikt algoritme in machine learning dat hieronder wordt geïllustreerd.

Matrixvermenigvuldigingsdiagram

Kortom, dit is wat we gaan doen:

1. Maak drie GPU-buffers (twee voor de te vermenigvuldigen matrices en één voor de resultaatmatrix)
2. Beschrijf de invoer en uitvoer voor de compute shader
3. Compute shader-code compileren
4. Zet een rekenpijplijn op
5. Verzend de gecodeerde opdrachten in batch naar de GPU
6. Lees de resultaatmatrix GPU-buffer

GPU-buffers maken

Omwille van de eenvoud worden matrices weergegeven als een lijst met getallen met drijvende komma. Het eerste element is het aantal rijen, het tweede element het aantal kolommen en de rest zijn de werkelijke getallen van de matrix.

Eenvoudige weergave van een matrix in JavaScript en het equivalent daarvan in wiskundige notatie

De drie GPU-buffers zijn opslagbuffers omdat we gegevens in de compute-shader moeten opslaan en ophalen. Dit verklaart waarom de GPU-buffergebruiksvlaggen GPUBufferUsage.STORAGE voor allemaal bevatten. De resultaatmatrixgebruiksvlag heeft ook GPUBufferUsage.COPY_SRC omdat deze naar een andere buffer wordt gekopieerd om te worden gelezen zodra alle GPU-wachtrijopdrachten allemaal zijn uitgevoerd.

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
});

Bind groepsindeling en bind groep

De concepten van de indeling van de bindgroep en de bindgroep zijn specifiek voor WebGPU. Een bindgroepindeling definieert de invoer/uitvoerinterface die door een shader wordt verwacht, terwijl een bindgroep de daadwerkelijke invoer/uitvoergegevens voor een shader vertegenwoordigt.

In het onderstaande voorbeeld verwacht de indeling van de bindgroep twee alleen-lezen opslagbuffers bij genummerde invoerbindingen 0 , 1 en een opslagbuffer bij 2 voor de compute-shader. De bindgroep daarentegen, gedefinieerd voor deze bindgroepindeling, koppelt GPU-buffers aan de items: gpuBufferFirstMatrix aan de binding 0 , gpuBufferSecondMatrix aan de binding 1 en resultMatrixBuffer aan de 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: {
        buffer: gpuBufferFirstMatrix
      }
    },
    {
      binding: 1,
      resource: {
        buffer: gpuBufferSecondMatrix
      }
    },
    {
      binding: 2,
      resource: {
        buffer: resultMatrixBuffer
      }
    }
  ]
});

Bereken shadercode

De compute shader-code voor het vermenigvuldigen van matrices is geschreven in WGSL , de WebGPU Shader Language, die triviaal vertaalbaar is naar SPIR-V . Zonder in detail te treden, zou u hieronder de drie opslagbuffers moeten vinden die zijn geïdentificeerd met var<storage> . Het programma gebruikt firstMatrix en secondMatrix als invoer en resultMatrix als uitvoer.

Merk op dat voor elke opslagbuffer een binding wordt gebruikt die overeenkomt met dezelfde index die is gedefinieerd in de bindgroeplay-outs en bindgroepen die hierboven zijn gedeclareerd.

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;
    }
  `
});

Pijpleiding instellen

De rekenpijplijn is het object dat feitelijk de rekenbewerking beschrijft die we gaan uitvoeren. Maak het door device.createComputePipeline() aan te roepen. Er zijn twee argumenten nodig: de indeling van de bindgroep die we eerder hebben gemaakt, en een rekenfase die het ingangspunt definieert van onze compute shader (de main WGSL-functie) en de daadwerkelijke compute shader-module die is gemaakt met device.createShaderModule() .

const computePipeline = device.createComputePipeline({
  layout: device.createPipelineLayout({
    bindGroupLayouts: [bindGroupLayout]
  }),
  compute: {
    module: shaderModule,
    entryPoint: "main"
  }
});

Opdrachten indienen

Na het instantiëren van een bindgroep met onze drie GPU-buffers en een rekenpijplijn met een bindgroepindeling, is het tijd om ze te gebruiken.

Laten we een programmeerbare compute pass-encoder starten met commandEncoder.beginComputePass() . We zullen dit gebruiken om GPU-opdrachten te coderen die de matrixvermenigvuldiging uitvoeren. Stel de pijplijn in met passEncoder.setPipeline(computePipeline) en de bindgroep ervan op index 0 met passEncoder.setBindGroup(0, bindGroup) . De index 0 komt overeen met de group(0) -decoratie in de WGSL-code.

Laten we het nu hebben over hoe deze compute shader op de GPU gaat draaien. Ons doel is om dit programma stap voor stap parallel uit te voeren voor elke cel van de resultaatmatrix. Voor een resultaatmatrix van grootte 16 bij 32 bijvoorbeeld, om de uitvoeringsopdracht te coderen, op een @workgroup_size(8, 8) , zouden we passEncoder.dispatchWorkgroups(2, 4) of passEncoder.dispatchWorkgroups(16 / 8, 32 / 8) . Het eerste argument "x" is de eerste dimensie, het tweede "y" is de tweede dimensie en het laatste argument "z" is de derde dimensie die standaard 1 is, omdat we deze hier niet nodig hebben. In de GPU-computerwereld wordt het coderen van een opdracht om een ​​kernelfunctie op een set gegevens uit te voeren, dispatching genoemd.

Parallelle uitvoering voor elke resultaatmatrixcel

De grootte van het werkgroepraster voor onze computershader is (8, 8) in onze WGSL-code. Daarom worden "x" en "y", die respectievelijk het aantal rijen van de eerste matrix en het aantal kolommen van de tweede matrix zijn, gedeeld door 8. Daarmee kunnen we nu een rekenaanroep verzenden met passEncoder.dispatchWorkgroups(firstMatrix[0] / 8, secondMatrix[1] / 8) . Het aantal werkgroeprasters dat moet worden uitgevoerd, zijn de dispatchWorkgroups() -argumenten.

Zoals te zien is in de bovenstaande tekening, heeft elke shader toegang tot een uniek builtin(global_invocation_id) object dat zal worden gebruikt om te weten welke resultaatmatrixcel moet worden berekend.

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

Als u de compute pass-encoder wilt beëindigen, roept u passEncoder.end() aan. Maak vervolgens een GPU-buffer die u kunt gebruiken als bestemming voor het kopiëren van de resultaatmatrixbuffer met copyBufferToBuffer . Voltooi ten slotte de coderingsopdrachten met copyEncoder.finish() en verzend deze naar de GPU-apparaatwachtrij door device.queue.submit() aan te roepen met de GPU-opdrachten.

// 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]);

Resultatenmatrix lezen

Het lezen van de resultaatmatrix is ​​net zo eenvoudig als het aanroepen van gpuReadBuffer.mapAsync() met GPUMapMode.READ en wachten op de terugkerende belofte die aangeeft dat de GPU-buffer nu in kaart is gebracht. Op dit punt is het mogelijk om het toegewezen bereik op te halen met gpuReadBuffer.getMappedRange() .

Resultaat van matrixvermenigvuldiging

In onze code is het resultaat dat in de DevTools JavaScript-console is vastgelegd "2, 2, 50, 60, 114, 140".

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

Gefeliciteerd! Je hebt het gehaald. Je kunt met het voorbeeld spelen .

Nog een laatste truc

Een manier om uw code leesbaarder te maken, is door de handige getBindGroupLayout methode van de rekenpijplijn te gebruiken om de bindgroepindeling af te leiden uit de shader-module . Met deze truc hoeft u geen aangepaste indeling van de bindingsgroep te maken en een pijplijnindeling in uw rekenpijplijn op te geven, zoals u hieronder kunt zien.

Er is een illustratie van getBindGroupLayout voor het vorige voorbeeld beschikbaar .

 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: [

Prestatiebevindingen

Dus hoe verhoudt het uitvoeren van matrixvermenigvuldiging op een GPU zich tot het uitvoeren ervan op een CPU? Om daar achter te komen, heb ik het zojuist beschreven programma voor een CPU geschreven. En zoals je in de onderstaande grafiek kunt zien, lijkt het gebruik van de volledige kracht van GPU een voor de hand liggende keuze als de grootte van de matrices groter is dan 256 bij 256.

GPU versus CPU-benchmark

Dit artikel was nog maar het begin van mijn reis door WebGPU te verkennen . Verwacht binnenkort meer artikelen met meer diepgaande informatie over GPU Compute en over hoe weergave (canvas, textuur, sampler) werkt in WebGPU.