Dentro del Streaming Multiprocessor

Anatomía de un SM

Un Streaming Multiprocessor (SM) es la unidad de cómputo fundamental de toda GPU de NVIDIA. El artículo 1 presentó el SM a alto nivel — un procesador autónomo con sus propias unidades de ejecución, registros y memoria local. Ahora abrimos la tapa y miramos lo que realmente hay dentro.

Un SM contiene varios tipos de unidades de ejecución, cada una especializada en una clase diferente de operaciones. El diagrama a continuación muestra la distribución de un SM en la arquitectura H100 (Hopper) . Los números exactos varían entre generaciones de GPU — arquitecturas anteriores tienen menos CUDA Cores, tamaños de archivo de registros diferentes, o carecen por completo de Tensor Cores — pero los principios estructurales se mantienen notablemente consistentes.

┌─────────────────────────────────────────────────────────────┐
│                    Streaming Multiprocessor                   │
├──────────────┬──────────────┬──────────────┬────────────────┤
│  Warp Sched. │  Warp Sched. │  Warp Sched. │  Warp Sched.   │
│  + Dispatch  │  + Dispatch  │  + Dispatch  │  + Dispatch    │
├──────────────┴──────────────┴──────────────┴────────────────┤
│                                                              │
│  ┌────────────┐ ┌────────────┐ ┌─────┐ ┌─────┐ ┌────────┐  │
│  │ 32 CUDA    │ │ 32 CUDA    │ │ SFU │ │ SFU │ │ LSU ×8 │  │
│  │ Cores (FP32)│ │ Cores (FP32)│ │     │ │     │ │        │  │
│  └────────────┘ └────────────┘ └─────┘ └─────┘ └────────┘  │
│  ┌────────────┐ ┌────────────┐ ┌─────┐ ┌─────┐ ┌────────┐  │
│  │ 32 CUDA    │ │ 32 CUDA    │ │ SFU │ │ SFU │ │ LSU ×8 │  │
│  │ Cores (FP32)│ │ Cores (FP32)│ │     │ │     │ │        │  │
│  └────────────┘ └────────────┘ └─────┘ └─────┘ └────────┘  │
│                                                              │
│  ┌─────────────┐ ┌─────────────┐ ┌─────────────┐            │
│  │ Tensor Core │ │ Tensor Core │ │ Tensor Core │ ...×4      │
│  │ (MMA 4×4)   │ │ (MMA 4×4)   │ │ (MMA 4×4)   │            │
│  └─────────────┘ └─────────────┘ └─────────────┘            │
│                                                              │
│  ┌──────────────────────────────────────────────────────┐    │
│  │              256 KB Register File                     │    │
│  │         (65,536 × 32-bit registers)                   │    │
│  └──────────────────────────────────────────────────────┘    │
│  ┌──────────────────────────────────────────────────────┐    │
│  │         256 KB L1 Data Cache / Shared Memory          │    │
│  │      (configurable split, e.g., 192 KB / 64 KB)      │    │
│  └──────────────────────────────────────────────────────┘    │
└──────────────────────────────────────────────────────────────┘

Recorramos cada componente:

• Warp Schedulers (×4): cada planificador selecciona un warp listo por ciclo y despacha su siguiente instrucción a la unidad de ejecución apropiada. Cuatro planificadores operando en paralelo es lo que permite al SM mantener alimentados múltiples pipelines de ejecución simultáneamente.
• CUDA Cores (128 en total, 4 grupos de 32): unidades de ejecución de propósito general para punto flotante y enteros. Cada uno realiza una operación multiply-add en FP32 por ciclo de reloj. Son los caballos de batalla del cómputo general.
• Tensor Cores (×4): unidades especializadas de multiplicación-acumulación de matrices. Cada una computa una multiplicación-suma de matrices $4 \times 4$ en un solo ciclo — aproximadamente 128 multiply-adds donde un CUDA Core realiza solo uno.
• SFUs — Unidades de Funciones Especiales (×4): hardware dedicado para matemáticas trascendentales (sin, cos, exp, log, raíz cuadrada recíproca). Estas operaciones requerirían muchos ciclos en ALUs de propósito general.
• LSUs — Unidades de Carga/Almacenamiento (×16): gestionan las solicitudes de memoria hacia memoria global (HBM) y memoria compartida. La coalescencia de memoria — combinar muchas solicitudes a nivel de hilo en menos transacciones amplias — ocurre en esta etapa.
• Archivo de Registros (256 KB): la memoria más rápida disponible, con acceso de latencia cero. Privada para cada hilo, dividida entre todos los hilos activos en el SM.
• Caché L1 / Memoria Compartida (256 KB): SRAM rápida on-chip con una división configurable entre caché L1 gestionada por hardware y memoria compartida gestionada por el programador.

