Comienza a usar GPU Compute en la Web

En esta publicación, se explora la API experimental de WebGPU con ejemplos y se ayuda comienzas a realizar cálculos de datos paralelos con la GPU.

François Beaufort
François Beaufort
GitHub

Información general

Como ya sabrás, la unidad de procesamiento gráfico (GPU) es una unidad en una computadora que originalmente se especializaba en procesar gráficos. Sin embargo, en los últimos 10 años, evolucionó que permite a los desarrolladores implementar muchos tipos de algoritmos, no solo representar gráficos en 3D y, al mismo tiempo, aprovechar la arquitectura única de la GPU: Estas capacidades se conocen como procesamiento de GPU, y usar una GPU como coprocesador para la computación científica de uso general se denomina Programación de GPU (GPGPU).

El procesamiento de GPU contribuyó significativamente al reciente auge del aprendizaje automático, ya que las redes neuronales de convolución y otros modelos pueden aprovechar arquitectura para ejecutarse de manera más eficiente en las GPU. Con la plataforma web actual sin capacidades de procesamiento de GPU, la "GPU para la Web" de W3C Grupo de la comunidad está diseñando una API para exponer las APIs de GPU modernas que están disponibles en la mayoría de dispositivos actuales. Esta API se llama WebGPU.

WebGPU es una API de bajo nivel, como WebGL. Es muy potente y bastante detallado, ya que en la nube. Pero está bien. Lo que buscamos es el rendimiento.

En este artículo, nos enfocaremos en la parte de procesamiento de GPU de WebGPU y, para Honestamente, te escribo desde la superficie para que puedas empezar a jugar en tu por sí solas. Profundizaremos más y trataremos la renderización de WebGPU (lienzo, textura, etc.) en los próximos artículos.

Accede a la GPU

Es fácil acceder a la GPU en WebGPU. Llamando a navigator.gpu.requestAdapter() Devuelve una promesa de JavaScript que se resolverá de forma asíncrona con una GPU. adaptador. Considera este adaptador como la tarjeta gráfica. Puede integrarse (en el mismo chip que la CPU) o discreta (por lo general, una tarjeta PCIe que es más tiene un buen rendimiento, pero usa más energía).

Una vez que tengas el adaptador de GPU, llama a adapter.requestDevice() para obtener una promesa que se resolverá con el dispositivo GPU que usarás para hacer cálculos de GPU.

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

Ambas funciones tienen opciones que te permiten ser específico sobre el tipo de adaptador (preferencia de alimentación) y el dispositivo (extensiones, límites) que desees. Para el por cuestiones de simplicidad, usaremos las opciones predeterminadas en este artículo.

Cómo escribir memoria de búfer

Veamos cómo usar JavaScript para escribir datos en la memoria para la GPU. Esta no es sencillo debido al modelo de zona de pruebas que se usa en la Web moderna navegadores.

En el siguiente ejemplo, se muestra cómo escribir cuatro bytes en la memoria de búfer accesible desde la GPU. Llama a device.createBuffer(), que toma el tamaño de la búfer y su uso. Si bien la marca de uso GPUBufferUsage.MAP_WRITE es no se requiere para esta llamada específica, seamos explícitos que queremos escribir a este búfer. El resultado es un objeto del búfer de la GPU asignado en la creación gracias a Se estableció mappedAtCreation como verdadero. Luego, el búfer de datos binarios sin procesar asociado se puede se puede recuperar llamando al método de búfer de la GPU getMappedRange().

Escribir bytes es similar si ya jugaste con ArrayBuffer; usa un TypedArray y copia los valores allí.

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

En este punto, se asigna el búfer de la GPU, lo que significa que es propiedad de la CPU. sean accesibles en lectura y escritura desde JavaScript. Para que la GPU pueda acceder a ellos, no debe estar mapeada, lo cual es tan sencillo como llamar a gpuBuffer.unmap().

El concepto de asignado/no asignado es necesario para evitar condiciones de carrera en las que la GPU y la memoria de acceso a la CPU al mismo tiempo.

Leer la memoria del búfer

Ahora, veamos cómo copiar un búfer de la GPU a otro y volver a leerlo.

Como escribimos en el primer búfer de la GPU y queremos copiarlo en una segunda Búfer de GPU, se requiere una nueva marca de uso GPUBufferUsage.COPY_SRC. El segundo El búfer de la GPU se crea en un estado sin asignar esta vez con device.createBuffer() Su marca de uso es GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ, ya que se usará como destino de la primera GPU. almacenar en búfer y leer en JavaScript una vez que se hayan ejecutado los comandos de copia de la 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
});

Como la GPU es un coprocesador independiente, todos los comandos de la GPU se ejecutan de forma asíncrona. Por eso hay una lista de comandos GPU compilados y enviados en por lotes cuando sea necesario. En WebGPU, el codificador de comandos de GPU que devuelve device.createCommandEncoder() es el objeto de JavaScript que compila un lote de "almacenado en búfer" que se enviarán a la GPU en algún momento. Los métodos de Por otro lado, GPUBuffer no están almacenados en búfer, lo que significa que se ejecutan de forma atómica. en el momento en que se llaman.

