English · Español

02 — Jerarquía de memoria en GPU¶

Jerarquía de memoria GPU: HBM (lenta, grande, ~1–3 TB/s, ~40–80 GiB) → L2 (compartida entre SMs, ~3–10 TB/s, ~40–100 MiB) → SMEM (manual, por bloque, ~10–20 TB/s, ~64–164 KiB) → registros (por hilo, ~muchos TB/s, ~256 KiB por SM repartidos entre todos los hilos residentes). Cada nivel es ~3–10× más rápido y ~10–100× más pequeño que el anterior.

Esta página da a la jerarquía de memoria de la GPU el mismo nivel de detalle que 02-memory-hierarchy.md de la Fase 1 dio a las caches de CPU — pero con las diferencias específicas de GPU señaladas explícitamente. Tras esta página deberías ser capaz de dibujar la jerarquía de memoria con los anchos de banda etiquetados.

La jerarquía, de arriba a abajo¶

Para una GPU de clase A100 (números del whitepaper de A100; representativos de "GPU moderna de centro de datos"):

Nivel	Tamaño	Ancho de banda	Latencia	Alcance	Gestionado por
HBM (DRAM)	40–80 GiB	1.5–3 TB/s	400–800 ciclos	Todo el dispositivo	Allocator (`cudaMalloc`)
L2 cache	40 MiB	~5 TB/s	200–250 ciclos	Compartida entre SMs	Hardware
SMEM (memoria compartida (shared memory))	164 KiB / SM	~19 TB/s	20–30 ciclos	Por bloque de hilos	Programador (`__shared__`)
L1 cache	compartida con SMEM (164 KiB)	~19 TB/s	20–30 ciclos	Por SM	Hardware (o repartición con SMEM)
Register file	256 KiB / SM	un registro / ciclo / lane	1 ciclo	Por hilo	Compilador

(Los números de H100 son ~30–50% mayores en todo; misma jerarquía.)

A modo de comparación, L1/L2/L3/DRAM de CPU son aproximadamente 32 KiB / 256 KiB / 30 MiB / 64 GiB, con anchos de banda ~1 TB/s / 200 GB/s / 70 GB/s / 20 GB/s en una CPU moderna. La HBM de la GPU es dos órdenes de magnitud más rápida que la DRAM de CPU. La L2 de la GPU es media magnitud más rápida que la L3 de CPU y 5× mayor. La GPU es una máquina de ancho de banda de memoria.

HBM: el nivel de memoria global¶

HBM (High Bandwidth Memory) es el equivalente en GPU de la "memoria principal". Todo lo que no cabe en cache vive aquí. Todas las reservas cudaMalloc salen de HBM. La KV cache (Fase 22) — cuando se porte a GPU en la Fase 24 — vive en HBM.

Tres cosas que saber:

HBM es rápida en ancho de banda, no en latencia. Una sola carga desde HBM tarda 400–800 ciclos. La latencia se oculta vía cambio de warp (theory/01). Para saturar el ancho de banda de HBM, necesitas muchas transacciones de memoria concurrentes en vuelo.
HBM es el limitador de tasa para kernels memory-bound. La intensidad de la decode-attention de la Fase 22 (~1 FLOP/byte fp16) está muy por debajo del machine balance (~150 FLOPs/byte en A100). Las FPUs de la GPU están ociosas esperando HBM. Mismo diagnóstico que en CPU, números más rápidos.
El ancho de banda de HBM es por dirección. Citado p.ej. como "2 TB/s" — eso es agregado bidireccional. Una carga de solo-lectura ve aproximadamente la mitad. Un read-modify-write ve menos.

L2: la cache compartida implícita¶

L2 está entre HBM y los SMs. Gestionada por hardware (como L2/L3 de CPU). Todos los SMs la comparten.

Propiedad clave: L2 es lo bastante grande como para contener una o dos capas de KV cache en tamaños de modelo modestos. Por eso "la tasa de aciertos en L2 para lecturas de cache" es algo que los sistemas de serving optimizan. Una cache bien dispuesta que relee las mismas filas K, V entre las capas de un mismo forward pass puede dejarlas en L2 para el acceso a partir de la segunda capa.

