Skip to content

English · Español

00 — Por qué una fase separada para el modelo mental de la GPU

La intuición de CPU (caches, predicción de saltos, ejecución especulativa, pocos hilos pesados) no se transfiere a GPU. La GPU es una máquina diferente — no un CPU más rápido. Esta fase construye el modelo mental nuevo antes de escribir cualquier kernel.

Esta página existe para combatir un único modo de fallo recurrente: tratar la GPU como "una CPU pero con más cores". Cada línea de theory/01..04 está calibrada contra esa idea equivocada. Si interiorizas esta página, lo demás es mecánico.

El cambiazo

Una CPU tiene 4–96 cores. Cada core es una unidad de ejecución gruesa: emisión out-of-order, predicción de saltos, ejecución especulativa, varios MB de L2 privada, 32+ KiB de L1, un motor complejo de desambiguación de memoria. La CPU está construida para hacer un hilo rápido sobre código impredecible. Para ir más rápido, paralelizas entre cores y lanes SIMD — ambos escasos.

Una GPU tiene 60–150 SMs. Cada SM tiene 32–128 unidades de ejecución finas (CUDA cores). Cada unidad fina tiene un estado privado mínimo. La GPU está construida para hacer correr millones de hilos simultáneamente sobre código predecible. Para ir más rápido, lanzas más hilos — que son abundantes.

Si llevas el modelo mental de CPU a una GPU, escribirás código que:

  • Lanza "solo unos pocos" hilos (1024 en vez de 1000000), pensando que es "eficiente". La GPU se queda al 1% de ocupación.
  • Tiene ramas dependientes de datos. La GPU ejecuta ambas ramas igualmente, tirando la mitad del trabajo.
  • Usa acceso a memoria indirecto / con persecución de punteros. La GPU hace una transacción de memoria por escalar en vez de una transacción por warp. Ancho de banda efectivo: 1/32 del pico.
  • Reserva y libera mucho. Cada cudaMalloc es una llamada síncrona de 50 microsegundos. Un bucle haciendo esto está muerto.
  • Mide wall-clock leyendo un timer de Python, sin cudaDeviceSynchronize(). Mides el overhead de lanzamiento, no el trabajo.

Cada uno de esos errores es invisible a la intuición de CPU. Cada uno cuesta 10–100×. La Fase 23 es la fase donde reemplazas la intuición. La Fase 24 es cuando, de otro modo, pagarías por no haberlo hecho.

Lo que se mantiene igual

Algunas intuiciones de la Fase 1 se transfieren directamente:

  • La jerarquía de memoria sigue siendo una jerarquía. La GPU tiene su propia versión: HBM (piensa en DRAM, más lenta que la cache, más grande que la cache), L2 (compartida entre SMs), SMEM (por bloque, controlada a mano — no gestionada automáticamente como la L1 de CPU), registros (por hilo). Los números son distintos. La forma es la misma.
  • La intensidad aritmética es el diagnóstico correcto. La ecuación del roofline \(\text{perf} = \min(\pi, I \beta)\) es idéntica. Solo cambian los números (peak FLOPS, peak BW) y la multiplicidad de dtypes (fp64/fp32/fp16/bf16/fp8/int8 cada uno tiene su propio \(\pi\) en una GPU con Tensor Cores).
  • El acceso coalescido importa. "Hilos adyacentes acceden a memoria adyacente" es la versión GPU de "el acceso lineal a memoria gana en CPU". Ambos van de cuenta de transacciones.
  • El razonamiento del roofline funciona. La decode-attention de la Fase 22 se sitúa en la pendiente del roofline de GPU en I≈1 (fp16), igual que en el roofline de CPU. Mismo diagnóstico (memory-bound), distintos números absolutos.

Lo que cambia