El resto de este artículo examina cada componente en mayor detalle, comenzando por los CUDA Cores.

CUDA Cores: Los Caballos de Batalla

Los CUDA Cores son las unidades de ejecución de propósito general para punto flotante y enteros dentro del SM. Cada uno puede realizar una operación multiply-add en FP32 por ciclo de reloj. Un SM del H100 tiene 128 CUDA Cores organizados en 4 grupos de 32, lo que le da 128 operaciones FP32 por ciclo.

¿Por qué 32 por grupo? Porque un warp es de 32 hilos, y cada planificador de warps despacha una instrucción a 32 CUDA Cores simultáneamente — un core por hilo. Los 4 grupos se corresponden con los 4 planificadores de warps: cada planificador puede emitir una instrucción de warp por ciclo, y los 4 operan en paralelo. Este es el vínculo fundamental entre el modelo de ejecución SIMT y el hardware físico.

Los CUDA Cores manejan una amplia gama de operaciones: suma, multiplicación, multiply-add fusionado (FMA), aritmética de enteros, comparaciones y operaciones de bits. Si tu kernel está haciendo matemáticas escalares — sumando vectores, computando funciones elemento a elemento, ejecutando flujo de control — esas operaciones probablemente están fluyendo a través de los CUDA Cores. Son la ruta de ejecución por defecto para la mayoría de las cargas de trabajo en GPU.

(Ver también "CUDA Core" en el Glosario de GPU de Modal para una definición concisa a nivel de hardware.)

Tensor Cores: Máquinas de Multiplicación de Matrices

Los Tensor Cores son hardware especializado para una operación específica: multiplicación-acumulación de matrices pequeñas (MMA). En un solo ciclo, un Tensor Core computa:

donde $A$ es $4 \times 4$, $B$ es $4 \times 4$, $C$ es $4 \times 4$ (el acumulador), y $D$ es el resultado $4 \times 4$. Eso es aproximadamente 128 operaciones multiply-add en un solo ciclo — comparado con un CUDA Core que realiza exactamente 1.

Desglosemos cada parte de esa ecuación:

• $A$ y $B$: las matrices de entrada, típicamente almacenadas en FP16 o BF16 por eficiencia. Menor precisión significa que cada elemento ocupa la mitad del tamaño de FP32, así que caben el doble de valores en el mismo espacio de registros, y los circuitos de multiplicación son más pequeños y rápidos.
• $C$: el acumulador, a menudo almacenado en FP32 para mayor precisión. Se multiplica en baja precisión pero se acumula en alta precisión — esta es la base del entrenamiento de precisión mixta . La idea es que los errores individuales de multiplicación en FP16 son pequeños, pero se acumulan a lo largo de miles de sumas. Acumular en FP32 mantiene la suma parcial precisa.
• $D$: el resultado, que sobrescribe $C$ para la siguiente iteración de bloques. En una multiplicación de matrices grande, $C$ acumula resultados de muchas multiplicaciones de bloques antes de escribirse de vuelta.

