La ausencia de compute: cómo se hacía GPGPU en WebGL y qué costaba

De todas las diferencias entre WebGL y WebGPU, la que cambia lo que se puede construir no es el rendimiento: es que WebGL no tiene forma de ejecutar un programa arbitrario en la GPU. Todo lo que se ha hecho de cómputo general en la web hasta ahora se hizo disfrazándolo de dibujado, con un conjunto de trucos que funcionaban, que costaban mucho y que tenían límites duros. Ver esos límites de cerca es lo que hace evidente qué significa tener cómputo de verdad.

Calcular dibujando

La técnica se llama ping-pong de texturas y todo el mundo que ha hecho simulación en WebGL la conoce. La idea es aprovechar que un fragment shader es, en el fondo, una función que se ejecuta una vez por píxel y escribe un valor.

El montaje es así. Se crea una textura de punto flotante de N por N y se interpreta como un array de N al cuadrado elementos: cada texel guarda el estado de un elemento, con hasta cuatro números por texel en los canales RGBA. Se crea una segunda textura idéntica. Se dibuja un cuadrilátero que cubre exactamente la textura de destino, con un framebuffer que apunta a ella. El fragment shader recibe implícitamente qué texel está calculando a través de las coordenadas del fragmento, lee el estado anterior de la textura de origen, calcula el nuevo y lo devuelve como color. Al terminar, se intercambian origen y destino y se repite.

// WebGL2: integrar posiciones de particulas guardadas en una textura.
#version 300 es
precision highp float;
uniform sampler2D posiciones;
uniform sampler2D velocidades;
uniform float dt;
out vec4 salida;

void main() {
  ivec2 celda = ivec2(gl_FragCoord.xy);
  vec3 p = texelFetch(posiciones, celda, 0).xyz;
  vec3 v = texelFetch(velocidades, celda, 0).xyz;
  salida = vec4(p + v * dt, 1.0);
}

Funciona. Se han construido simulaciones de fluidos, sistemas de partículas de millones de elementos y hasta redes neuronales sobre esta base. Pero cada línea de ese shader esconde una concesión.

Los tres límites duros

Cada invocación escribe exactamente un texel, en una posición que no elige. El rasterizador decide qué fragmento le toca a quién. Un fragment shader no puede escribir en una dirección calculada. Esto elimina de un plumazo toda una familia de algoritmos: cualquier cosa que necesite dispersar resultados —construir un histograma, insertar elementos en celdas de una rejilla espacial, compactar un array quitando huecos, ordenar— no se puede expresar directamente. Los rodeos existen y son caros: para construir un histograma se dibujan puntos con mezcla aditiva; para compactar se hace una suma de prefijos con varias pasadas de reducción; para ordenar se implementa una red bitónica con una pasada de dibujado por etapa.

No hay memoria compartida ni sincronización entre invocaciones. Los fragmentos se ejecutan de forma independiente y no pueden comunicarse. Cualquier algoritmo que quiera que un grupo de hilos coopere sobre un bloque de datos —que es la técnica central de la programación en GPU— hay que descomponerlo en varias pasadas completas de dibujado, escribiendo el resultado intermedio a una textura entre cada una. Cada pasada es un viaje entero a la memoria principal.

Los datos tienen que caber en el modelo de textura. Un array de estructuras no se guarda: se descompone en varias texturas de cuatro canales, y el código de indexación se vuelve una traducción constante entre índice lineal y coordenada bidimensional. Los tipos enteros son incómodos, los índices grandes chocan con el tamaño máximo de textura, y el empaquetado consume tiempo de desarrollo que no tiene nada que ver con el problema.

Qué añade un compute shader

Un @compute de WGSL tiene cuatro cosas que un fragment shader no puede tener, y cada una elimina uno de los límites anteriores.

Escritura arbitraria en búferes de almacenamiento. Una variable declarada var<storage, read_write> se puede indexar con cualquier expresión y escribir. La invocación decide dónde escribe. Eso hace expresables la dispersión, la compactación, la construcción de estructuras de datos y la ordenación in situ.

Memoria compartida por grupo de trabajo. Una variable var<workgroup> es visible por todas las invocaciones del mismo grupo y vive en memoria dentro del chip, con una latencia de un orden de magnitud menor que la memoria global. Es lo que permite que un grupo de hilos cargue un bloque de datos una vez y lo reutilice muchas veces, que es el fundamento de la multiplicación de matrices eficiente, de las convoluciones y de las reducciones.

Sincronización dentro del grupo. workgroupBarrier() garantiza que todas las invocaciones del grupo han llegado a ese punto y que sus escrituras en memoria compartida son visibles. Sin barrera no hay cooperación posible.

Operaciones atómicas. Un atomic<u32> en memoria de almacenamiento o compartida admite incrementos, intercambios y comparaciones indivisibles. Es lo que hace posible que mil invocaciones cuenten en el mismo contador sin perder cuentas.

Y una quinta que es de forma más que de fondo: la geometría de la ejecución la eliges tú. dispatchWorkgroups(x, y, z) lanza una rejilla tridimensional de grupos, y cada invocación conoce su posición con @builtin(global_invocation_id). No hay que fingir que un problema unidimensional es una imagen.

@group(0) @binding(0) var<storage, read> entrada: array<f32>;
@group(0) @binding(1) var<storage, read_write> salida: array<f32>;
@group(0) @binding(2) var<storage, read_write> total: atomic<u32>;

@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
  let i = gid.x;
  if (i >= arrayLength(&entrada)) { return; }
  let v = entrada[i] * 2.0;
  salida[i] = v;
  if (v > 1.0) { atomicAdd(&total, 1u); }
}

Ese fragmento hace tres cosas que en WebGL requerirían tres pasadas de dibujado distintas y una textura intermedia: leer, escribir en la misma posición, y contar cuántos elementos cumplen una condición. En cómputo es un solo lanzamiento.

Qué pasa a ser rutinario

La lista de lo que cambia de categoría es larga y conviene tenerla presente porque orienta lo que merece la pena aprender.

Simulación de partículas con interacción. No solo mover puntos, que ya se hacía, sino que se vean entre sí: colisiones, atracción, cohesión. Requiere estructuras espaciales, que requieren dispersión.

Fluidos y materiales. Los métodos de partículas requieren buscar vecinos; los de rejilla requieren resolver sistemas iterativos con reducciones globales. Las dos cosas dependen de cooperación entre hilos.

Álgebra lineal. Multiplicar matrices con tiling en memoria compartida es el ejemplo canónico de por qué existe var<workgroup>: reduce los accesos a memoria global por un factor igual al tamaño del bloque.

Ordenación y búsqueda. Un radix sort en GPU necesita histogramas, sumas de prefijos y dispersión. Los tres eran acrobacias.

Renderizado dirigido por GPU. Descartar objetos no visibles en un compute shader y escribir los parámetros de dibujado en un búfer que drawIndirect consume. La CPU deja de saber cuántos objetos se dibujan.

Procesado de imagen serio. Un desenfoque separable con memoria compartida es varias veces más rápido que el mismo desenfoque hecho con dos pasadas de fragment shader, porque cada píxel de la ventana se lee una vez en lugar de N.

Con el diagnóstico completo, queda ordenar la respuesta: qué gana WebGPU exactamente, pieza por pieza.