Elegir el tamaño de workgroup con criterio

Casi todo el código de ejemplo de WebGPU que circula usa @workgroup_size(64) y casi ninguno explica por qué. No es una convención estética ni un número redondo: es el menor valor que no desperdicia carriles de ejecución en ninguna de las arquitecturas de GPU que hay en el mercado, y a la vez el mayor que deja al planificador libertad para llenar el chip. Entender de dónde sale ese número es entender cómo ejecuta una GPU, y a partir de ahí sabrás cuándo apartarte de él.

El ancho SIMD es la unidad real de ejecución

Una GPU no ejecuta invocaciones sueltas. Agrupa un número fijo de ellas y las ejecuta en bloqueo de paso: un único contador de programa, una única instrucción por ciclo, y tantos conjuntos de registros como invocaciones en el grupo. NVIDIA lo llama warp y mide 32. AMD lo llama wave: las arquitecturas GCN usan wave64 y las RDNA pueden usar wave32 o wave64 según lo que decida el compilador. Apple lo llama grupo SIMD y mide 32. Intel lo llama subgrupo y su ancho lo elige el compilador entre 8, 16 y 32. Qualcomm Adreno usa waves de 64 o 128 según la generación.

Ese grupo es indivisible. Si tu workgroup tiene 40 invocaciones y el hardware ejecuta en warps de 32, el planificador asigna dos warps: el primero con sus 32 carriles ocupados y el segundo con 8 ocupados y 24 apagados. Esos 24 carriles no se aprovechan para nada; no se les puede dar trabajo de otro workgroup, porque el warp entero pertenece a un solo workgroup. Has pagado 64 carriles de hardware para hacer el trabajo de 40. La eficiencia es del 62,5 por ciento antes de que tu shader ejecute su primera instrucción.

La cuenta general es directa. Con un ancho SIMD w y un tamaño de workgroup n, el número de grupos SIMD es el techo de n / w, los carriles reservados son ese techo por w, y la fracción útil es n dividido entre esos carriles.

Ahí está la primera razón para el 64: es el menor número que da 100 por ciento en las dos columnas a la vez. Con 32 desperdicias la mitad de una wave64; con cualquier valor que no sea múltiplo de 64 desperdicias en alguna de las dos familias. Y como el mismo shader se ejecuta en el portátil con NVIDIA del desarrollador y en el móvil con Adreno del usuario, el valor que hay que elegir es el que no es malo en ninguna parte.

El otro extremo: la ocupación

Si múltiplos de 64 son buenos, ¿por qué no usar el máximo, 256, siempre? Porque un workgroup entero se asigna a una única unidad de cómputo y no se libera hasta que todas sus invocaciones terminan, y eso tiene dos costes.

El primero es de recursos. Cada unidad de cómputo tiene un banco de registros y una memoria compartida de tamaño fijo. Un workgroup consume registros proporcionalmente a su tamaño y consume la memoria workgroup que declares. Cuantos más recursos consuma un grupo, menos grupos caben a la vez en la misma unidad. Y el número de grupos concurrentes es lo que la GPU usa para ocultar la latencia: cuando un warp se bloquea esperando una lectura de memoria —cientos de ciclos— el planificador cambia a otro warp que tenga trabajo. Si solo hay un workgroup residente, no hay a quién cambiar y la unidad se queda parada.

El segundo es de granularidad. Con un dispatch de 5000 elementos y grupos de 256, salen 20 workgroups. Si la GPU tiene 30 unidades de cómputo, diez se quedan sin trabajo. Con grupos de 64 salen 79 workgroups y todas las unidades tienen algo que hacer, y además la cola se vacía de forma más pareja al final. Este efecto —la última tanda incompleta— se llama cola de dispatch y es la causa de que subir el tamaño del workgroup a veces empeore el tiempo total en dominios pequeños.

El compromiso, resumido: grupos pequeños llenan mejor el chip; grupos grandes comparten más datos. Si tu kernel no usa memoria compartida ni barreras, no hay ninguna ventaja en agrandar el grupo, y sí una desventaja. Si tu kernel carga datos en memoria compartida para reutilizarlos, cada duplicación del tamaño reduce a la mitad el número de barreras y aumenta la reutilización, y ahí sí compensa.

El criterio operativo

Para un kernel elemento a elemento sin memoria compartida —escalar un array, aplicar una función, mezclar dos texturas— usa 64. No hay nada que ganar subiendo y sí ocupación que perder.

Para un kernel de reducción o scan dentro del workgroup, usa 256. Cada invocación acaba haciendo log2(256) = 8 pasos con barrera en vez de 6, pero el número de resultados parciales que hay que consolidar en un segundo dispatch se divide por cuatro, y ese segundo dispatch es puro sobrecoste. La cuenta sale a favor de 256 en cuanto el array tiene más de unas decenas de miles de elementos.

Para un kernel con tiles en memoria compartida —multiplicación de matrices, convolución con halo, estarcido— el tamaño lo dicta la geometría del tile, no el ancho SIMD, y sale casi siempre 16 por 16, que son 256. Ahí el tamaño no es negociable porque el tile y el workgroup son la misma cosa.

Para un kernel de imagen usa 8 por 8 o 16 por 16. Ocho por ocho son 64 con forma cuadrada, que es lo mejor de los dos mundos.

// Kernel elemento a elemento: 64, una dimension, sin memoria compartida.
@compute @workgroup_size(64)
fn escalar(@builtin(global_invocation_id) gid: vec3u) { /* ... */ }

// Kernel de reduccion: 256, para minimizar el numero de parciales.
@compute @workgroup_size(256)
fn reducir(@builtin(local_invocation_index) li: u32) { /* ... */ }

// Kernel de imagen: 8x8, cuadrado para minimizar el halo.
@compute @workgroup_size(8, 8)
fn desenfocar(@builtin(global_invocation_id) gid: vec3u) { /* ... */ }

// Kernel con tile: el tile manda, 16x16 = 256.
@compute @workgroup_size(16, 16)
fn gemm(@builtin(local_invocation_id) lid: vec3u) { /* ... */ }

Un apunte sobre la extensión de subgrupos, disponible en Chromium desde principios de 2025 tras su periodo de prueba: cuando está activa, el shader puede leer subgroup_size y usar operaciones colectivas como subgroupAdd. Es tentador ajustar el tamaño del workgroup al ancho SIMD real del dispositivo, y funciona, pero conviene recordar dos cosas. Una, que subgroup_size es un valor de ejecución y no puede alimentar a @workgroup_size, que es constante de pipeline; lo que se puede leer es subgroupMinSize y subgroupMaxSize de la información del adaptador, antes de crear el pipeline. Dos, que la extensión no está en los tres motores, así que el camino sin subgrupos tiene que seguir existiendo.