¿Por qué $4 \times 4$? Parece ser un punto óptimo de diseño de hardware — lo suficientemente pequeño para implementarlo eficientemente en silicio (el número de cables y unidades de multiplicación escala con el cuadrado de la dimensión), lo suficientemente grande para que dividir una multiplicación de matrices grande en bloques de $4 \times 4$ amortice la sobrecarga de control. En la práctica, el programador rara vez piensa directamente en bloques de $4 \times 4$. Bibliotecas como cuBLAS y Triton dividen matrices grandes en bloques (a menudo de $16 \times 16$ o más) y mapean cada bloque a secuencias de instrucciones MMA de Tensor Core.

Un SM del H100 tiene 4 Tensor Cores. En precisión FP16, cada uno realiza aproximadamente 128 multiply-adds por ciclo, por lo que el SM entrega:

A lo largo del GPU H100 completo con 132 SMs a aproximadamente 1.8 GHz:

Por esto los Tensor Cores son tan importantes para el aprendizaje profundo. La multiplicación de matrices domina el cómputo tanto en entrenamiento como en inferencia (atención, capas lineales, convoluciones expresadas como matmuls), y los Tensor Cores la aceleran en aproximadamente dos órdenes de magnitud sobre los CUDA Cores solos.

La simulación a continuación demuestra la operación MMA de $4 \times 4$. Usamos entradas en FP16 y un acumulador en FP32, reflejando lo que el hardware realmente hace:

import numpy as np

# Tensor Core operation: D = A @ B + C (4×4 matrices)
A = np.array([[1, 0, 2, 1],
              [0, 1, 1, 0],
              [2, 0, 1, 1],
              [1, 1, 0, 2]], dtype=np.float16)

B = np.array([[1, 2, 0, 1],
              [0, 1, 1, 0],
              [1, 0, 2, 1],
              [2, 1, 0, 1]], dtype=np.float16)

C = np.zeros((4, 4), dtype=np.float32)  # accumulator in FP32

# One Tensor Core cycle: D = A @ B + C
D = (A.astype(np.float32) @ B.astype(np.float32)) + C

print("Tensor Core MMA: D = A × B + C")
print(f"\nA (FP16):\n{A}")
print(f"\nB (FP16):\n{B}")
print(f"\nC (FP32 accumulator):\n{C}")
print(f"\nD = A @ B + C (FP32 result):\n{D}")
print(f"\nOperations in one cycle: {4*4*(2*4-1)} multiply-adds")
print(f"A CUDA Core does 1 per cycle. A Tensor Core does ~128.")

(Ver también "Tensor Core" en el Glosario de GPU de Modal para más detalles sobre cómo estas unidades se mapean a las instrucciones MMA a nivel de warp.)

Unidades de Funciones Especiales y Unidades de Carga/Almacenamiento

Dos tipos más de unidades de ejecución completan el conjunto computacional del SM.

Las SFUs (Unidades de Funciones Especiales) manejan matemáticas trascendentales — sin , cos , exp , log y raíz cuadrada recíproca ( rsqrt ). Computar estas funciones con multiply-adds simples requeriría aproximaciones iterativas (series de Taylor, algoritmos CORDIC) que toman muchos ciclos. El hardware dedicado de las SFUs las maneja en significativamente menos ciclos — típicamente alrededor de 8 ciclos por operación, comparado con docenas para una implementación por software. Un SM típicamente tiene 4 SFUs.

¿Dónde aparecen las SFUs en el aprendizaje profundo? Las funciones de activación como GELU y Swish involucran exp y tanh ; softmax requiere exp a lo largo de toda una fila de logits; la normalización de capa usa rsqrt . Estas operaciones no son el cuello de botella en la mayoría de los modelos (las matmuls dominan), pero sin SFUs serían notablemente más lentas.

Las LSUs (Unidades de Carga/Almacenamiento) gestionan las solicitudes de memoria. Cuando un hilo necesita datos de memoria global (HBM de la GPU) o memoria compartida, una LSU emite la solicitud y gestiona la transacción. Un SM tiene aproximadamente 32 LSUs, permitiéndole manejar múltiples solicitudes de memoria por ciclo.