Tres cosas cambian cualitativamente y necesitan modelos mentales nuevos. Cada una es una página de teoría separada:

  1. Modelo de ejecución: SIMT, no OoO. Un warp de 32 hilos ejecuta la misma instrucción en paso de bloqueo. Las ramas desvían lanes; los lanes que no están en la rama activa se detienen. Es el concepto más contra-CPU y se cubre en 01-gpu-vs-cpu-mental-model.md.

  2. Jerarquía de memoria: SMEM manual. La SMEM (memoria compartida (shared memory)) no es una cache auto-gestionada. El programador elige qué cachear en SMEM, cuándo cargar, cuándo desalojar. Esto es potencia; también es la fuente de la mayoría de errores "este kernel es lento porque olvidé usar SMEM". Cubierto en 02-gpu-memory-hierarchy.md.

  3. Rendimiento vía ocupación, no vía "menos hilos más rápidos". El rendimiento de la GPU viene de que el SM se mantenga ocupado alternando entre warps a cada ciclo. La ocupación está limitada por la cantidad de registros por hilo, SMEM por bloque e hilos por bloque. Cada uno de estos es un presupuesto que asignas. Cubierto en 03-warps-and-occupancy.md.

Qué deberías ser capaz de hacer al final de la Fase 23

  1. Indicar, a partir del perfil del dispositivo de tu GPU cloud específica: peak fp16 TFLOPS (Tensor Core), peak HBM bandwidth, tamaño de L2, máximo hilos/bloque, registros/SM, SMEM/SM, compute capability.
  2. Calcular el machine balance \(I_\text{crit}\) para fp16 en tu GPU. Comparar con los 10 FLOPs/byte del i5-8250U y el ~1 FLOP/byte del operador de decode de la Fase 22. Indicar el régimen.
  3. Esbozar la jerarquía de memoria con los anchos de banda anotados.
  4. Predecir, solo a partir de las specs del dispositivo, lo que tardará un cudaMemcpy de N bytes en microsegundos, para N pequeño y N grande.
  5. Identificar dos razones por las que un warp puede correr al <100% de eficiencia (divergencia, acceso no-coalescido) y una razón por la que un SM puede correr al <100% de ocupación (presión de registros).
  6. Aprovisionar una instancia GPU en la nube, correr un benchmark, terminarla, todo en menos de 30 minutos, y producir un log de coste.

Ese último ítem es operativo, no intelectual — pero es lo que hace posible la Fase 24+ sin desfases de coste.

La conexión con la Fase 22

Lo más útil de la Fase 22 (KV cache) para la Fase 23 es su menú de operadores. Al final de la Fase 22 Borja ha medido:

  • Prefill: compute-bound en CPU, con intensidad escalando con \(P\).
  • Decode-step FFN: weight-read-bound en CPU.
  • Decode-step attention: ~0.5–1.0 FLOPs/byte, memory-bound en CPU.

La Fase 23 toma esos mismos operadores y los coloca sobre el roofline de GPU. El diagnóstico (cuál es compute-bound, cuál es memory-bound) es el mismo cualitativamente. Los números (cuán rápido, qué fracción del pico) son radicalmente distintos.

Esa continuidad es importante pedagógicamente. El roofline no es un ritual por-máquina — es el modelo mental que explica ambas máquinas. Misma ecuación, distintas constantes. Para la Fase 24, cuando Borja escriba un kernel fused-softmax y lo mida, el punto de ese kernel aterrizará sobre este roofline de la Fase 23. La Fase 23 construye el plot; la Fase 24 empieza a usarlo como herramienta.

Lo que esta fase deliberadamente no hace

  • No escribe un kernel. Fase 24. La Fase 23 solo usa kernels pre-construidos (cudaMemcpy, GEMM de cuBLAS).
  • No importa PyTorch. Fase 24. La comodidad de array-en-GPU de la Fase 23 es cupy (NumPy-en-GPU), que es una librería substrato, no un framework.
  • No ajusta. La Fase 23 mide; ajustar es Fase 24.
  • No hace multi-GPU. Fase 35.

La fase es intencionadamente una fase de "medida y orientación", exactamente como la Fase 1 lo fue para CPU.

Siguiente: theory/01-gpu-vs-cpu-mental-model.md — las diferencias del modelo de ejecución, en detalle.