Skip to content

English · Español

01 — GPU vs CPU: el modelo de ejecución

SIMT (Single Instruction, Multiple Threads) significa que la GPU corre 32 hilos en paso de bloqueo dentro de un warp. Si el código tiene un if con ramas distintas en lanes diferentes, la GPU ejecuta ambas ramas y enmascara las lanes inactivas. No es "ramas baratas como en CPU OoO". Es "ramas que te cuestan los lanes que descartan".

Esta página explica el modelo de ejecución SIMT y tres de sus consecuencias que violan la intuición de CPU. Al final deberías ser capaz de predecir, dado un esbozo de kernel, qué líneas perjudican el rendimiento y por qué.

CPU: ejecución out-of-order

Un core moderno de CPU es un motor de ejecución especulativa con una estructura grande de contabilidad (reorder buffer, register file con renombrado, branch predictor, desambiguación de memoria, etc.). Un hilo emite instrucciones; la CPU encuentra oportunidades para solaparlas, predecirlas y reordenarlas. Una rama es barata si se predice correctamente y solo levemente cara (10–20 ciclos) en caso de fallo de predicción. Las cargas de memoria empiezan pronto y el resto del pipeline continúa. Los hilos en la CPU son pesados: cada uno tiene su propio register file, su pila gestionada por el kernel, su entrada en el scheduler.

El eje dominante de paralelismo por core de CPU es el paralelismo a nivel de instrucción (ILP): hacer que el stream de un hilo se solape lo máximo posible. Para escalar más allá de un core, lanzas otro hilo pesado.

GPU: SIMT

Un SM de GPU contiene 4–8 warp schedulers, cada uno gestionando un pool de warps (32 hilos cada uno). A cada ciclo, un scheduler elige un warp que esté listo (sus operandos están en registros, sin stall) y emite una instrucción en los 32 hilos en paso de bloqueo. Cada hilo tiene su propio program counter (formalmente — en realidad comparten un PC excepto entre ramas), su propia asignación de register file (recortada del register file del SM), y un thread-id que determina qué datos procesa.

El eje dominante de paralelismo por SM es el paralelismo a nivel de hilo (TLP) a granularidad de warp: mantener muchos warps disponibles para que cuando uno se detenga (esperando memoria), otro corra en el ciclo siguiente. Sin reorder buffer, sin branch predictor en el sentido CPU, sin ejecución especulativa.

Rendimiento-por-SM = instrucciones-warp-por-ciclo × FLOPs-por-instrucción-warp.

Para maximizar esto, el SM necesita warps residentes suficientes como para que cada ciclo tenga al menos uno listo para emitir. Ése es todo el juego.

Consecuencia 1: Divergencia de ramas

CPU: if (x > 0) a; else b; — el predictor adivina, el camino erróneo se descarta si falla. Coste (peor caso): 10 ciclos para un hilo.

GPU: mismo código en un warp. Si la mitad de los hilos tiene x > 0 y la otra mitad no, el warp ejecuta a para la primera mitad (16 lanes), enmascarando los demás (su resultado se descarta), luego ejecuta b para la segunda mitad. Ambas ramas corren. El rendimiento cae 2× para ese warp durante la región ramificada.

Si anidas ramas tres niveles con divergencia total en cada nivel, caes a ⅛ del rendimiento. Ejemplo real: un kernel que despacha "este token atiende, este otro es padding" vía rama en tiempo de ejecución será ~2× más lento que el mismo kernel procesando solo tokens reales.

Fix: estructura tu kernel para que los warps sean coherentes — los 32 hilos del warp toman la misma rama. Si la divergencia es inherente (p.ej., diferentes secuencias en un batch tienen distintas longitudes), usa despacho a nivel de bloque: asignar cada bloque a una secuencia para que los warps dentro de un bloque sean coherentes.

Consecuencia 2: Acceso coalescido a memoria

CPU: las lecturas dispersas de DRAM son lentas porque cada cache-line fill trae 64 bytes que puede que no uses, pero el subsistema de memoria sigue manejándolas eficientemente — el prefetcher aprende los patrones de acceso, la cache cachea todo lo que llega.

GPU: un warp haciendo 32 cargas de memoria en una instrucción pide al subsistema de memoria 32 direcciones. Si esas direcciones son adyacentes (el hilo i lee addr + i * 4), la GPU las consolida en una transacción de 128 bytes. Si están dispersas (el hilo i lee addr + perm[i] * 4 para alguna permutación perm), la GPU emite 32 transacciones separadas. El ancho de banda efectivo cae a ≈1/32.

Fix: estructura los datos para que hilos adyacentes accedan a memoria adyacente. El layout "structure-of-arrays" (SoA) se prefiere sobre "array-of-structures" (AoS) precisamente porque SoA da este coalescing gratis. El layout (B, H, S, d_h) de la KV cache de la Fase 22 fue elegido en parte pensando en coalescing — hilos adyacentes trabajando sobre d_h adyacentes obtienen lecturas coalescidas.

El ejemplo estándar de memory-coalescing que verás de nuevo en la Fase 24:

// Coalesced — good.
__global__ void copy_coalesced(float* dst, float* src, int N) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < N) dst[tid] = src[tid];          // thread i reads src[i], adjacent
}