La coalescencia de memoria ocurre a este nivel: cuando muchos hilos en un warp solicitan direcciones que caen dentro de la misma línea de caché de 128 bytes, las LSUs pueden combinarlas en una sola transacción de memoria amplia en lugar de emitir 32 solicitudes separadas. Esta es una de las consideraciones de rendimiento más importantes en la programación CUDA — los patrones de acceso coalescentes pueden ser un orden de magnitud más rápidos que los dispersos. Volveremos a la coalescencia en detalle en un artículo posterior sobre patrones de acceso a memoria.

El Archivo de Registros: La Memoria Más Rápida

El archivo de registros es la memoria más rápida disponible para los hilos — sin latencia adicional, accesible cada ciclo, y privada para cada hilo. Un SM del H100 tiene 256 KB de registros, lo que equivale a 65,536 registros individuales de 32 bits. Es una cantidad sustancial de almacenamiento — mayor que toda la caché L1 de la mayoría de los cores de CPU — y necesita serlo, porque se comparte entre todos los hilos activos en el SM.

Esto crea un compromiso fundamental. Si cada hilo usa muchos registros (algo común en kernels complejos que mantienen resultados intermedios a mano), entonces menos hilos pueden estar activos simultáneamente, reduciendo la ocupancia — la proporción de warps activos respecto al máximo que el SM soporta. Menor ocupancia significa que los planificadores de warps tienen menos warps para elegir cuando uno se detiene por una solicitud de memoria, reduciendo la capacidad del SM para ocultar latencia alternando entre warps.

Esto es la presión de registros — uno de los precipicios de rendimiento más sutiles en la programación de GPU. Un kernel que usa 128 registros por hilo podría ejecutarse al 25% de ocupancia, mientras que uno que usa 24 registros por hilo puede llenar el SM por completo. Si la ocupancia reducida realmente perjudica el rendimiento depende de si el kernel está limitado por cómputo o por memoria. Exploraremos este compromiso en profundidad en el artículo 5.

La siguiente simulación muestra cómo el presupuesto de registros limita los hilos activos:

total_registers = 65536
max_threads = 2048

