Ocupación: por qué a veces más hilos van más rápido
Qué limita cuántos grupos de trabajo caben a la vez en una unidad de cómputo, por qué la ocupación esconde la latencia y por qué maximizarla no siempre es la decisión correcta.
La ocupación es la métrica que explica los resultados más contraintuitivos de la programación en GPU: que un shader al que le quitas una variable local vaya un 40% más rápido, que reducir el tamaño de workgroup mejore el rendimiento, y que lanzar el doble de invocaciones cueste lo mismo que lanzar la mitad. Es también la métrica que se malinterpreta con más frecuencia, porque «más ocupación es mejor» es cierto hasta un punto y falso después.
- Definir ocupación y explicar el mecanismo por el que esconde latencia.
- Enumerar los tres recursos que limitan cuántos grupos caben a la vez.
- Calcular la ocupación de un kernel a partir de su consumo de registros y memoria compartida.
- Decidir cuándo conviene sacrificar ocupación a cambio de otra cosa.
Qué es y por qué funciona
Una unidad de cómputo puede tener varios grupos de trabajo residentes a la vez. Residente significa que sus registros están asignados, su memoria compartida reservada y sus invocaciones listas para ejecutarse en cuanto el planificador les dé paso. La ocupación es la proporción entre las invocaciones residentes y el máximo que la unidad admitiría.
El mecanismo por el que eso acelera es el mismo que ya vimos: cuando un subgrupo emite una lectura de memoria global y se queda esperando 400 ciclos, el planificador da paso a otro subgrupo residente. Si hay ocho grupos residentes, hay mucho trabajo con el que rellenar esos 400 ciclos y las ALU no paran. Si hay uno solo, la unidad de cómputo se queda mirando la pared.
Con números redondos: un kernel donde una de cada diez instrucciones es una lectura de memoria global de 400 ciclos necesita del orden de cuarenta subgrupos con trabajo listo para tapar cada espera. Si el hardware permite 64 subgrupos residentes y tu configuración solo deja entrar 8, estás usando una octava parte de la capacidad de ocultación y el kernel va a ir varias veces más lento de lo que podría, con exactamente la misma aritmética.
Los tres límites
Tres recursos finitos deciden cuántos grupos caben. El que se agote primero manda.
El fichero de registros. Es el limitante más frecuente. Cada invocación viva tiene sus registros asignados en exclusiva mientras dura, y el fichero es de tamaño fijo. Si una unidad de cómputo tiene 65536 registros de 32 bits y tu shader usa 64 registros por invocación, caben 1024 invocaciones. Si el compilador necesita 128, caben 512: la ocupación se ha reducido a la mitad por una decisión que ni siquiera está en tu código, sino en cómo el compilador asignó registros.
Qué consume registros: variables locales vivas simultáneamente, resultados intermedios de expresiones complejas, bucles desenrollados, y funciones incorporadas que necesitan temporales. WGSL no te deja controlar esto directamente —no hay nada equivalente al maxrregcount de CUDA— pero sí puedes influir: acortar la vida de las variables, evitar mantener muchos valores vivos a la vez, y no forzar el desenrollado de bucles grandes.
La memoria compartida. Si tu workgroup declara 8 KiB de var<workgroup> y la unidad de cómputo dispone de 32 KiB, caben cuatro workgroups. Si declaras 16 KiB, caben dos. Es el límite más fácil de calcular porque el número está en tu código, y es el que hay que vigilar cuando se usa la técnica de tiling.
Las ranuras de planificación. Hay un máximo duro de grupos y de subgrupos residentes por unidad de cómputo, independientemente de que sobren registros y memoria. Es la razón de que workgroups muy pequeños desperdicien capacidad: con @workgroup_size(16) en un hardware que admite 16 grupos residentes, tienes 256 invocaciones vivas aunque hubiera sitio para 2048.
No hay ninguna API en WebGPU para consultar cuántos registros usa un shader ni cuántos grupos serán residentes. Esa información existe en las herramientas nativas de los fabricantes y no llega a la web. En la práctica, la ocupación en WebGPU se estudia por experimento: cambiar el tamaño de workgroup, simplificar el shader, medir. Es menos cómodo y el modelo mental sigue siendo igual de necesario para saber qué experimentos vale la pena hacer.
El cálculo, con un ejemplo
Un kernel de desenfoque con tiling, sobre una unidad de cómputo hipotética con 64 KiB de memoria compartida, 65536 registros y un máximo de 2048 invocaciones residentes.
Versión A, con @workgroup_size(16, 16) y un tile de 18 por 18 floats en memoria compartida:
- Invocaciones por grupo: 256. Memoria compartida por grupo: 18 por 18 por 4 bytes, o sea 1296 bytes.
- Límite por memoria compartida: 64 KiB entre 1296 son 50 grupos. No es el limitante.
- Si el shader usa 40 registros por invocación: 65536 entre 40 son 1638 invocaciones, o sea 6 grupos de 256.
- Límite por ranuras: 2048 entre 256 son 8 grupos.
- Manda el más restrictivo: 6 grupos, 1536 invocaciones, ocupación del 75%.
Versión B, idéntica pero el shader usa 72 registros porque alguien añadió tres variables vivas a la vez:
- 65536 entre 72 son 910 invocaciones, o sea 3 grupos.
- Ocupación del 37,5%. La mitad de capacidad de ocultación de latencia por tres variables.
Ese es el mecanismo por el que un cambio aparentemente inocuo cambia el rendimiento de forma desproporcionada. Y es también por qué el efecto es escalonado y no gradual: pasar de 40 a 41 registros puede no cambiar nada, y pasar de 43 a 44 puede tirar un grupo entero.
Cuándo no maximizar
Aquí es donde la mayoría del material se queda corto. La ocupación no es el objetivo: es un medio para esconder latencia. Si la latencia ya está escondida, subirla más no aporta nada, y hay tres casos en que conviene bajarla a propósito.
Cuando el kernel tiene mucha aritmética por acceso a memoria. Si la intensidad aritmética es alta, hay pocos ciclos que rellenar y una ocupación del 25% basta. En ese régimen, dejar que cada invocación use más registros —y por tanto mantener más valores en el nivel más rápido de la jerarquía en lugar de recalcularlos o volcarlos— sale ganando.
Cuando el tiling aporta más de lo que cuesta. Un tile más grande en memoria compartida reduce la ocupación y aumenta la reutilización de datos. Si la reutilización sube por un factor 4 y la ocupación baja un 30%, el cambio es bueno. Esta es exactamente la disyuntiva de la multiplicación de matrices.
Cuando el problema está limitado por ancho de banda. Más grupos residentes leyendo a la vez no aumentan el ancho de banda del bus: solo hacen más cola. Aquí subir la ocupación es literalmente inútil.
El síntoma que separa los casos es el experimento del nivel anterior: si añadir paralelismo mejora, faltaba ocupación; si no cambia nada, el cuello está en otro sitio.
Las palancas que tienes en WebGPU
Sin acceso a contadores de hardware, quedan cuatro decisiones, todas medibles por experimento directo.
El tamaño de workgroup. Múltiplo de 64, entre 64 y 256. Es el rango donde casi todo funciona bien. Vale la pena probar 64, 128 y 256 con el kernel real y quedarse con el que mida mejor; la diferencia entre ellos puede ser del 30% y el ganador cambia entre GPUs.
El tamaño de la memoria compartida. Está en tu código y el efecto es calculable. Reducirlo a la mitad puede duplicar los grupos residentes.
La presión de registros. Indirecta pero real: menos variables vivas a la vez, expresiones más cortas, no acumular resultados intermedios que se usan mucho después. Un shader escrito en pasos pequeños y secuenciales suele usar menos registros que uno con una expresión gigante.
El trabajo por invocación. Que cada invocación procese cuatro elementos en lugar de uno reduce el número total de invocaciones y aumenta el trabajo útil por unidad de ocupación. A veces es la optimización más eficaz de todas, porque amortiza el coste fijo de arranque de cada invocación entre varios elementos.
Hay una trampa en la que cae casi todo el mundo que descubre este concepto: tratar la ocupación como el objetivo, reescribir shaders para bajar registros, y acabar con código ilegible que no va más rápido. La ocupación no se optimiza: se optimiza la ocultación de latencia, y la ocupación es una de las tres formas de conseguirla.
Las otras dos se olvidan sistemáticamente y a menudo son mejores. La primera es el paralelismo a nivel de instrucción: si dentro de una misma invocación emites cuatro lecturas independientes antes de usar ninguna, las cuatro latencias se solapan entre sí, y eso oculta tanto como cuadruplicar los grupos residentes. Escribir let a = buf[i]; let b = buf[j]; let c = buf[k]; let d = buf[l]; y usarlos después es medible y no cuesta legibilidad.
La segunda es procesar varios elementos por invocación. Un kernel donde cada invocación maneja cuatro píxeles tiene una cuarta parte de invocaciones —peor ocupación nominal— y con frecuencia va más rápido, porque las cuatro lecturas se solapan, los cálculos comunes se hacen una vez y el coste fijo de arranque se amortiza. Hay kernels donde este cambio da un factor 2 con la ocupación bajando.
Y el consejo de método que evita perder días: la ocupación no se mide en WebGPU, así que trátala como una hipótesis, no como un dato. Si crees que un kernel sufre de baja ocupación, la prueba es reducir a la mitad su memoria compartida o su tamaño de grupo y volver a medir. Si el tiempo baja, tenías razón. Si no cambia, tu problema era otro y acabas de ahorrarte una tarde de reescribir un shader que ya estaba bien.
Con el hardware entendido, ya se puede escribir código sin ir a ciegas. El punto de entrada de cualquier programa de WebGPU es navigator.gpu y la detección de soporte.