Hilos, Warps y el Modelo de Programación CUDA

La Jerarquía de Hilos

CUDA organiza el trabajo paralelo en una jerarquía de tres niveles: hilos → bloques de hilos → grids . Comprender esta jerarquía es esencial porque se mapea directamente al hardware: los hilos se ejecutan en CUDA Cores, los bloques de hilos se asignan a Streaming Multiprocessors (SMs) y un grid abarca la GPU completa. Si el mapeo parece mecánico, es intencional — el modelo de programación de CUDA es una abstracción delgada sobre el silicio, y mantenerlo cerca del metal es lo que hace rápidas a las GPUs.

Grid (la GPU completa)
├── Block (0,0)              ← se asigna a un SM
│   ├── Warp 0 (hilos 0-31)      ← 32 hilos ejecutan en sincronía
│   ├── Warp 1 (hilos 32-63)
│   └── ...
├── Block (0,1)              ← se asigna a otro SM
│   ├── Warp 0
│   └── ...
└── Block (1,0)
    └── ...

Desglosemos cada nivel de abajo hacia arriba.

Hilo (Thread): la unidad de ejecución más pequeña. Cada hilo tiene sus propios registros y contador de programa, pero ejecuta el mismo código que todos los demás hilos — solo que con datos diferentes. Este es el modelo SIMT (Single Instruction, Multiple Threads) de NVIDIA: un flujo de instrucciones, múltiples flujos de datos. Si conoces SIMD en CPUs, SIMT es un primo cercano, pero con la diferencia crucial de que cada hilo tiene su propio flujo de control (más sobre esto en la sección de divergencia).

Warp: un grupo de exactamente 32 hilos que se ejecutan en sincronía. El hardware siempre planifica y ejecuta warps, nunca hilos individuales. Cuando un planificador de warps emite una instrucción, los 32 hilos del warp la ejecutan simultáneamente. Puedes pensar en el warp como la verdadera unidad atómica de ejecución en GPU — todo por debajo es una conveniencia de programación, y todo por encima es una herramienta organizativa.

Bloque de Hilos (Cooperative Thread Array / CTA): un grupo de warps — hasta 1,024 hilos en la mayoría de las arquitecturas — que se ejecutan en el mismo SM. Los hilos dentro de un bloque pueden cooperar mediante memoria compartida y barreras de sincronización ( __syncthreads() ). Este es el nivel donde ocurre la comunicación entre hilos. Los hilos en bloques diferentes no pueden comunicarse directamente — son independientes por diseño, lo que permite al planificador de la GPU distribuir bloques libremente entre los SMs.

Grid: la colección de todos los bloques de hilos lanzados por una sola llamada al kernel. El planificador de bloques de la GPU distribuye los bloques entre los SMs disponibles. Un grid puede contener millones de hilos, y el hardware se encarga de toda la planificación — el programador simplemente declara las dimensiones del grid y deja que la GPU determine cuándo y dónde se ejecuta cada bloque.

¿Por Qué 32? El Warp como Unidad de Ejecución

El número 32 está grabado en el hardware de las GPUs de NVIDIA. Cada instrucción se emite a exactamente 32 hilos simultáneamente — ni más, ni menos. Esta constante aparentemente arbitraria tiene implicaciones profundas para cómo escribes y optimizas código de GPU.

Cuando lanzas un kernel con, digamos, 100 hilos en un bloque, el hardware no ejecuta exactamente 100 hilos. Redondea hacia arriba al siguiente múltiplo de 32 — en este caso, 4 warps (128 hilos). Los últimos 28 hilos en el warp 3 se enmascaran : pasan por los movimientos de ejecutar cada instrucción, pero sus resultados se descartan. Esto significa que lanzar 33 hilos desperdicia casi tantos recursos como lanzar 64 — un hecho que sorprende a muchos recién llegados a la programación GPU.

Cuantifiquemos el desperdicio para diferentes cantidades de hilos:

import numpy as np

def warps_needed(num_threads):
    return (num_threads + 31) // 32  # ceiling division by 32

# Collect all row data first
thread_counts = [1, 32, 33, 64, 100, 128, 256, 1000]
rows = []
for n in thread_counts:
    warps = warps_needed(n)
    total = warps * 32
    wasted = total - n
    efficiency = n / total * 100
    rows.append((n, warps, total, wasted, efficiency))

# Compute max width per column for alignment
w_eff = max(len(f"{r[4]:.0f}") for r in rows)

# Print aligned table
for n, warps, total, wasted, eff in rows:
    print(f"  {n:>4} threads → {warps:>2} warps ({total:>4} slots, {wasted:>3} wasted, {eff:>{w_eff}.0f}% efficient)")

print()
print("Takeaway: always use multiples of 32 for thread counts.")
print("33 threads wastes 46% of warp capacity!")

