SIMT: por qué tus invocaciones no son hilos independientes

WGSL se escribe como si cada invocación fuese un hilo con su propio flujo de control. Es una ficción muy útil y muy peligrosa: el hardware no ejecuta hilos independientes, sino bloques fijos de 32 o 64 invocaciones que comparten un único contador de programa. Casi todas las anomalías de rendimiento que verás en un shader se explican por la distancia entre esa ficción y esa realidad.

Una instrucción, muchos datos, un solo contador de programa

El modelo se llama SIMT, single instruction, multiple threads, y está a medio camino entre dos cosas que ya conoces.

No es SIMD clásico, como las instrucciones vectoriales de una CPU. Ahí escribes explícitamente que quieres sumar cuatro floats a la vez y el vector es visible en el código. En SIMT escribes código escalar, para una invocación, y el hardware se encarga de ejecutar 32 copias a la vez.

Tampoco es multihilo. Los hilos de una CPU tienen cada uno su contador de programa y pueden estar en puntos completamente distintos del código. Las invocaciones de un mismo subgrupo comparten el contador: en cada ciclo, todas ejecutan la misma instrucción, sobre datos distintos.

El tamaño del subgrupo es una propiedad del hardware y no se elige. En las GPU de NVIDIA son 32 invocaciones y se llama warp. En AMD son 64 en las arquitecturas GCN y configurable entre 32 y 64 en RDNA, y se llama wavefront. En las de Apple son 32. En muchas GPU móviles varía entre 8 y 32 según el modelo. WebGPU expone esa cifra de forma indirecta a través de adapter.info.subgroupMinSize y adapter.info.subgroupMaxSize, y solo garantiza un rango, no un valor.

De ese modelo salen tres consecuencias inmediatas.

Si una invocación tiene que esperar, esperan las 32. No hay forma de que media docena avance mientras el resto se bloquea. La unidad de planificación es el subgrupo entero.

Si el código toma dos caminos distintos, se ejecutan los dos. Es la divergencia, y merece su propia lección.

Las invocaciones vecinas leen memoria a la vez. El hardware combina los accesos de un subgrupo en el menor número posible de transacciones de memoria. Si las 32 leen posiciones contiguas, es una transacción; si leen posiciones dispersas, pueden ser 32. La diferencia de rendimiento entre esos dos casos es de un orden de magnitud.

Workgroup y subgrupo: dos cosas distintas que se confunden

Es la confusión más frecuente del nivel y merece una separación explícita.

El subgrupo es hardware. Fijo, no elegible, entre 8 y 64 invocaciones según el chip. Es la unidad de ejecución.

El workgroup es API. Lo declaras tú en el shader con @workgroup_size(x, y, z), tiene un máximo de 256 invocaciones por defecto en WebGPU, y es la unidad de cooperación: dentro de un workgroup hay memoria compartida y barreras; entre workgroups, nada.

Un workgroup se ejecuta como varios subgrupos consecutivos sobre la misma unidad de cómputo. Un workgroup de 64 en una GPU con subgrupos de 32 son dos subgrupos. Uno de 256 son ocho.

De ahí sale la primera regla práctica de este nivel, y es una regla que se puede aplicar sin entender nada más: declara siempre el tamaño de workgroup como múltiplo de 64. Con 64 encajas exactamente en un subgrupo de AMD y en dos de NVIDIA o Apple. Con un tamaño de, digamos, 100, el hardware redondea hacia arriba al múltiplo del subgrupo: reserva cuatro subgrupos de 32, o sea 128 invocaciones, de las cuales 28 no hacen absolutamente nada pero ocupan registros y ranuras de planificación. Has tirado el 22% de la capacidad por escribir un número redondo en base diez.

// Mal: 100 no es multiplo de 32 ni de 64.
@compute @workgroup_size(100)
fn peor(@builtin(global_invocation_id) gid: vec3<u32>) { /* ... */ }

// Bien: 64 encaja en cualquier tamano de subgrupo real.
@compute @workgroup_size(64)
fn mejor(@builtin(global_invocation_id) gid: vec3<u32>) { /* ... */ }

// Bien para trabajo bidimensional: 8 por 8 son 64.
@compute @workgroup_size(8, 8)
fn imagen(@builtin(global_invocation_id) gid: vec3<u32>) { /* ... */ }

La segunda regla es sobre el reparto. Como dispatchWorkgroups lanza grupos enteros, casi nunca cubre exactamente tu problema:

const N = 100000;
const TAM = 64;
pase.dispatchWorkgroups(Math.ceil(N / TAM));  // 1563 grupos, 100032 invocaciones

Sobran 32 invocaciones, y esas 32 van a indexar fuera del array si no las paras. La guarda es obligatoria y va siempre en la primera línea del kernel:

@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
  let i = gid.x;
  if (i >= arrayLength(&datos)) { return; }
  // ...
}

Ese if diverge, sí, pero solo en el último workgroup del dispatch, y ahí no hay nada que hacer. Es divergencia aceptable.

Dónde se nota el lockstep en la práctica

Cuatro sitios concretos, para que el modelo deje de ser abstracto.

El coste de un if no es el que parece. Un condicional donde todas las invocaciones del subgrupo toman la misma rama es casi gratis: el hardware salta y ya está. Un condicional donde se dividen ejecuta las dos ramas. La diferencia entre esos dos casos no está en el código, está en los datos.

Un bucle dura lo que la iteración más larga del subgrupo. Si escribes while (i < n) y n vale 3 para 31 invocaciones y 500 para una, el subgrupo entero itera 500 veces. Las 31 están enmascaradas y no producen resultado, pero ocupan la máquina.

Las derivadas de los fragment shaders dependen del lockstep. Funciones como dpdx y dpdy calculan la diferencia entre píxeles vecinos, y solo pueden hacerlo porque esos píxeles se ejecutan a la vez en el mismo subgrupo, en bloques de 2 por 2. Por eso textureSample —que necesita las derivadas para elegir el nivel de mipmap— tiene que llamarse en control de flujo uniforme: si unas invocaciones del bloque han salido por una rama distinta, las derivadas no están definidas. WGSL rechaza en tiempo de compilación las llamadas a esas funciones dentro de flujo no uniforme, y el error de «uniformity» que verás alguna vez viene exactamente de aquí.

El rasterizador siempre produce cuadrados de 2 por 2. Un triángulo que cubre un solo píxel activa cuatro invocaciones de fragmento; tres de ellas se descartan al final. Es la razón profunda de que la geometría muy pequeña sea desproporcionadamente cara y de que valga la pena usar niveles de detalle: un modelo de diez mil triángulos que ocupa cien píxeles en pantalla puede ejecutar más invocaciones de fragmento que píxeles tiene la pantalla entera.

Con el lockstep entendido, el siguiente paso es ver exactamente qué pasa cuando las invocaciones de un subgrupo dejan de estar de acuerdo: la divergencia.