Una vez que tengas el codificador de comandos de la GPU, llama a copyEncoder.copyBufferToBuffer(). como se muestra a continuación para agregar este comando a la cola de comandos para ejecutarlo más tarde. Por último, para terminar de codificar los comandos, llama a copyEncoder.finish() y envía a la cola de comandos del dispositivo GPU. La cola es responsable de administrar envíos realizados mediante device.queue.submit() con los comandos de la GPU como argumentos Esto ejecutará atómicamente todos los comandos almacenados en el array en orden.

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

En este punto, se enviaron los comandos de la cola de GPU, pero no necesariamente se ejecutaron. Para leer el segundo búfer de la GPU, llama a gpuReadBuffer.mapAsync() con GPUMapMode.READ Devuelve una promesa que se resolverá cuando se agote el búfer de la GPU. asignado. Luego, obtén el rango asignado con gpuReadBuffer.getMappedRange() que contiene los mismos valores que el primer búfer de GPU una vez que todos los comandos de GPU en cola se ejecutaron.

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

Puedes probar esta muestra.

En resumen, esto es lo que debes recordar sobre las operaciones de memoria de búfer:

  • Los búferes de la GPU no deben estar asignados para usarse en el envío de colas de dispositivos.
  • Cuando se asignan, los búferes de la GPU se pueden leer y escribir en JavaScript.
  • Los búferes de la GPU se asignan cuando mapAsync() y createBuffer() con Se llama a mappedAtCreation establecido como verdadero.

Programación de sombreadores

Los programas que se ejecutan en la GPU y que solo realizan cálculos (y no generan dibujos) triángulos) se denominan sombreadores de cómputos. Son ejecutadas en paralelo por cientos de de núcleos de GPU (más pequeños que los núcleos de CPU) que funcionan juntos para procesar de datos no estructurados. Su entrada y salida son búferes en WebGPU.

Para ilustrar el uso de sombreadores de cómputos en WebGPU, jugaremos con los modelos de valores de multiplicación, un algoritmo común en el aprendizaje automático que se ilustra a continuación.

Diagrama de multiplicación de matrices
Diagrama de multiplicación de matrices

En resumen, esto es lo que haremos:

  1. Crear tres búferes de GPU (dos para que se multipliquen las matrices y uno para la matriz de resultados)
  2. Describe las entradas y salidas para el sombreador de cómputos
  3. Compila el código del sombreador de cómputos
  4. Configura una canalización de cómputos
  5. Envía por lotes los comandos codificados a la GPU
  6. Lee el búfer de GPU de la matriz de resultados

Creación de búferes de GPU

Para simplificar, las matrices se representarán como una lista de elementos y números de punto. El primer elemento es el número de filas, el segundo elemento de columnas y el resto son los números reales de la matriz.

Representación simple de una matriz en JavaScript y su equivalente en notación matemática
Representación simple de una matriz en JavaScript y su equivalente en notación matemática

Los tres búferes de la GPU son búferes de almacenamiento, ya que necesitamos almacenar y recuperar datos en el sombreador de cómputos. Esto explica por qué las marcas de uso del búfer de la GPU incluyen GPUBufferUsage.STORAGE para todos ellos. La marca de uso de la matriz de resultados también tiene GPUBufferUsage.COPY_SRC porque se copiará en otro búfer para una vez que se hayan ejecutado todos los comandos de la cola de 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
});

Cómo vincular el diseño y el grupo de vinculaciones

Los conceptos de diseño y grupo de vinculaciones son específicos de WebGPU. Una vinculación de grupo define la interfaz de entrada/salida que espera un sombreador, mientras que un Bind group representa los datos de entrada y salida reales para un sombreador.

En el siguiente ejemplo, el diseño del grupo de vinculaciones espera dos búferes de almacenamiento de solo lectura en vinculaciones de entrada numeradas 0, 1 y un búfer de almacenamiento en 2 para el sombreador de cómputos. Por otro lado, el grupo de vinculaciones, definido para este diseño de grupo, asocia La GPU almacena en búfer las entradas: gpuBufferFirstMatrix a la vinculación 0. gpuBufferSecondMatrix a la vinculación 1 y resultMatrixBuffer a la Vinculación de 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
      }
    }
  ]
});

Calcula el código del sombreador

El código del sombreador de cómputos para multiplicar matrices se escribe en WGSL, la WebGPU Shader Language, que se puede traducir de forma trivial a SPIR-V Sin Al entrar en detalle, encontrarás a continuación los tres búferes de almacenamiento identificados con var<storage>. El programa usará firstMatrix y secondMatrix como y resultMatrix como su salida.

Ten en cuenta que cada búfer de almacenamiento tiene una decoración binding utilizada que corresponde a el mismo índice definido en los diseños de grupos de vinculaciones y los grupos de vinculaciones declarados anteriormente.

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