La regla práctica es simple: elige siempre cantidades de hilos que sean múltiplos de 32 . Las opciones comunes son 128, 256 o 512 hilos por bloque. Ir más alto (hasta 1,024) aumenta la ocupación pero también aumenta la presión de registros, así que la elección óptima depende del uso de recursos de tu kernel.

Divergencia de Warps

Ya que los 32 hilos en un warp ejecutan la misma instrucción simultáneamente, ¿qué pasa cuando diferentes hilos necesitan tomar diferentes ramas? Esta situación — llamada divergencia de warps — es una de las trampas de rendimiento más importantes en la programación GPU.

Considera este pseudocódigo para un kernel CUDA:

# CUDA kernel (pseudocode):
if threadIdx.x % 2 == 0:
    result = expensive_path(data)    # even threads
else:
    result = cheap_path(data)        # odd threads

El hardware no puede ejecutar ambas ramas simultáneamente. En cambio, las serializa : primero, los 32 hilos ejecutan expensive_path (con los hilos impares enmascarados), luego los 32 ejecutan cheap_path (con los hilos pares enmascarados). El tiempo de ejecución es la suma de ambas ramas, no el máximo — efectivamente duplicando el costo comparado con un warp donde todos los hilos toman el mismo camino.

import numpy as np

# Simulate 32 threads in a warp
warp_size = 32
thread_ids = np.arange(warp_size)

# Divergent: even/odd threads take different paths
even_mask = (thread_ids % 2 == 0)
odd_mask = ~even_mask

# Cost without divergence (all take same path)
cost_uniform = 1  # one pass through the branch

# Cost with divergence (serialised)
cost_divergent = 2  # pass 1: even threads, pass 2: odd threads

print(f"Warp of {warp_size} threads:")
print(f"  Even threads: {even_mask.sum()}, Odd threads: {odd_mask.sum()}")
print(f"  Without divergence: {cost_uniform} pass (all threads same branch)")
print(f"  With divergence:    {cost_divergent} passes (serialised branches)")
print(f"  Slowdown: {cost_divergent}×")
print()
print("Worst case: 32 threads, each taking a different branch → 32 serial passes")
print("Best practice: restructure code so all threads in a warp take the same path")

La divergencia es más dañina cuando las ramas están desbalanceadas (un camino es costoso, el otro barato) o cuando los hilos dentro del mismo warp frecuentemente toman caminos diferentes. Una estrategia de mitigación común es reestructurar los datos para que los hilos en el mismo warp sigan naturalmente el mismo flujo de control — por ejemplo, ordenar los elementos de trabajo por tipo antes de distribuirlos a los hilos.

La Jerarquía de Memoria

La memoria de la GPU es una jerarquía de niveles progresivamente más grandes pero más lentos, no muy diferente de la jerarquía de caché en CPUs — pero con algunas diferencias importantes. En una GPU, algunos niveles son explícitamente gestionados por el programador, lo que te da más control pero también más responsabilidad. De más rápida a más lenta:

┌─────────────────────────────────────────────────────────┐
│ Registros        │ ~0 ciclos │ Por hilo       │ ~256 KB │
│ (más rápido)     │           │ (privado)      │ por SM  │
├──────────────────┼───────────┼────────────────┼─────────┤
│ Memoria          │ ~5 ciclos │ Por bloque     │ ~164 KB │
│ Compartida / L1  │           │ (compartida)   │ por SM  │
├──────────────────┼───────────┼────────────────┼─────────┤
│ Caché L2         │ ~30 cic.  │ Todos los SMs  │ ~50 MB  │
├──────────────────┼───────────┼────────────────┼─────────┤
│ Memoria Global   │ ~300 cic. │ Todos los SMs  │ 80 GB   │
│ (HBM / más lento)│           │ (RAM de GPU)   │ ~3 TB/s │
└─────────────────────────────────────────────────────────┘

Cada nivel sirve un propósito distinto:

