El ratio aritmético/memoria

Hay un número que predice, antes de escribir una línea, si tu kernel va a estar limitado por la memoria o por el cálculo: las operaciones de coma flotante que haces por cada byte que traes de memoria global. Se calcula con lápiz en dos minutos, se compara con una constante que depende solo de la máquina, y el resultado te dice cuáles de las optimizaciones que estabas considerando no van a servir absolutamente para nada. Es la métrica más rentable de todo el nivel porque su valor está en lo que te ahorra hacer.

Cómo se calcula la intensidad aritmética

La definición es literal: operaciones de coma flotante dividido entre bytes leídos o escritos en memoria global. Se cuenta por elemento de trabajo y las unidades son FLOP por byte.

Primer ejemplo, hecho a mano. Una suma de vectores:

@group(0) @binding(0) var<storage, read>       a: array<f32>;
@group(0) @binding(1) var<storage, read>       b: array<f32>;
@group(0) @binding(2) var<storage, read_write> c: array<f32>;

@compute @workgroup_size(64)
fn sumar(@builtin(global_invocation_id) gid: vec3u) {
  let i = gid.x;
  if (i >= arrayLength(&a)) { return; }
  c[i] = a[i] + b[i];
}

Por cada elemento: dos lecturas y una escritura de 4 bytes, doce bytes en total, y una sola suma. La intensidad es 1/12, es decir 0,083 FLOP por byte. No hay nada que puedas hacer dentro de ese kernel para cambiarla: no hay ninguna operación que quitar y no hay ningún byte que ahorrar mientras los datos sigan siendo f32. Está irremediablemente limitado por memoria, y cualquier tiempo que dediques a optimizar su aritmética es tiempo tirado. El único margen es el formato de los datos y fusionar la operación con la que venga después.

Segundo ejemplo, el opuesto. La multiplicación de matrices con bloques de lado T que has visto en reutilización de datos. Por cada elemento del resultado se hacen 2K operaciones —K multiplicaciones y K sumas— y se leen 2K/T valores de memoria global, que son 8K/T bytes. Divide:

intensidad = 2K / (8K/T) = T/4  FLOP por byte

La intensidad es proporcional al lado del bloque, y eso es exactamente por qué el tiling funciona: no reduce el cálculo, mueve el kernel hacia la derecha en el eje de la intensidad. Con T igual a 16 son 4 FLOP por byte, dieciséis veces más que la versión directa, que sale a 0,25. Con bloqueo en registros y un bloque efectivo de 64, son 16.

Un panorama de kernels reales, con las cuentas hechas sobre el tráfico obligatorio, es decir contando cada dato una sola vez:

Salvo los dos últimos, que son juguetes de demostración, todo lo demás vive por debajo de 20.

El modelo roofline sin gráficos

El modelo se enuncia en una línea y no necesita dibujo. El rendimiento que puedes alcanzar es el menor de dos techos:

rendimiento alcanzable = min( pico de cálculo , intensidad · ancho de banda )

El primer término es una constante horizontal: por muchos bytes que traigas, la GPU no hace más FLOP por segundo de los que tiene unidades. El segundo es una rampa que crece con tu intensidad: si tu kernel hace pocas operaciones por byte, tu rendimiento está atado a la velocidad del bus. Los dos se cruzan en un punto que se llama intensidad crítica y que se calcula dividiendo:

intensidad crítica = pico de cálculo / ancho de banda

Con números redondos de una GPU dedicada de gama media, 10 TFLOP/s y 400 GB/s, sale 10·10¹² / 400·10⁹ igual a 25 FLOP por byte. Y lo interesante es lo estable que es ese número en todo el rango de hardware:

Todas las cifras son órdenes de magnitud, no especificaciones de producto, pero la conclusión aguanta cualquier ajuste razonable: la intensidad crítica de una GPU moderna está entre 20 y 50 FLOP por byte, y lleva décadas subiendo porque la capacidad de cálculo crece más deprisa que el ancho de banda de la memoria. Es una propiedad estructural de la clase de máquina, no una peculiaridad de un modelo.

Ahora vuelve a la tabla anterior y compara. Si tu kernel está por debajo de 25, ninguna optimización aritmética te va a servir: puedes eliminar la mitad de las operaciones y el tiempo no se moverá, porque la unidad de ejecución ya está esperando a la memoria la mayor parte del tiempo. La mayoría de los kernels reales están uno o dos órdenes de magnitud por debajo del cruce. Esa es la conclusión incómoda de todo el nivel, y explica por qué las optimizaciones que la gente intenta primero —quitar una raíz cuadrada, reemplazar una división— casi nunca cambian nada medible.

