Introdução à computação de GPU na Web

Esta postagem explora a API WebGPU experimental usando exemplos e ajuda você vai começar a realizar cálculos paralelos de dados usando a GPU.

François Beaufort
François Beaufort
GitHub

Contexto

Como você já deve saber, a Unidade de Processamento Gráfico (GPU) é um sistema em um computador que era originalmente especializado para processamento gráficos. No entanto, nos últimos 10 anos, evoluiu para um modelo que permite aos desenvolvedores implementar muitos tipos de algoritmos, não apenas renderize gráficos 3D, aproveitando a arquitetura exclusiva da GPU. Esses recursos são chamados de computação de GPU e usam uma GPU como coprocessador para computação científica de uso geral é chamado de programação com GPU (GPGPU).

A computação GPU contribuiu significativamente para a recente expansão do machine learning, já que as redes neurais de convolução e outros modelos podem aproveitar a mais eficiente em GPUs. Com a plataforma da Web atual sem recursos de computação da GPU, a "GPU para a Web" do W3C Grupo da comunidade está desenvolvendo uma API para expor as APIs de GPU modernas disponíveis na maioria dispositivos atuais. Essa API é chamada de WebGPU.

A WebGPU é uma API de baixo nível, como a WebGL. Ele é muito poderoso e bastante detalhado, em mais detalhes. Mas tudo bem. O que buscamos é desempenho.

Neste artigo, vou me concentrar na parte de computação de GPU da WebGPU e, como Estou passando o básico para você começar a tocar no seu por conta própria. Vou me aprofundar e abordar a renderização da WebGPU (tela, textura, etc.) nos próximos artigos.

Acessar a GPU

É fácil acessar a GPU na WebGPU. Ligando para navigator.gpu.requestAdapter() retorna uma promessa de JavaScript que será resolvida de forma assíncrona com uma GPU por um adaptador. Pense nesse adaptador como uma placa de vídeo. Ele pode ser integrado (no mesmo chip que a CPU) ou discreta (geralmente uma placa PCIe com tem um bom desempenho, mas usa mais energia).

Quando você tiver o adaptador de GPU, chame adapter.requestDevice() para receber uma promessa que será resolvido com um dispositivo GPU que você usará para calcular a GPU.

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

Ambas as funções têm opções que permitem que você seja específico sobre o tipo de adaptador (preferência de energia) e dispositivo (extensões, limites) desejados. Para o Para simplificar, usaremos as opções padrão neste artigo.

Gravar memória do buffer

Vamos ver como usar o JavaScript para gravar dados na memória para a GPU. Isso processo não é simples por causa do modelo de sandbox usado nos modelos modernos da navegadores da Web.

O exemplo abaixo mostra como gravar quatro bytes no buffer de memória acessível da GPU. Ele chama device.createBuffer(), que assume o tamanho do e o uso dele. Mesmo que a sinalização de uso GPUBufferUsage.MAP_WRITE necessária para esta chamada específica, vamos deixar claro que queremos escrever para esse buffer. Isso resulta em um objeto de buffer de GPU mapeado na criação, graças à mappedAtCreation definido como verdadeiro. Depois, o buffer de dados binários brutos associados ser recuperados chamando o método de buffer da GPU getMappedRange().

A gravação de bytes é familiar se você já jogou com ArrayBuffer. use uma TypedArray e copie os valores nele.

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

Nesse ponto, o buffer da GPU é mapeado, o que significa que ele pertence à CPU. ela é acessível em leitura/gravação do JavaScript. Para que a GPU possa acessá-lo, precisa ser desmapeado, o que é tão simples quanto chamar gpuBuffer.unmap().

O conceito de mapeado/não mapeado é necessário para evitar disputas em que a GPU e a memória de acesso da CPU ao mesmo tempo.

Ler memória do buffer

Agora vamos ver como copiar e ler um buffer de GPU em outro.

Como estamos escrevendo no primeiro buffer da GPU e queremos copiá-lo para um segundo buffer da GPU, é necessária uma nova sinalização de uso GPUBufferUsage.COPY_SRC. A segunda O buffer de GPU é criado em um estado não mapeado, desta vez com device.createBuffer(): A flag de uso é GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ, já que será usada como destino da primeira GPU. e ler em JavaScript depois que os comandos de cópia da GPU forem executados.

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

Como a GPU é um coprocessador independente, todos os comandos da GPU são executados de forma assíncrona. É por isso que há uma lista de comandos de GPU criados e enviados lotes quando necessário. Na WebGPU, o codificador de comandos da GPU retornado pela device.createCommandEncoder() é o objeto JavaScript que cria um lote de “armazenado em buffer” comandos que serão enviados à GPU em algum momento. Os métodos no GPUBuffer, por outro lado, "sem buffer", o que significa que são executados atomicamente no momento em que são chamados.

