Rendimiento: El Modelo Roofline y Cuellos de Botella

Dos Cuellos de Botella: Cómputo y Memoria

Cada kernel de GPU está limitado por uno de dos recursos: cómputo (qué tan rápido los núcleos pueden procesar números) o ancho de banda de memoria (qué tan rápido los datos pueden fluir entre la memoria y los núcleos).

• Limitado por cómputo (compute-bound): los núcleos son el cuello de botella. Hacer la memoria más rápida no ayudaría — los núcleos no pueden procesar datos más rápido. Ejemplo: una multiplicación de matrices grande donde los datos caben en caché y los Tensor Cores están completamente utilizados.
• Limitado por memoria (memory-bound): el ancho de banda es el cuello de botella. Agregar más núcleos no ayudaría — están esperando datos. Ejemplo: operaciones elemento a elemento como ReLU o normalización de capa, donde cada elemento requiere muy poco cómputo pero debe cargarse desde y escribirse de vuelta a HBM.

La mayoría de las operaciones de deep learning están limitadas por memoria. Esto es contraintuitivo — ¡las GPUs tienen miles de núcleos! Pero las GPUs modernas tienen tanta capacidad de cómputo (cientos de TFLOPS) que solo las operaciones con mayor intensidad aritmética (multiplicaciones de matrices grandes, convoluciones) pueden mantener ocupados los núcleos.

Intensidad Aritmética

¿Cómo determinas si un kernel está limitado por cómputo o por memoria? Calcula la intensidad aritmética :

FLOPs es el número de operaciones de punto flotante que realiza el kernel. Bytes transferidos es el total de datos movidos entre la memoria de la GPU y el SM (lecturas + escrituras). La razón $I$ nos dice cuánto cómputo hacemos por byte de dato movido, medido en FLOPs/byte.

Consideremos dos extremos:

• $I$ baja (p. ej., 0.5 FLOPs/byte): muy poco cómputo por byte. El kernel pasa la mayor parte del tiempo esperando datos. Ejemplo: suma de vectores — 1 suma por 12 bytes cargados (4 bytes por cada una de 2 entradas + 4 bytes escritos = 12 bytes, 1 FLOP). $I = 1/12 \approx 0.08$ .
• $I$ alta (p. ej., 100 FLOPs/byte): mucho cómputo por byte. El kernel mantiene los núcleos ocupados. Ejemplo: multiplicación de matrices grande — para matrices de $n \times n$ , $2n^3$ FLOPs pero solo $3 \times 4n^2$ bytes (leer dos matrices + escribir una, en FP32). $I = 2n^3 / 12n^2 = n/6$ . Para $n = 4096$ , $I \approx 683$ FLOPs/byte.

Calculemos la intensidad aritmética para varias operaciones comunes:

import numpy as np

n = 4096
elem = 4  # bytes per float32
vec_n = 1_000_000
seq, d = 2048, 128

# Collect all row data first
rows = [
    ("Vector add (1M)",            vec_n,              3 * vec_n * elem),
    ("ReLU (1M)",                  vec_n,              2 * vec_n * elem),
    ("LayerNorm (1M)",             5 * vec_n,          2 * vec_n * elem),
    (f"Matmul ({n}\u00d7{n})",    2 * n**3,           3 * n**2 * elem),
    (f"Attention QK^T ({seq}\u00d7{d})", 2*seq**2*d,  (2*seq*d + seq**2) * elem),
]

# Compute max widths per column
w_name  = max(len(r[0]) for r in rows)
w_flops = max(len(f"{r[1]:,.0f}") for r in rows)
w_bytes = max(len(f"{r[2]:,.0f}") for r in rows)
w_ratio = max(len(f"{r[1]/r[2]:.1f}") for r in rows)

print("Arithmetic Intensity Examples (FP32):")
print()
for name, flops, bts in rows:
    I = flops / bts
    print(f"  {name:<{w_name}} {flops:>{w_flops},.0f} FLOPs / {bts:>{w_bytes},.0f} bytes = {I:>{w_ratio}.1f} FLOPs/byte")

