Anatomía de una GPU: por qué no es una CPU con muchos núcleos
La arquitectura interna de una GPU moderna, la diferencia entre optimizar latencia y optimizar rendimiento agregado, y qué hay realmente dentro de una unidad de cómputo.
Todo el material de este nivel existe para responder una pregunta que aparece en cuanto empiezas a medir: por qué un cambio de dos líneas en un shader duplica el tiempo de fotograma, y por qué otro cambio que parece igual de grande no hace nada. Las respuestas no están en la API: están en cómo está construida la máquina. Un modelo mental correcto del hardware convierte esas sorpresas en predicciones.
- Contrastar el objetivo de diseño de una CPU con el de una GPU.
- Describir la jerarquía de unidades de cómputo, planificadores y bancos de ALU.
- Explicar cómo una GPU esconde la latencia de memoria en lugar de reducirla.
- Traducir los conceptos de hardware a los términos que usa WebGPU y WGSL.
Dos máquinas con objetivos opuestos
Una CPU está diseñada para terminar una secuencia de instrucciones lo antes posible. Todo su presupuesto de transistores va a eso: predicción de saltos para no parar en un if, ejecución fuera de orden para no parar cuando una instrucción depende de otra, ejecución especulativa para adelantar trabajo que quizá haga falta, y cachés grandes y de baja latencia para que la memoria no la frene. En un procesador de escritorio moderno, la parte del chip dedicada a ejecutar aritmética es una fracción pequeña; la mayoría es maquinaria para que esa aritmética no se detenga.
Una GPU está diseñada para terminar un millón de secuencias independientes en el menor tiempo total. El objetivo no es que ninguna acabe pronto, sino que el conjunto acabe pronto. Con ese objetivo, casi toda la maquinaria de la CPU es un desperdicio: si tienes un millón de tareas independientes y una se para esperando memoria, no necesitas predecirla ni adelantarla, necesitas simplemente ejecutar otra. Así que la GPU gasta los transistores en lo contrario: muchísimas unidades aritméticas, un fichero de registros enorme para mantener miles de tareas vivas a la vez, y un planificador que cambia de tarea sin coste.
La consecuencia es la frase que hay que interiorizar: una GPU no reduce la latencia, la esconde. Un acceso a memoria de vídeo tarda cientos de ciclos igual que en una CPU. La diferencia es que la GPU tiene otras miles de tareas listas para ocupar esos ciclos, mientras que la CPU tiene dos o cuatro por núcleo.
flowchart TB host[CPU y memoria del sistema] -->|bus PCIe o memoria unificada| vram[VRAM memoria global] vram --> l2[Cache L2 compartida] l2 --> cu1[Unidad de computo 1] l2 --> cu2[Unidad de computo 2] l2 --> cuN[Unidad de computo N] cu1 --> sch[Planificador de grupos de invocaciones] sch --> regs[Fichero de registros muy grande] sch --> alu[Bancos de ALU que operan en lockstep] sch --> smem[Memoria compartida en chip] sch --> tex[Unidades de textura y cache L1] style host fill:#89b4fa,color:#11111b style vram fill:#94e2d5,color:#11111b style l2 fill:#94e2d5,color:#11111b style cu1 fill:#fab387,color:#11111b style cu2 fill:#fab387,color:#11111b style cuN fill:#fab387,color:#11111b style sch fill:#cba6f7,color:#11111b style regs fill:#f9e2af,color:#11111b style alu fill:#a6e3a1,color:#11111b style smem fill:#f9e2af,color:#11111b style tex fill:#f9e2af,color:#11111b
Dentro de una unidad de cómputo
El nombre cambia según el fabricante y eso confunde más de lo necesario. NVIDIA la llama streaming multiprocessor, AMD compute unit, Intel Xe-core, Apple GPU core, Arm shader core. Es la misma idea: la unidad replicable de la que una GPU tiene entre cuatro y ciento cincuenta.
Dentro hay cinco piezas, y las cinco tienen consecuencias directas sobre cómo se escribe un shader.
Bancos de unidades aritméticas. Un grupo de ALU que ejecutan la misma instrucción sobre datos distintos en el mismo ciclo. No son núcleos independientes: comparten el decodificador de instrucciones. Esto es lo que obliga a la ejecución en lockstep y lo que hace que la divergencia cueste.
Un fichero de registros desproporcionadamente grande. Decenas o cientos de kilobytes por unidad de cómputo, frente a unos pocos cientos de bytes por núcleo en una CPU. Los registros no se guardan y restauran al cambiar de tarea: cada tarea viva tiene los suyos asignados de forma permanente mientras dura. Por eso el cambio de contexto es gratis, y por eso el número de tareas que caben depende de cuántos registros consuma cada una.
Un planificador de grupos. Elige en cada ciclo qué grupo de invocaciones tiene sus datos listos y le da paso. Es hardware, no software, y su decisión cuesta cero ciclos.
Memoria compartida en chip. Del orden de decenas de kilobytes, direccionable explícitamente desde el shader, con una latencia de unos pocos ciclos frente a los cientos de la memoria global. En WGSL es var<workgroup>. Es el recurso que hace que un algoritmo bien escrito vaya diez veces más rápido que el mismo algoritmo mal escrito.
Unidades de función fija. Muestreo de texturas con filtrado e interpolación, interpolación de atributos, operaciones de rasterización. Hacen en un ciclo cosas que en ALU costarían decenas. Por eso una lectura de textura con textureSample puede ser más barata que emular ese filtrado a mano.
El coste de las cosas, en órdenes de magnitud
Estas cifras varían mucho entre arquitecturas, pero los órdenes de magnitud relativos son estables y son lo que hay que tener en la cabeza:
| Operación | Coste relativo aproximado |
|---|---|
| Suma o multiplicación en punto flotante de 32 bits | 1 |
| Multiplicación y suma fusionadas | 1 |
| Raíz cuadrada inversa, seno, coseno, exponencial | 4 a 8 |
| División en punto flotante | 4 a 10 |
| Lectura de un registro | prácticamente 0 |
| Lectura de memoria compartida sin conflicto de bancos | 2 a 5 |
| Lectura de caché L1 o de textura | 20 a 40 |
| Lectura de memoria global con acierto de L2 | 100 a 200 |
| Lectura de memoria global sin acierto de caché | 300 a 800 |
Léelo dos veces. Un acceso a memoria global cuesta lo que varios cientos de multiplicaciones. Esa sola línea reordena todas las intuiciones de optimización que traes de la CPU. En una GPU, recalcular un valor casi siempre es más barato que guardarlo en memoria y volver a leerlo. Una tabla de consulta precalculada, que en CPU es una optimización clásica, en GPU suele ser más lenta que evaluar la función.
Cuando dudes entre dos implementaciones de un shader, cuenta accesos a memoria, no operaciones aritméticas. La que hace menos viajes a memoria gana casi siempre, aunque haga el doble de cuentas.
Cómo se llaman estas cosas en WebGPU
El mapa entre hardware y API es directo y conviene fijarlo ya, porque el resto del nivel lo usa constantemente.
Una invocación es la instancia de ejecución de tu shader para un elemento: un vértice, un fragmento o un punto de la rejilla de cómputo. Es lo más parecido a un hilo, con la salvedad de que no es independiente.
Un grupo de trabajo —workgroup— es el conjunto de invocaciones que declaras con @workgroup_size(x, y, z) y que se ejecutan juntas en una unidad de cómputo, compartiendo la memoria de var<workgroup> y pudiendo sincronizarse con workgroupBarrier(). Es una construcción de la API que se apoya en una realidad del hardware.
Un grupo de invocaciones en lockstep —warp en NVIDIA, wavefront en AMD, subgroup en la terminología portable— es la unidad real de ejecución del hardware: 32 o 64 invocaciones que avanzan a la vez sobre la misma instrucción. WGSL no lo expone en su núcleo, aunque hay extensiones de subgrupos, y sin embargo determina el rendimiento de todo lo que escribas.
El dispatch es la rejilla completa: dispatchWorkgroups(x, y, z) lanza esa cantidad de grupos, y el total de invocaciones es el producto por el tamaño del grupo.
Y en el camino de render, la correspondencia es que el rasterizador es un generador de invocaciones: convierte triángulos en fragmentos y los agrupa en bloques que se ejecutan en lockstep, exactamente igual que un dispatch de cómputo, solo que el reparto lo decide el hardware.
Hay una imagen mental que hace daño: la GPU como una fábrica de cálculo a pleno rendimiento donde el objetivo es reducir el número de operaciones. La realidad medida es la contraria. En un shader típico de una aplicación real, las unidades aritméticas están ociosas la mayor parte del tiempo, esperando datos. La utilización aritmética efectiva de muchos shaders de producción está por debajo del 20%.
Eso reordena por completo el orden de las optimizaciones. Quitar una multiplicación de un shader que espera memoria no acelera nada: la ALU tenía tiempo de sobra. Lo que acelera es reducir el número de viajes a memoria, mejorar el patrón de acceso para que las lecturas de invocaciones vecinas caigan en la misma línea de caché, o aumentar el número de grupos vivos para que haya más trabajo con el que rellenar la espera.
El corolario más contraintuitivo, y el que da sentido al resto del nivel: añadir cálculo a un shader puede salir gratis. Si el shader está limitado por memoria, las operaciones extra se ejecutan en ciclos que de todas formas estaban perdidos. Es la razón por la que técnicas que suenan caras —evaluar un ruido procedural en lugar de leer una textura, recalcular una normal en vez de almacenarla— salen ganando en la práctica. Y es también la razón por la que hay que medir antes de optimizar: la intuición de coste de la CPU no solo no sirve, sino que apunta en la dirección equivocada.
Con la máquina descrita, toca mirar de cerca la pieza que más consecuencias tiene: la ejecución en grupos de invocaciones en lockstep.