Configuración de la canalización

La canalización de cómputos es el objeto que describe la operación de cómputos. que realizaremos. Para crearlo, llama a device.createComputePipeline(). Toma dos argumentos: el diseño del grupo de vinculaciones que creamos anteriormente y una instancia de etapa que define el punto de entrada de nuestro sombreador de cómputos (la función WGSL main) y el módulo de sombreador de cómputos real creado con device.createShaderModule().

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

Envío de comandos

Después de crear una instancia de un grupo de vinculaciones con nuestros tres búferes de GPU y una instancia con un diseño de grupos de vinculaciones, es hora de usarlos.

Comencemos un codificador de pases de cómputos programable con commandEncoder.beginComputePass() Lo usaremos para codificar comandos de GPU que realizará la multiplicación de matrices. Configura su canalización con passEncoder.setPipeline(computePipeline) y su grupo de vinculaciones en el índice 0 con passEncoder.setBindGroup(0, bindGroup) El índice 0 corresponde al Decoración group(0) en el código WGSL.

Ahora, hablemos sobre cómo se ejecutará este sombreador de cómputos en la GPU. Nuestro es ejecutar este programa en paralelo para cada celda de la matriz de resultados, paso a paso. Para una matriz de resultados de 16 por 32, por ejemplo, para codificar el comando de ejecución, en un @workgroup_size(8, 8), llamaremos passEncoder.dispatchWorkgroups(2, 4) o passEncoder.dispatchWorkgroups(16 / 8, 32 / 8). El primer argumento "x" es la primera dimensión, la segunda, "y" es la segunda dimensión, y al último, “z” es la tercera dimensión, cuyo valor predeterminado es 1, ya que no la necesitamos aquí. En el mundo de la computación con GPU, codificar un comando para ejecutar una función de kernel en un conjunto de datos se denomina envío.

Ejecución en paralelo para cada celda de la matriz de resultados
Ejecución en paralelo para cada celda de la matriz de resultados

El tamaño de la cuadrícula del grupo de trabajo para nuestro sombreador de cómputos es (8, 8) en nuestro WGSL. código. Por ese motivo, "x" y "y" que son respectivamente el número de filas de se dividirán la primera matriz y el número de columnas de la segunda 8. Con eso, podemos enviar una llamada de procesamiento con passEncoder.dispatchWorkgroups(firstMatrix[0] / 8, secondMatrix[1] / 8) El la cantidad de cuadrículas de grupo de trabajo que se ejecutarán son los argumentos dispatchWorkgroups().

Como se ve en el dibujo anterior, cada sombreador tendrá acceso a un builtin(global_invocation_id) objeto que se usará para saber qué resultado celda de matriz para procesar.

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 finalizar el codificador de pases de cómputos, llama a passEncoder.end(). Luego, crea un Búfer de la GPU que se usará como destino para copiar el búfer de la matriz de resultados con copyBufferToBuffer Por último, termina de codificar los comandos con copyEncoder.finish() y envíalos a la cola de dispositivos de la GPU mediante una llamada device.queue.submit() con los comandos de la 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]);

Leer la matriz de resultados

Leer la matriz de resultados es tan fácil como llamar a gpuReadBuffer.mapAsync() con GPUMapMode.READ y esperando a que se resuelva la promesa que indica se asigna el búfer de la GPU. En este punto, es posible obtener el mapa rango con gpuReadBuffer.getMappedRange().

Resultado de la multiplicación de matrices
Resultado de la multiplicación de matrices

En nuestro código, el resultado registrado en la consola de JavaScript de Herramientas para desarrolladores es “2, 2, 50, 60, 114, 140".

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

¡Felicitaciones! Lo lograste. Puedes jugar con la muestra.

Un último truco

Una forma de facilitar la lectura del código es usar Método getBindGroupLayout de la canalización de cómputos para inferir el grupo de vinculaciones diseño del módulo de sombreador. Este truco elimina la necesidad de crear un diseño de grupos de vinculaciones personalizado y especificación de un diseño de canalización en tu procesamiento de procesamiento, como puedes ver a continuación.

Una ilustración de getBindGroupLayout para el ejemplo anterior está disponible.

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

Hallazgos de rendimiento

Entonces, ¿cómo se compara la ejecución de la multiplicación de matrices en una GPU con la ejecución en una ¿CPU? Para averiguarlo, escribí el programa que acabamos de describir para una CPU. Y, como pueden como se ve en el siguiente gráfico, usar toda la potencia de la GPU parece una opción obvia cuando el tamaño de las matrices es mayor que 256 por 256.

Comparativas de GPU y CPU
Comparativas de GPU y CPU

Este artículo fue solo el comienzo de mi recorrido por explorar WebGPU. Espera más próximamente, en los que se explicarán más detalles sobre el procesamiento con GPU y cómo (lienzo, textura, muestra) funciona en WebGPU.