print()
print("Low intensity → memory-bound (waiting for data)")
print("High intensity → compute-bound (cores are the limit)")

El Modelo Roofline

El modelo roofline (Williams et al., 2008) es una herramienta visual que muestra, de un vistazo, si un kernel está limitado por cómputo o por memoria. Es posiblemente el modelo de estimación rápida más útil para el análisis de rendimiento en GPU (consulta el Modal GPU Glossary para un tratamiento interactivo).

El modelo grafica dos límites de hardware:

• Techo de cómputo (línea horizontal): el rendimiento pico de cómputo de la GPU en FLOPS/s (p. ej., ~990 TFLOPS FP16 para H100 SXM).
• Pendiente del ancho de banda de memoria (línea diagonal): el ancho de banda pico de memoria de la GPU en bytes/s (p. ej., ~3.35 TB/s para H100 HBM3).

El rendimiento alcanzable de un kernel está dado por:

Esta fórmula captura los dos regímenes:

• Si $I \times \text{Peak BW} < \text{Peak FLOPS/s}$ : el kernel está limitado por memoria . El rendimiento está limitado por la velocidad a la que llegan los datos, no por el cómputo. Aumentar $I$ (p. ej., fusionando operaciones) ayudaría.
• Si $I \times \text{Peak BW} \geq \text{Peak FLOPS/s}$ : el kernel está limitado por cómputo . Ya está saturando los núcleos. Solo hardware más rápido o mejoras algorítmicas ayudan.

El punto de cresta (ridge point) es la intensidad aritmética donde los dos techos se encuentran: $I_{\text{ridge}} = \text{Peak FLOPS/s} / \text{Peak BW}$ . Los kernels por debajo de esta intensidad están limitados por memoria; por encima, limitados por cómputo.

Calculemos el punto de cresta para un H100 y clasifiquemos varias operaciones comunes:

import json, js

# H100 SXM specs (approximate)
peak_flops = 990e12    # ~990 TFLOPS (FP16 Tensor Core)
peak_bw = 3.35e12      # ~3.35 TB/s (HBM3)

ridge_point = peak_flops / peak_bw
print(f"H100 SXM Roofline:")
print(f"  Peak compute:  {peak_flops/1e12:.0f} TFLOPS (FP16)")
print(f"  Peak bandwidth: {peak_bw/1e12:.2f} TB/s")
print(f"  Ridge point:    {ridge_point:.0f} FLOPs/byte")
print()

# Classify operations
ops = [
    ("Vector add",      0.08),
    ("ReLU",            0.12),
    ("LayerNorm",       0.62),
    ("Softmax",         1.5),
    ("Attention QK^T",  64),
    ("Matmul 4096²",    683),
]
rows = []
for name, intensity in ops:
    achievable = min(peak_flops, intensity * peak_bw)
    pct = achievable / peak_flops * 100
    bound = "memory" if intensity < ridge_point else "compute"
    rows.append([name, f"{intensity:.1f}", bound, f"{pct:.1f}%"])

js.window.py_table_data = json.dumps({
    "headers": ["Operation", "Intensity (FLOPs/byte)", "Bottleneck", "% of Peak"],
    "rows": rows
})

print("Most DL ops (element-wise, normalization, softmax) are memory-bound.")
print("Only large matmuls and convolutions cross the ridge point.")

Ley de Little: El Paralelismo Oculta la Latencia

La Ley de Little, tomada de la teoría de colas, proporciona la relación fundamental entre rendimiento (throughput), paralelismo y latencia:

Para lograr un throughput dado, necesitas suficiente paralelismo para mantener el pipeline lleno mientras las operaciones individuales se completan. Si una solicitud de memoria toma 300 ciclos de reloj (latencia) y quieres emitir una solicitud por ciclo (throughput = 1/ciclo), necesitas 300 solicitudes en vuelo simultáneamente (paralelismo = 300).

Esta es exactamente la razón por la que las GPUs necesitan miles de hilos concurrentes: ocultan la latencia de memoria de ~300 ciclos teniendo suficientes warps listos para ejecutarse mientras otros esperan. Si el SM tiene solo unos pocos warps activos, se detiene frecuentemente (baja ocupación ). Con muchos warps, siempre hay uno listo para ejecutarse (alta ocupación = buena ocultación de latencia).