Nota práctica: no puedes reservar "en L2" — solo puedes reservar en HBM y esperar que L2 atrape un working set. Hay hints de acceso L2-persistente (cudaStreamAttribute*) para control avanzado. La Fase 23 no los usa; la Fase 24 podría.

SMEM: el scratchpad gestionado por el programador¶

Éste es el que rompe la intuición de CPU.

La SMEM es on-chip, rápida (~19 TB/s agregada por SM), por bloque de hilos, y explícitamente gestionada. Declaras __shared__ float tile[32][32] en el kernel; el compilador reserva 4 KiB de SMEM por bloque; escribes un bucle de carga que copia un tile de datos globales a tile, haces __syncthreads(), y ahora todos los hilos del bloque pueden acceder a tile a velocidad SMEM.

No hay un mecanismo "SMEM autocargada desde HBM". La SMEM es lo que tú pones ahí. Si no cargas nada, no tienes nada. Si cargas el tile equivocado, tu kernel es incorrecto.

Esto es potencia y trampa. Potencia: puedes estructurar el movimiento de datos exactamente para tu patrón de acceso, evitando los bytes desperdiciados que las caches de CPU siempre traen. Trampa: equivócate con el tamaño del tile y o desperdicias SMEM o causas thrashing con demasiadas cargas globales.

Todo kernel GPU "rápido" — GEMM, attention, convolución — tiene una danza de SMEM. Los labs de la Fase 24 construyen esa danza desde cero sobre un fused-softmax.

Peculiaridad: SMEM y L1 comparten la misma SRAM física en la mayoría de GPUs modernas. Configuras el split en el lanzamiento (p.ej., 100 KiB SMEM + 28 KiB L1, vs 64 KiB SMEM + 64 KiB L1). Los defaults suelen estar bien; los kernels avanzados ajustan esto.

Registros: estado por hilo¶

Cada hilo tiene sus propios registros, recortados del register file de 256 KiB del SM. El compilador decide cuántos registros usa cada hilo.

Compromiso: - Más registros por hilo = más estado retenido por hilo = potencialmente más rápido (evita el spilling) pero menos hilos pueden ser residentes en el SM (ocupación menor). - Menos registros por hilo = algunas variables locales se vuelcan a "memoria local" (¡lenta!) pero más hilos pueden ser residentes (ocupación mayor).

No hay respuesta perfecta. cuBLAS y similares afinan esto por profiling. Para kernels escritos a mano, aprenderás a leer la salida de --ptxas-options=-v para el conteo de registros y ajustar vía __launch_bounds__. Fase 24, no Fase 23.

Para la Fase 23, la lección es: la ocupación depende del conteo de registros, y el conteo de registros depende de la lógica del kernel.

La pirámide de ancho de banda (para una A100)¶

                  ▲
        Registers  │  256 KiB/SM, per-thread, 1-cycle access
                  │
            SMEM  │  164 KiB/SM, per-block, ~19 TB/s, ~20 cycles
                  │
            L2    │  40 MiB, device-wide, ~5 TB/s, ~200 cycles
                  │
            HBM   │  40 GiB, device-wide, ~1.5-3 TB/s, ~500 cycles
                  ▼

El factor entre niveles adyacentes es aproximadamente: - HBM → L2: 2× - L2 → SMEM: 4× - SMEM → reg: 100×+

Así que la mayor discontinuidad de ancho de banda está entre HBM y el resto. La mayor discontinuidad de latencia está entre L2/HBM y SMEM/registros. Juntas explican por qué "tile a SMEM" es la optimización universal en GPU.

Memory coalescing: la regla formal¶

(Planteada en theory/01; formalizada aquí.)

La unidad de load-store de la GPU sirve un warp emitiendo transacciones de memoria. Una transacción es un acceso a una cache line de HBM (cache line = 128 bytes en GPUs modernas). Para servir las 32 instrucciones de carga de un warp:

Si las 32 direcciones caen dentro de una cache line de 128 bytes, una transacción sirve a los 32 hilos. Coalescido. Eficiente en ancho de banda.
Si caen en dos cache lines, dos transacciones. 2× menos eficiente.
Si caen en N cache lines (peor caso N=32), N transacciones. No coalescido. Eficiencia en ancho de banda 1/N.

