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.
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()
ycreateBuffer()
con Se llama amappedAtCreation
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](https://developer.chrome.com/static/docs/capabilities/web-apis/gpu-compute/image/matrix-multiplication-dia-18ebcce8f6445.jpeg?authuser=19&hl=es)
En resumen, esto es lo que haremos:
- Crear tres búferes de GPU (dos para que se multipliquen las matrices y uno para la matriz de resultados)
- Describe las entradas y salidas para el sombreador de cómputos
- Compila el código del sombreador de cómputos
- Configura una canalización de cómputos
- Envía por lotes los comandos codificados a la GPU
- 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](https://developer.chrome.com/static/docs/capabilities/web-apis/gpu-compute/image/simple-representation-a-a54c48f6c0ddd.jpeg?authuser=19&hl=es)
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](https://developer.chrome.com/static/docs/capabilities/web-apis/gpu-compute/image/execution-parallel-each-6e443814da856.jpeg?authuser=19&hl=es)
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](https://developer.chrome.com/static/docs/capabilities/web-apis/gpu-compute/image/matrix-multiplication-res-c04b7fceb1ce8.jpeg?authuser=19&hl=es)
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](https://developer.chrome.com/static/docs/capabilities/web-apis/gpu-compute/image/gpu-vs-cpu-benchmark-391d7651d455d.jpeg?authuser=19&hl=es)
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.