• Registros: almacenamiento por hilo con latencia cero. El compilador asigna variables locales a registros automáticamente. Son la memoria más rápida del chip, pero son limitados — en la mayoría de las arquitecturas, cada SM tiene aproximadamente 65,536 registros de 32 bits compartidos entre todos sus hilos activos. Kernels complejos que necesitan muchas variables por hilo pueden causar presión de registros , forzando al compilador a derramar valores a memoria local más lenta.
• Memoria compartida: almacenamiento por bloque de baja latencia (~5 ciclos) gestionado explícitamente por el programador. Los hilos en el mismo bloque pueden leer, escribir y sincronizarse a través de la memoria compartida. Es el caballo de batalla de los algoritmos por tiles: cargar un tile de datos desde la memoria global a la memoria compartida, computar sobre él, y luego escribir los resultados de vuelta. En Hopper (H100), cada SM tiene hasta 228 KB de memoria compartida y caché L1 combinados.
• Caché L2: latencia media (~30 ciclos), compartida entre todos los SMs. Es gestionada por el hardware — el programador no la controla directamente. Actúa como un buffer entre los SMs y la memoria global, almacenando en caché los datos accedidos recientemente. El H100 tiene alrededor de 50 MB de L2.
• Memoria global (HBM): la latencia más alta (~300 ciclos de reloj), la mayor capacidad. Esta es la RAM principal de la GPU — 80 GB en el H100. Todos los datos empiezan y terminan aquí. El ancho de banda es impresionante (~3 TB/s en H100), pero la latencia es aproximadamente 60× mayor que la memoria compartida. La mayoría de los kernels están limitados por qué tan rápido pueden mover datos de HBM a las unidades de cómputo, no por el cómputo en sí.

La regla general es: mantener los datos lo más cerca posible de los núcleos . Mover datos de HBM a registros es, en muchas cargas de trabajo, el verdadero cuello de botella — no la aritmética. Esta es precisamente la razón por la que el modelo roofline (que exploraremos más adelante en este track) grafica el rendimiento contra la intensidad aritmética: si tu kernel hace muy poca matemática por byte cargado, estás limitado por memoria, y ninguna cantidad de núcleos más rápidos ayudará.

Coalescencia de Memoria

Cuando los hilos en un warp acceden a la memoria global, el hardware intenta coalescer sus solicitudes individuales en un número menor de transacciones de memoria amplias. Este es posiblemente el concepto de optimización más importante en la programación GPU, porque puede hacer la diferencia entre usar el 3% del ancho de banda disponible y el 100%.

Así funciona. Si los 32 hilos en un warp acceden a direcciones consecutivas de 4 bytes — el hilo 0 lee la dirección 0, el hilo 1 lee la dirección 4, el hilo 2 lee la dirección 8, y así sucesivamente hasta el hilo 31 leyendo la dirección 124 — el hardware combina las 32 solicitudes en una sola transacción de 128 bytes. Eso es 1 transacción en lugar de 32 , un ahorro masivo de ancho de banda.

Si los hilos acceden a direcciones dispersas o con stride, cada solicitud puede disparar una transacción separada, desperdiciando la mayoría de los bytes en cada una. Por eso los patrones de acceso a memoria importan enormemente en GPUs — el mismo algoritmo puede ejecutarse 10–30× más rápido simplemente reorganizando cómo los hilos acceden a la memoria.

import numpy as np

warp_size = 32
element_size = 4  # bytes (float32)
transaction_size = 128  # bytes (GPU memory transaction width)
elements_per_txn = transaction_size // element_size  # 32

# Coalesced: consecutive addresses
coalesced_addrs = np.arange(warp_size) * element_size
txns_coalesced = 1  # all fit in one 128-byte transaction

# Strided: every other element (stride = 2)
strided_addrs = np.arange(warp_size) * element_size * 2
txns_strided = 2  # spans 256 bytes → 2 transactions

# Random: scattered addresses
np.random.seed(42)
random_addrs = np.random.choice(10000, warp_size, replace=False) * element_size
# Worst case: up to 32 separate transactions
unique_txn_blocks = len(set(addr // transaction_size for addr in random_addrs))
txns_random = unique_txn_blocks

# Collect rows for aligned output
rows = [
    ("Coalesced (consecutive)", txns_coalesced, "baseline"),
    ("Strided (every other)",   txns_strided,   "slower"),
    ("Random (scattered)",      txns_random,    "slower"),
]
w_name = max(len(r[0]) for r in rows)
w_txn  = max(len(str(r[1])) for r in rows)
w_mult = max(len(f"{r[1]}") for r in rows)

print("Memory Coalescing: 32 threads accessing global memory")
print(f"  Transaction size: {transaction_size} bytes (fits {elements_per_txn} float32s)")
print()
for name, txns, label in rows:
    suffix = "transaction" if txns == 1 else "transactions"
    print(f"  {name:<{w_name}}  {txns:>{w_txn}} {suffix:<12} → {txns:>{w_mult}}× {label}")
print()
print("Coalesced access uses bandwidth efficiently.")
print("Scattered access wastes most of each transaction.")

La lección es clara: al diseñar estructuras de datos para computación en GPU, favorece Structure of Arrays (SoA) sobre Array of Structures (AoS) . Si cada hilo necesita un campo de un registro, el layout SoA significa que esos campos son contiguos en memoria y pueden coalescerse. El layout AoS intercala campos de diferentes registros, produciendo patrones de acceso con stride que desperdician ancho de banda.

Quiz

Pon a prueba tu comprensión del modelo de programación CUDA — jerarquías de hilos, warps, niveles de memoria y patrones de acceso.