Depois de ter o codificador de comando da GPU, chame copyEncoder.copyBufferToBuffer() conforme mostrado abaixo, para adicionar esse comando à fila para execução posterior. Por fim, conclua os comandos de codificação chamando copyEncoder.finish() e envie para a fila de comandos do dispositivo da GPU. A fila é responsável por envios feitos via device.queue.submit() com os comandos da GPU como argumentos. Isso executará atomicamente todos os comandos armazenados na matriz em ordem.

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

Neste ponto, os comandos da fila da GPU já foram enviados, mas não necessariamente executados. Para ler o segundo buffer de GPU, chame gpuReadBuffer.mapAsync() com GPUMapMode.READ. Ele retorna uma promessa que será resolvida quando o buffer da GPU for mapeadas. Em seguida, acesse o intervalo mapeado com gpuReadBuffer.getMappedRange() que contém os mesmos valores do primeiro buffer de GPU depois que todos os comandos de GPU na fila foram executados.

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

Você pode testar esta amostra.

Resumindo, veja o que é preciso lembrar sobre operações de memória de buffer:

  • Os buffers de GPU não podem ser mapeados para serem usados no envio da fila do dispositivo.
  • Quando mapeados, os buffers de GPU podem ser lidos e gravados em JavaScript.
  • Os buffers de GPU são mapeados quando mapAsync() e createBuffer() com mappedAtCreation definido como verdadeiro são chamados.

Programação de sombreador

Programas em execução na GPU que realizam apenas cálculos (e não desenham) triângulos) são chamados de sombreadores de computação. Elas são executadas em paralelo por centenas de núcleos de GPU (menores que os núcleos de CPU) que operam juntos para processar dados. A entrada e a saída são buffers na WebGPU.

Para ilustrar o uso de sombreadores de computação na WebGPU, vamos praticar com o modelo multiplicação, um algoritmo comum no machine learning ilustrado abaixo.

Diagrama de multiplicação de matrizes
Diagrama de multiplicação de matrizes

Em resumo, vamos fazer o seguinte:

  1. Crie três buffers de GPU (dois para as matrizes multiplicarem e um para matriz de resultados)
  2. Descrever a entrada e a saída do sombreador de computação
  3. Compilar o código do sombreador de computação
  4. Configurar um pipeline de computação
  5. Enviar em lote os comandos codificados à GPU
  6. Ler o buffer de GPU da matriz de resultado

Criação de buffers de GPU

Para simplificar, as matrizes serão representadas como uma lista de valores números de pontos. O primeiro elemento é o número de linhas; o segundo elemento número de colunas, e o restante são os números reais da matriz.

Representação simples de uma matriz em JavaScript e seu equivalente em notação matemática
Representação simples de uma matriz em JavaScript e seu equivalente em notação matemática

Os três buffers de GPU são de armazenamento, já que precisamos armazenar e recuperar dados o sombreador de computação. Isso explica por que as flags de uso do buffer da GPU incluem GPUBufferUsage.STORAGE para todas elas. A flag de uso da matriz de resultado também tem GPUBufferUsage.COPY_SRC porque ele será copiado para outro buffer para leitura depois que todos os comandos da fila da GPU forem executados.

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

Vincular layout e grupo de vinculação

Os conceitos do layout de grupo de vinculação e do grupo de vinculação são específicos da WebGPU. Uma vinculação layout de grupo define a interface de entrada/saída esperada por um sombreador, enquanto um O grupo de vinculação representa os dados reais de entrada/saída de um sombreador.

No exemplo abaixo, o layout do grupo de vinculação espera dois buffers de armazenamento somente leitura em Vinculações de entrada numeradas 0, 1 e um buffer de armazenamento em 2 para o sombreador de computação. Por outro lado, o grupo de vinculação, definido para esse layout, associa A GPU armazena em buffer às entradas: gpuBufferFirstMatrix para a vinculação 0. gpuBufferSecondMatrix para a vinculação 1, e resultMatrixBuffer para a vinculação 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
      }
    }
  ]
});

Código do sombreador de computação

O código do sombreador de computação para multiplicar matrizes é escrito em WGSL, a Linguagem de sombreador da WebGPU, que pode ser traduzida para SPIR-V. Sem mais detalhadamente, você encontra abaixo dos três buffers de armazenamento identificados com var<storage>. O programa vai usar firstMatrix e secondMatrix como de entrada e resultMatrix como saída.

Observe que cada buffer de armazenamento tem uma decoração binding usada que corresponde o mesmo índice definido nos layouts de grupo de vinculação e nos grupos de vinculação declarados acima.

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