import numpy as np

mem_latency_cycles = 300  # typical HBM latency
clock_ghz = 1.8           # H100 SM clock

# How many warps needed to hide memory latency?
instructions_per_warp_per_cycle = 1
warps_needed = mem_latency_cycles * instructions_per_warp_per_cycle

print("Little's Law on a GPU SM:")
print(f"  Memory latency:   {mem_latency_cycles} cycles")
print(f"  To hide latency:  need {warps_needed} warps in flight")
print(f"  Max warps per SM: 64 (H100)")
print(f"  Minimum occupancy for full hiding: {warps_needed/64:.0%}")
print()
print("This is why occupancy matters: more active warps = more")
print("choices for the scheduler = fewer stall cycles.")
print()

# Collect rows for aligned output
rows = []
for occupancy_pct in [25, 50, 75, 100]:
    active_warps = int(64 * occupancy_pct / 100)
    hiding_ratio = min(1.0, active_warps / warps_needed)
    stall_pct = (1 - hiding_ratio) * 100
    rows.append((occupancy_pct, active_warps, stall_pct))

w_stall = max(len(f"{r[2]:.0f}") for r in rows)

for occ, warps, stall in rows:
    print(f"  {occ:>3}% occupancy ({warps:>2} warps): "
          f"~{stall:>{w_stall}.0f}% potential stall cycles")

Cuellos de Botella Comunes

A continuación se presenta una lista de verificación de los problemas de rendimiento más comunes en GPUs. Cada uno representa una forma diferente en que los recursos de hardware quedan ociosos o subutilizados.

• Presión de registros: el kernel usa demasiados registros por hilo, lo que significa que menos hilos activos caben en el SM, lo que reduce la ocupación, lo que disminuye la capacidad del SM para ocultar la latencia de memoria. Solución: simplificar el kernel, usar __launch_bounds__ para orientar al compilador, o aceptar el compromiso si cada hilo realiza suficiente cómputo para compensar.
• Conflictos de bancos: cuando múltiples hilos en un warp acceden al mismo banco de memoria compartida simultáneamente, los accesos se serializan. Un conflicto de bancos de 32 vías (todos los hilos accediendo al mismo banco) hace la memoria compartida 32× más lenta. Solución: rellenar los arreglos de memoria compartida o reestructurar los patrones de acceso para que hilos adyacentes accedan a bancos diferentes.
• Divergencia de warps: los hilos en un warp toman diferentes ramas, causando que ambos caminos se ejecuten secuencialmente con los hilos inactivos enmascarados. Solución: reestructurar el código para que los hilos en el mismo warp tomen el mismo camino (p. ej., ordenar datos por categoría antes de procesarlos).
• Acceso a memoria no coalescido: los hilos acceden a direcciones dispersas de memoria global, generando muchas transacciones pequeñas en lugar de pocas grandes. Solución: usar Structure of Arrays (SoA) en lugar de Array of Structures (AoS), y asegurar que hilos adyacentes accedan a direcciones de memoria adyacentes.
• Sobrecarga de lanzamiento de kernels: cada lanzamiento de kernel tiene aproximadamente 5–10 µs de sobrecarga (CPU → driver → planificación de GPU). Para kernels diminutos, esta sobrecarga domina el tiempo real de cómputo. Solución: fusionar kernels ( torch.compile hace esto automáticamente) o usar CUDA Graphs para agrupar lanzamientos en una sola submisión.
• Transferencia host-dispositivo: copiar datos entre CPU y GPU ( cudaMemcpy ) pasa por PCIe (~32 GB/s para Gen4 x16), que es órdenes de magnitud más lento que el ancho de banda de memoria de la GPU (~3 TB/s). Solución: minimizar transferencias, usar memoria fijada (pinned memory), y solapar cómputo con transferencias usando streams de CUDA.

Quiz

Pon a prueba tu comprensión del análisis de rendimiento en GPU — intensidad aritmética, el modelo roofline, ocultación de latencia y cuellos de botella comunes.