Los dos números que necesitas para tu propia máquina se miden, porque WebGPU no los expone. El ancho de banda real se obtiene con un kernel de flujo puro: leer un buffer grande y escribir otro, sin cálculo, con acceso perfectamente coalescente; divides los bytes movidos entre el tiempo con timestamp-query y tienes tus GB/s efectivos, que estarán entre el 60 y el 85 por ciento del pico anunciado. El pico de cálculo se obtiene con un bucle de fma encadenadas en varias cadenas independientes sobre valores en registros, sin tocar memoria. Con esos dos números tienes tu intensidad crítica real, que es más útil que la del fabricante.

Qué hacer en cada zona

Si estás por debajo del cruce, limitado por memoria, solo hay dos palancas: mover menos bytes y reutilizar más lo que mueves.

Reducir bytes empieza por el formato. La feature shader-f16 permite usar medias precisiones en el shader y en los buffers, lo que divide el tráfico por dos y en la mayoría del hardware moderno además duplica el rendimiento aritmético, porque las unidades procesan dos f16 por carril:

enable f16;

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

@compute @workgroup_size(64)
fn escalar(@builtin(global_invocation_id) gid: vec3u) {
  let i = gid.x;
  if (i >= arrayLength(&pesos)) { return; }
  salida[i] = pesos[i] * 2.0h;
}

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

El empaquetado va más lejos donde el rango lo permite. pack4x8unorm mete cuatro valores normalizados entre cero y uno en un solo u32, y unpack4x8unorm los recupera: cuatro bytes en vez de dieciséis, un factor de cuatro en tráfico a cambio de una instrucción de empaquetado que es gratis cuando estás limitado por memoria. Para pares de medias precisiones sin activar la feature están pack2x16float y unpack2x16float.

// Colores, normales o pesos en el rango [0,1]: 4 bytes en vez de 16.
let empaquetado: u32 = pack4x8unorm(vec4f(r, g, b, a));
let recuperado: vec4f = unpack4x8unorm(empaquetado);

Las texturas comprimidas son la misma idea con soporte de hardware: texture-compression-bc en escritorio, texture-compression-etc2 y texture-compression-astc en móvil. BC7 ocupa un byte por téxel frente a los cuatro de RGBA8 y BC1 medio byte, así que el factor va de cuatro a ocho. Y lo que las hace especiales no es el ahorro en disco sino que la GPU las muestrea comprimidas: el ahorro se aplica al bus y a la caché, no solo al almacenamiento. Ninguna de las tres está garantizada, así que comprueba adapter.features.has(...) y ten un camino alternativo.

Las otras dos palancas de esta zona ya las conoces: fusionar pasadas para eliminar el viaje completo de un buffer intermedio, y subir la reutilización con bloques en memoria compartida.

Si estás por encima del cruce, limitado por cálculo, la lista es distinta. Reduce operaciones sacando invariantes de los bucles. Usa fma(a, b, c) en vez de a * b + c donde el redondeo único no te moleste, porque es una sola instrucción en vez de dos. Sustituye funciones caras por aproximaciones: inverseSqrt en vez de dividir por sqrt, una multiplicación por un recíproco precalculado en vez de una división, un polinomio en vez de un pow con exponente constante. Y vigila la divergencia dentro del subgrupo, porque una rama que toman la mitad de los carriles ejecuta las dos ramas en serie y duplica el coste sin que nada lo indique.

La lista de comprobación, ordenada por impacto

No por elegancia ni por lo satisfactorio que sea escribirla, sino por el factor que suele dar. Bájala en orden y para cuando el kernel deje de ser el cuello de botella.

1. No calcular lo que no se ve. Culling, niveles de detalle, media resolución para efectos de baja frecuencia. Ahorra el cien por cien de lo que elimines y ninguna otra optimización compite con eso.
2. Fusionar pasadas. Cada buffer intermedio que desaparece se lleva una escritura completa y una lectura completa. Factor de 1,5 a 3 en cadenas de post-procesado.
3. Reducir los bytes por elemento. Formatos más pequeños, f16, empaquetado, texturas comprimidas. Factor de 2 a 8, y es el que más veces está disponible.
4. Arreglar el patrón de acceso. Coalescencia, estructura de arrays, alineación. Hasta factor 32 en el caso patológico y no cuesta más que una tarde de reorganizar los datos.
5. Reutilizar con bloques. Factor igual al lado del bloque. Solo aplica si el mismo dato lo leen varios hilos del mismo workgroup.
6. Ajustar el tamaño del workgroup y la ocupación. Factor de 1,1 a 2, y a veces cero. Es barato de probar con un barrido, así que hazlo, pero no esperes milagros.
7. Micro-optimizar la aritmética. Factor de 1,05 a 1,3, y exactamente cero si estás limitado por memoria, que es el caso probable. Va al final por eso.

El orden no es arbitrario: los primeros puntos reducen trabajo, los del medio reducen bytes, y los últimos reducen ciclos. En una GPU moderna, el trabajo vale más que los bytes y los bytes valen mucho más que los ciclos.