# Collect rows for aligned output
scenarios = [
    ("Simple kernel",  24),
    ("Complex kernel", 128),
]
rows = []
for name, regs in scenarios:
    threads = min(total_registers // regs, max_threads)
    warps = threads // 32
    occupancy = threads / max_threads
    rows.append((name, regs, threads, warps, occupancy))

# Compute column widths
w_name = max(len(r[0]) for r in rows)
w_regs = max(len(str(r[1])) for r in rows)
w_thr  = max(len(str(r[2])) for r in rows)
w_wrp  = max(len(str(r[3])) for r in rows)
w_occ  = max(len(f"{r[4]:.0%}") for r in rows)

print("Register budget tradeoff:")
print(f"  Total registers: {total_registers}")
print(f"  Max threads/SM:  {max_threads}")
print()
for name, regs, threads, warps, occ in rows:
    print(f"  {name:<{w_name}} ({regs:>{w_regs}} regs/thread): "
          f"{threads:>{w_thr}} threads, {warps:>{w_wrp}} warps, {f'{occ:.0%}':>{w_occ}} occupancy")
print()
print("Fewer active warps → less latency hiding → potential stalls.")
print("This is the register pressure problem (article 5).")

Observa la diferencia dramática: 100% de ocupancia vs. 25%. El kernel simple da a los planificadores de warps 64 warps para ciclar, mientras que el kernel complejo les da solo 16. Si esos 16 warps son suficientes para mantener ocupadas las unidades de ejecución depende de la intensidad aritmética del kernel — un tema que formalizaremos con el modelo roofline más adelante en el track.

Caché L1 y Memoria Compartida

Cada SM tiene un total combinado de 256 KB de SRAM rápida on-chip que cumple un doble propósito:

• Caché de Datos L1: cachea automáticamente datos obtenidos de memoria global (HBM). El hardware gestiona qué se cachea y qué se desaloja — el programador no lo controla directamente. Esto ayuda con patrones de acceso irregulares donde hilos en diferentes warps tocan la misma región de memoria.
• Memoria Compartida: memoria explícitamente gestionada por el programador que se comparte entre todos los hilos de un bloque de hilos (no de todo el SM — solo dentro de un único bloque). Los hilos pueden leer, escribir y sincronizarse a través de la memoria compartida, habilitando algoritmos cooperativos como la multiplicación de matrices por bloques y las reducciones paralelas.

La división entre L1 y memoria compartida es configurable. En el H100, los 256 KB pueden dividirse como, por ejemplo, 192 KB de memoria compartida / 64 KB de L1, o desplazarse en la otra dirección según las necesidades del kernel. El runtime de CUDA proporciona una API ( cudaFuncSetAttribute ) para solicitar una división preferida antes de lanzar un kernel.

La regla general es directa: kernels que necesitan cooperación entre hilos — matmul por bloques, reducciones, cómputos de stencil — se benefician de asignar más memoria compartida. Kernels con patrones de acceso irregulares y dependientes de datos se benefician de una caché L1 más grande. La mayoría de los kernels de aprendizaje profundo caen en la primera categoría, razón por la cual la memoria compartida es tan central para la programación de GPU de alto rendimiento.

(Ver también "Shared Memory" en el Glosario de GPU de Modal para más información sobre conflictos de banco y patrones de acceso óptimos.)

Planificadores de Warps: Manteniendo el SM Ocupado

Cada SM tiene 4 planificadores de warps. En cada ciclo de reloj, cada planificador puede:

• 1. Seleccionar uno de sus warps asignados que esté listo para ejecutarse — es decir, que no esté esperando datos de memoria, no esté detenido en una barrera de sincronización, y tenga su siguiente instrucción decodificada y lista.
• 2. Emitir una instrucción de ese warp a las unidades de ejecución apropiadas — CUDA Cores para aritmética, Tensor Cores para MMA, SFUs para funciones trascendentales, o LSUs para operaciones de memoria.

Con 4 planificadores operando en paralelo, un SM puede emitir hasta 4 instrucciones de warp por ciclo — potencialmente a diferentes unidades de ejecución. Por ejemplo, en un solo ciclo, un planificador podría enviar un warp a los CUDA Cores para un FMA, otro podría despachar un warp a un Tensor Core para un MMA, un tercero podría emitir una carga de memoria a través de las LSUs, y el cuarto podría enviar un warp a una SFU para un cómputo de exp . Este paralelismo entre tipos de unidades de ejecución es una fuente clave de throughput.

Este es el mecanismo central detrás de la ocultación de latencia . Cuando el warp A emite una carga de memoria y debe esperar cientos de ciclos hasta que los datos lleguen desde HBM, el planificador no se queda inactivo — selecciona el warp B, que está listo para ejecutarse. Cuando el warp B se detiene en una barrera __syncthreads() , selecciona el warp C. Mientras haya suficientes warps listos para alternar, las unidades de ejecución se mantienen ocupadas y el SM desperdicia muy pocos ciclos.

Por eso la ocupancia importa . Más warps activos significa más candidatos para que el planificador elija, y menos ciclos donde todos los warps están detenidos simultáneamente. Un SM completamente ocupado con 64 warps puede tolerar latencias de memoria largas porque casi siempre hay al menos un warp listo. Un SM con solo 4 warps activos tiene mucho menos margen de maniobra — si los 4 están esperando memoria, las unidades de ejecución quedan completamente inactivas hasta que llegan los datos.

Dicho esto, mayor ocupancia no siempre se traduce en mayor rendimiento. Un kernel limitado por cómputo que mantiene las unidades de ejecución ocupadas con aritmética puede rendir igual de bien al 50% de ocupancia que al 100%. El beneficio de alta ocupancia es específicamente ocultar latencia de memoria — y solo ayuda cuando hay latencia que ocultar.

(Ver también "Warp Scheduler" en el Glosario de GPU de Modal para más información sobre políticas de planificación y capacidades de emisión dual.)

Quiz

Pon a prueba tu comprensión de la arquitectura interna del Streaming Multiprocessor.