Ocupación y sus límites

Una GPU esconde la latencia de la memoria cambiando a otro grupo de hilos mientras uno espera, así que todo depende de cuántos grupos haya residentes a la vez. Esa proporción se llama ocupación y la determinan tres recursos finitos que se reparten entre los hilos: los registros, la memoria compartida y el tamaño del workgroup. Ninguno de los tres se puede consultar desde WebGPU. Lo que sí puedes es calcular los dos primeros con lápiz, controlar el tercero por completo, y medir el resultado.

Qué es la ocupación y por qué decide el rendimiento

Una unidad de ejecución de una GPU puede tener varios workgroups residentes a la vez: sus registros están asignados, su memoria compartida reservada, su estado vivo. El planificador elige en cada ciclo, de entre todos los grupos SIMD residentes, uno que tenga una instrucción lista y la emite. La ocupación es la razón entre los grupos residentes y el máximo que el hardware admite, y se suele expresar en porcentaje.

Importa por un único motivo, y ya lo has visto en el coste de la memoria: es el mecanismo con el que la GPU tapa los cientos de ciclos de una lectura de VRAM. Con dieciséis grupos residentes, que uno se bloquee 500 ciclos es irrelevante porque quedan quince con trabajo. Con dos grupos residentes y ambos esperando, la unidad se queda muda y esos 500 ciclos se pierden enteros. El shader es el mismo, la latencia es la misma; lo único que cambia es cuánta gente hay en la cola.

Cambiar de grupo cuesta cero ciclos, y esa es la propiedad que lo hace funcionar: cada grupo residente conserva sus propios registros, así que no hay estado que guardar ni restaurar. El precio de esa gratuidad es justamente el que estudia esta lección: los registros de todos los grupos residentes tienen que caber simultáneamente en el banco físico de la unidad. De ahí sale todo lo demás.

Los tres techos

Registros

Cada unidad de ejecución tiene un banco de registros de tamaño fijo que se reparte entre todos los hilos residentes. El orden de magnitud es de 64 KB por unidad en muchas arquitecturas, y llega a 256 KB en las de gama alta; ninguna GPU lo publica en una tabla accesible desde el navegador, así que trabaja con un modelo y valida midiendo.

Toma un modelo concreto: banco de 16384 registros de 32 bits, que son 64 KB, y un máximo de 1024 invocaciones residentes por unidad.

Si tu shader usa 64 registros por hilo, caben la mitad de hilos que si usa 32. Y la curva real no es suave sino escalonada, porque los registros se asignan en bloques y por grupo SIMD entero: pasar de 32 a 33 registros puede tirarte un escalón completo y costarte un tercio de la ocupación por una variable.

Lo que gasta registros no es el número de variables que declares sino el máximo de variables vivas a la vez, es decir cuántos valores tienen que existir simultáneamente en el punto más denso del shader. Tres fuentes lo disparan:

• Bucles desenrollados. Cuando el compilador desenrolla un bucle de ocho iteraciones, los ocho resultados intermedios pueden acabar vivos a la vez. El desenrollado es casi siempre bueno para la latencia y casi siempre malo para la ocupación.
• Arrays locales indexados dinámicamente. Un array declarado dentro de una función e indexado con un valor de ejecución no puede vivir en registros, porque los registros no son direccionables. El compilador lo coloca en memoria privada respaldada por VRAM, y cada acceso pasa de cero ciclos a los 300 u 800 de una lectura global. Es el peor accidente de rendimiento que se puede cometer en un shader, y el código no lo insinúa.
• La propia ocultación de latencia dentro del hilo. El compilador adelanta las cargas para solaparlas con el cálculo, y adelantar una carga alarga el intervalo de vida de su destino. Está cambiando registros por latencia escondida en tu nombre, y no le puedes decir que no.

Memoria compartida

maxComputeWorkgroupStorageSize vale 16384 bytes por defecto y es el total que puede pedir un workgroup. Lo que decide la ocupación es cuánta tiene la unidad de ejecución en total —del orden de 64 KB, otra vez dependiente del hardware— dividido entre lo que pide cada grupo.

Si tu workgroup pide los 16 KB del límite, solo caben cuatro. Y si además cada uno tiene 256 invocaciones, son 1024 hilos residentes: puede que te baste, o puede que acabes de renunciar a la mitad de la capacidad de la unidad por un tile que no reutilizabas tanto.

Tamaño del workgroup

maxComputeInvocationsPerWorkgroup vale 256 por defecto, maxComputeWorkgroupSizeX y maxComputeWorkgroupSizeY valen 256 cada uno y maxComputeWorkgroupSizeZ vale 64. Dentro de ese margen mandan dos efectos.

El primero es el desperdicio de carriles: un tamaño que no sea múltiplo del tamaño del subgrupo —32 en muchas GPU, 64 en otras— deja carriles apagados en el último grupo SIMD, y esos carriles no se pueden rellenar con trabajo de otro workgroup.