Regla práctica para fp32: el hilo \(i\) de un warp accediendo a la dirección base + i * 4 (es decir, floats contiguos) está siempre coalescido.

Para fp16/bf16 (\(i \cdot 2\)): también coalescido — 32 × 2 = 64 bytes, media cache line, una transacción.

Para int8 (\(i \cdot 1\)): 32 × 1 = 32 bytes, un cuarto de línea, una transacción. Incluso más eficiente.

La trampa es el acceso con stride: base + (i * S) * sizeof(T) para stride \(S > 1\). Ahora las 32 direcciones se reparten en S cache lines (mejor caso). Ancho de banda efectivo: 1/S del pico. Por eso a veces vale la pena desperdiciar memoria con padding (para evitar strides malos).

Lo que esta jerarquía significa para los operadores de la Fase 22¶

Re-coloca los operadores de la Fase 22 sobre la jerarquía GPU:

Prefill attention (\(P \times P\) por capa). Las matrices \(K, V\) son tile-friendly. El truco completo de Flash-Attention es mantener el working set en SMEM. El tiling HBM → SMEM reduce las lecturas de HBM en un factor O(P). Lab de la Fase 24.
Decode attention (\(1 \times S\) por capa). El \(K, V\) de cache para la secuencia actual es potencialmente 100 MiB → no cabe en SMEM. Stream a través de L2 (a veces acierta) y HBM (la mayoría falla). HBM-bound; la ocupación y el coalescing determinan el factor constante.
FFN matmul (\(1 \times d\) contra \(d \times 4d\)). Para un modelo de 7B, \(4d = 16384\), la matriz de pesos FFN es 4096×16384 fp16 = 134 MiB — no cabe en L2. Stream desde HBM. Mismo diagnóstico: bandwidth-bound en la lectura de pesos.
Append a KV cache. Pequeña escritura de la nueva fila K, V a HBM. Despreciable.
Sampling (argmax / multinomial sobre el vocab). Lee logits (vector del tamaño del vocab, p.ej., 32K fp16 = 64 KiB) — cabe en L2. Cómputo mínimo. Despreciable.

Fíjate en que para todos los operadores, el diagnóstico es "dónde vive el working set" — HBM, L2, SMEM. El plot del roofline es la respuesta; la jerarquía es el vocabulario.

Problemas de práctica¶

La decode-attention de la Fase 22 lee \(2 L S d s\) bytes de cache por step. En A100 fp16 HBM (2 TB/s), Llama-2-7B (L=32, d=4096), S=4096: ¿cuánto dura la pata de lectura de cache? Compara con la latencia por token que la gente cita (10–20 ms).
Un kernel hace dst[i] = src[i * 7] con stride 7, fp32. ¿Fracción coalescida? ¿Ancho de banda efectivo?
La cache de la Fase 22 para MiniGPT es 16 MiB (contexto de 4k). ¿Cabe en la L2 (40 MiB) de A100? ¿Y para Llama-2-7B con contexto 4k (2 GiB)?
¿Por qué "global atomic increment" es 100× más lento que "shared atomic increment" en la mayoría de GPUs? (Pista: ¿dónde se serializa la operación atómica?)

Lo que ahora deberías ser capaz de hacer¶

Dibujar la jerarquía con números.
Indicar la regla de coalescing y aplicarla a un fragmento de código.
Predecir en qué nivel cae el working set de cada operador de la Fase 22.
Explicar por qué SMEM la gestiona el programador y no se autocachea, y por qué eso es potente.

Lo que esta página NO cubre¶

Detalles específicos de HBM3 / HBM3e. Los números de H100 / B100 están señalados donde difieren del baseline de A100, pero esta página apunta a la forma de la jerarquía, no al filo de la innovación.
Conflictos de banco en SMEM en profundidad. Mencionados; el análisis completo es de la Fase 24, cuando reserves SMEM de verdad.
TMA (Tensor Memory Accelerator) en H100+. Fase 27.
Memoria multi-GPU (NVLink, NCCL). Fase 35.

Siguiente: theory/03-warps-and-occupancy.md — el modelo de ejecución de warps formalizado; ocupación como problema de asignación de recursos.