// Uncoalesced — bad.
__global__ void copy_strided(float* dst, float* src, int N, int stride) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < N) dst[tid] = src[tid * stride]; // thread i reads src[i*stride]
}

Con stride=32, el segundo kernel corre 32× más lento en la pata de memoria del trabajo — cada carga es su propia transacción.

Consecuencia 3: Ocultar la latencia mediante cambio de warp

CPU: ocultar la latencia es trabajo de la CPU, vía ILP + cache + ejecución especulativa. El programador no piensa en ello (hasta que perfila).

GPU: ocultar la latencia es también automático, pero vía cambio de warp. Cuando el warp A emite una carga de memoria que falla L2 y debe ir a HBM (~400 ciclos), el warp scheduler cambia al warp B (ya residente, registros listos) y emite su siguiente instrucción. Para cuando los warps B, C, D, E, F, G, H, I han emitido cada uno una instrucción, la carga del warp A ha vuelto, y el scheduler puede emitir la siguiente instrucción de A.

Este mecanismo solo funciona si hay suficientes warps residentes. Ocupación = (warps activos / warps máximos por SM). Por debajo de ~50% de ocupación en un kernel memory-bound, el SM se detiene durante el acceso a memoria — no hay suficientes otros warps a los que cambiar. Por encima de ~50% de ocupación, la latencia se oculta completamente (en principio).

Los kernels compute-bound son menos sensibles a la ocupación: incluso al 25% de ocupación, puedes saturar las FPUs si el kernel hace suficiente aritmética por cada acceso a memoria.

Lo que nada de esto significa

Algunas lecturas erróneas de lo anterior que se repiten:

  • "Las ramas en GPU son caras, así que escribiré código sin ramas." Solo importa dentro de un warp. Las ramas entre warps son gratis (los warps corren independientemente). Y muchos trucos "sin ramas" (enmascarado, predicación) son exactamente lo que la GPU hace internamente para la divergencia — no estás evitando el coste, solo lo estás escribiendo explícitamente.
  • "A la GPU le encanta el acceso denso, así que rediseñaré mi estructura de datos." Sí — pero el rediseño a veces añade cómputo. Corre el análisis de roofline: si tu kernel es compute-bound, el acceso no-coalescido puede no ser el cuello de botella.
  • "Más hilos = más velocidad." Solo hasta saturar la ocupación. Pasado eso, más hilos no ayuda (el SM ya está manteniendo los warps intercambiándose). Por debajo de la saturación, más hilos residentes = más velocidad, pero más hilos en vuelo no ayuda necesariamente.
  • "SIMT == SIMD." Cercanos pero no idénticos. SIMD es una instrucción, lanes de ancho fijo; SIMT es una instrucción, lanes con su propio thread-id y estado por hilo (p.ej., registros). La diferencia práctica: SIMT puede hacer ramas divergentes (lentamente); SIMD no. SIMT puede hacer scatter/gather (lentamente); SIMD a menudo no puede en absoluto.

El modelo CUDA en un diagrama

Grid (host launches it)
└── Blocks (each block runs on ONE SM, can't migrate)
    └── Warps (32 threads, lockstep execution within a warp)
        └── Threads (own register-file allocation; thread-id determines data)

SM (Streaming Multiprocessor)
├── Warp schedulers (4–8)
├── CUDA cores (FP32 ALUs, e.g., 128 per SM on A100)
├── Tensor Cores (matrix-multiply units, e.g., 4 per SM on A100)
├── Register file (e.g., 64K 32-bit registers per SM)
├── SMEM (e.g., 96–164 KiB per SM, programmer-managed)
└── L1 cache (sometimes unified with SMEM)

Un grid se lanza; CUDA asigna bloques a SMs; cada SM elige warps de sus bloques asignados; los warp schedulers emiten. Un bloque nunca migra entre SMs (para que SMEM y __syncthreads() funcionen). Los hilos en bloques distintos no pueden sincronizarse (salvo saliendo del kernel y arrancando otro).

Todo este modelo es explícito en el modelo de programación CUDA. Declaras la forma de grid + block en el lanzamiento (sintaxis <<<grid, block>>>). La CPU no tiene nada análogo — no dices "este bucle itera como un grid de tamaño 1024×32". En la GPU, sí.

Lo que ahora deberías ser capaz de hacer

  1. Leer un kernel CUDA e identificar posible divergencia de warp.
  2. Identificar el layout de datos que maximiza el coalescing para un patrón de acceso dado.
  3. Decidir si un kernel dado está limitado por ocupación y cuál recurso (registros, SMEM, hilos) lo limita.
  4. Explicar por qué "correr más hilos siempre es más rápido en una GPU" es falso.
  5. Indicar la diferencia entre SIMT y SIMD sin tropezar con el vocabulario.

Lo que esta página NO cubre

  • Sintaxis de kernel. Nada de CUDA C++, nada de DSL Triton. Eso es Fase 24.
  • Arquitecturas GPU específicas. Volta vs Turing vs Ampere vs Hopper difieren en detalles (generaciones de Tensor Cores, async copy, TMA) — cubierto cuando haga falta en Fase 27/36. Esta página construye el modelo mental transversal a arquitecturas.
  • Multi-GPU. Fase 35.

Siguiente: theory/02-gpu-memory-hierarchy.md — HBM, L2, SMEM, registros, con anchos de banda y la regla de coalescing formalizada.