El segundo es la granularidad de la residencia, que casi nadie tiene en cuenta y que es el argumento fuerte a favor de los grupos pequeños. Los workgroups se asignan enteros o no se asignan. Si tu consumo de registros permite 700 invocaciones residentes, con grupos de 256 caben dos —512 hilos, el 73 por ciento de lo disponible— y con grupos de 64 caben diez —640 hilos, el 91 por ciento—. El mismo shader, la misma presión de registros, dieciocho puntos de ocupación de diferencia solo por el tamaño del bloque.

Elegir el tamaño con criterio

Por qué 256 no siempre es mejor, aunque agrupe más trabajo: la granularidad de residencia empeora, la presión de registros de un grupo grande es mayor y la cola del dispatch se nota más —con 5000 elementos y grupos de 256 salen 20 workgroups, y si tu GPU tiene 30 unidades de ejecución hay diez sin trabajo—. Además, cada workgroupBarrier() obliga a esperar a las 256 invocaciones en vez de a 64, y el tiempo de una barrera lo marca la más lenta.

Y ahora el matiz que casi nadie cuenta: la ocupación máxima no es el objetivo. Es un medio, y solo compra una cosa: latencia escondida. No compra ancho de banda. Un kernel que usa muchos registros porque cada hilo calcula cuatro salidas y las mantiene en registros puede correr al 25 por ciento de ocupación y aun así ganarle a la versión de una salida por hilo al 100 por ciento, porque mueve cuatro veces menos bytes por resultado. Es un resultado clásico —lo popularizó Vasily Volkov en 2010 con el título Better Performance at Lower Occupancy— y sigue siendo verdad quince años después.

Hay incluso un efecto en el que más ocupación empeora las cosas: más hilos residentes significa un conjunto de trabajo instantáneo mayor, y un conjunto de trabajo mayor expulsa líneas de la caché antes de que se reutilicen. En un kernel con localidad temporal, subir la ocupación puede aumentar los fallos de L2. Como regla operativa, alrededor del 50 por ciento de ocupación ya suele bastar para saturar la ocultación de latencia, y a partir de ahí la curva es plana o descendente.

Medirla en WebGPU: la parte honesta

No existe ninguna API en WebGPU para consultar la ocupación, el número de registros que usa tu shader ni el tráfico de memoria que genera. Ni una. Lo único que la plataforma te da como pista real es el tamaño del subgrupo, en GPUAdapterInfo:

const adapter = await navigator.gpu.requestAdapter();
const info = adapter.info;              // tambien accesible como device.adapterInfo
console.log(info.vendor, info.architecture);
console.log(info.subgroupMinSize, info.subgroupMaxSize);

const device = await adapter.requestDevice({
  requiredFeatures: adapter.features.has('timestamp-query') ? ['timestamp-query'] : [],
});

Si subgroupMinSize y subgroupMaxSize valen los dos 32, sabes que el hardware ejecuta en grupos de 32 y que 64 es un tamaño seguro. Si valen 32 y 64, el compilador puede elegir, y te conviene un múltiplo de 64. No te dice nada sobre registros, pero te quita la única variable que podías estar adivinando mal por completo.

Todo lo demás se infiere midiendo, y la forma de hacerlo es un barrido. WGSL permite que @workgroup_size tome una constante sobreescribible, así que un mismo módulo genera todos los pipelines del barrido:

override TAM: u32 = 64u;

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

@compute @workgroup_size(TAM)
fn principal(@builtin(global_invocation_id) gid: vec3u) {
  let i = gid.x;
  if (i >= arrayLength(&entrada)) { return; }
  salida[i] = sqrt(abs(entrada[i])) * 1.0001;
}

const modulo = device.createShaderModule({ code: wgsl });

for (const tam of [32, 64, 128, 256]) {
  const pipeline = await device.createComputePipelineAsync({
    layout: 'auto',
    compute: { module: modulo, entryPoint: 'principal', constants: { TAM: tam } },
  });
  const grupos = Math.ceil(n / tam);
  console.log(tam, await medirConTimestamps(pipeline, grupos));
}

La restricción a recordar: una override puede dimensionar @workgroup_size, pero no puede dimensionar un array en el espacio workgroup, cuyo tamaño tiene que ser una expresión constante del módulo. En cuanto tu kernel tenga memoria compartida atada al tamaño del grupo, el barrido necesita varios módulos, uno por tamaño.

Y esto es lo que dice la forma de la curva, que es la única inferencia legítima que puedes hacer:

• Plana desde 64 en adelante. Estás limitado por ancho de banda. La ocupación no es tu problema y subirla no va a hacer nada; ve a reducir bytes.
• Mejora clara hasta 256. Estabas limitado por latencia con pocos hilos en vuelo. Sube el tamaño o dale a cada hilo más lecturas independientes.
• Empeora a partir de 128. Presión de registros o granularidad. Prueba a partir el kernel en dos o a reducir variables vivas.
• 32 compite con 64. Tu kernel tiene poco estado vivo y mucho trabajo independiente por hilo; probablemente ya estás saturando algo que no es la ocupación.