Configuração do pipeline

O pipeline de computação é o objeto que descreve a operação de computação vamos realizar. Crie-a chamando device.createComputePipeline(). São necessários dois argumentos: o layout de grupo de vinculação criado anteriormente e uma fase que define o ponto de entrada do sombreador de computação (a função main da WGSL). e o módulo do shader de computação real criado com device.createShaderModule().

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

Envio de comandos

Depois de instanciar um grupo de vinculação com nossos três buffers de GPU e um pipeline com um layout de grupo de vinculação, é hora de usá-los.

Vamos começar um codificador programável de passagem de computação commandEncoder.beginComputePass(): Ele será usado para codificar comandos da GPU que realizará a multiplicação de matrizes. Defina o pipeline com passEncoder.setPipeline(computePipeline) e o grupo de vinculação no índice 0 com passEncoder.setBindGroup(0, bindGroup). O índice 0 corresponde ao Decoração group(0) no código da WGSL.

Agora, vamos falar sobre como esse sombreador de computação será executado na GPU. Nossos o objetivo é executar esse programa em paralelo para cada célula da matriz de resultado, etapa por etapa. Para uma matriz de resultado de tamanho 16 por 32, por exemplo, para codificar o comando de execução, em um @workgroup_size(8, 8), chamaríamos passEncoder.dispatchWorkgroups(2, 4) ou passEncoder.dispatchWorkgroups(16 / 8, 32 / 8). O primeiro argumento "x" é a primeira dimensão, e a segunda é "y" é a segunda dimensão, e o último "z" é a terceira dimensão cujo padrão é 1, pois não precisamos dela aqui. No mundo da computação da GPU, a codificação de um comando para executar uma função kernel em um conjunto de dados é chamada de despacho.

Execução em paralelo para cada célula da matriz de resultado
Execução em paralelo para cada célula da matriz de resultado

O tamanho da grade do grupo de trabalho para nosso sombreador de computação é (8, 8) na nossa WGSL o código-fonte. Por isso, "x" e "y" que são respectivamente o número de linhas de a primeira matriz e o número de colunas da segunda matriz serão divididos até 8. Com isso, agora podemos enviar uma chamada de computação com passEncoder.dispatchWorkgroups(firstMatrix[0] / 8, secondMatrix[1] / 8): A número de grades de grupo de trabalho a serem executadas são os argumentos dispatchWorkgroups().

Como mostrado no desenho acima, cada shader terá acesso a um nome Objeto builtin(global_invocation_id) que será usado para saber qual resultado célula de matriz a ser calculada.

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

Para encerrar o codificador de passagem de computação, chame passEncoder.end(). Em seguida, crie uma Buffer de GPU a ser usado como destino para copiar o buffer da matriz de resultado com copyBufferToBuffer: Por fim, finalize os comandos de codificação com copyEncoder.finish() e envie-os para a fila de dispositivos de GPU chamando device.queue.submit() pelos comandos da 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]);

Matriz de resultados de leitura

Ler a matriz de resultados é tão fácil quanto chamar gpuReadBuffer.mapAsync() com GPUMapMode.READ e aguardando a resolução da promessa de retorno, o que indica o buffer da GPU está mapeado. Neste ponto, é possível receber o conjunto com gpuReadBuffer.getMappedRange().

Resultado de multiplicação de matrizes
Resultado da multiplicação de matrizes

Em nosso código, o resultado registrado no console JavaScript do DevTools é "2, 2, 50, 60, 114, 140".

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

Parabéns! Você conseguiu. Você pode usar o sample.

Um último truque

Uma forma de facilitar a leitura do código é usar a ferramenta Método getBindGroupLayout do pipeline de computação para inferir o grupo de vinculação o layout do módulo de shader. Esse truque elimina a necessidade de criar o layout do grupo de vinculação personalizado e a especificação de um layout de pipeline na computação pipeline, como mostrado abaixo.

Uma ilustração de getBindGroupLayout para a amostra anterior está disponível.

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

Descobertas de desempenho

Qual é a diferença entre executar a multiplicação de matrizes em uma GPU e executá-la em um CPU? Para descobrir, escrevi o programa que acabamos de descrever para uma CPU. E, da maneira como você como mostra no gráfico abaixo, usar toda a capacidade da GPU parece ser uma escolha óbvia quando o tamanho das matrizes for maior que 256 por 256.

Comparativo de mercado de GPU e CPU
Comparativo de mercado de GPU e CPU

Este artigo foi apenas o começo da minha jornada explorando a WebGPU. Espere mais artigos em breve com mais detalhes sobre computação de GPU e como a renderização (canvas, textura, sampler) funciona na WebGPU.