Skip to content

English · Español

03 — Warps y ocupación

Ocupación = (warps activos en un SM) / (warps máximos por SM). No es lo mismo que "muchos hilos en vuelo"; es "suficientes warps residentes para que el scheduler siempre tenga uno listo cuando otro espera memoria". Se limita por registros/hilo, SMEM/bloque y hilos/bloque — tres presupuestos que compiten.

Esta página hace de la ocupación una noción precisa, deriva los tres recursos que la acotan, y explica la relación con el rendimiento.

Qué mide la ocupación

Un SM tiene un cap hardware de warps residentes — típicamente 48 o 64 warps por SM (1536 o 2048 hilos). "Residente" = registros reservados, estado asignado, listo para correr cuando el scheduler lo elija.

Ocupación = (warps residentes) / (warps máximos por SM).

100% de ocupación significa que el SM tiene el pool de warps máximo permitido. 25% significa que solo un cuarto de ese pool está cargado.

Por qué importa: el warp scheduler elige un warp por ciclo para emitir una instrucción. Cuando un warp se detiene (carga de memoria aún sin volver, dependencia con una instrucción lenta), el scheduler elige otro warp. Con warps suficientes, cada ciclo tiene un candidato listo; el SM nunca queda inactivo. Con muy pocos, el SM queda inactivo durante los stalls.

Los kernels memory-bound necesitan alta ocupación para ocultar la latencia de HBM (~500 ciclos). Los kernels compute-bound pueden correr a menor ocupación porque no paran tan a menudo.

Qué acota la ocupación

Tres recursos, todos por SM:

1. Register file

Cada hilo retiene cierto número de registros (decidido por el compilador a partir de la lógica del kernel). Total de registros reservados = hilos × registros/hilo. Esto debe caber en el register file del SM (p.ej., 64K registros de 32 bits = 256 KiB en A100).

Si tu kernel usa 64 registros/hilo, y el SM tiene 64K registros, entonces hilos residentes máximos = 64K / 64 = 1024. Si el cap del SM es 2048 hilos, estás a 1024/2048 = 50% de ocupación.

Para subir ocupación: usa menos registros por hilo. Compromiso: la reutilización forzada de registros puede volcar locales a "memoria local" (¡que en realidad es HBM!) — un stall esperando una carga de 500 ciclos cada vez que tocas una variable volcada. Así que bajar el conteo de registros puede perjudicar un kernel compute-bound.

2. SMEM por bloque

Cada bloque reserva algo de SMEM __shared__ en el lanzamiento. Total de SMEM reservada = bloques/SM × SMEM/bloque. Debe caber en la asignación de SMEM del SM (típicamente 100–164 KiB).

Si tu kernel usa 100 KiB de SMEM/bloque y el SM tiene 100 KiB, obtienes exactamente un bloque por SM. Con 128 hilos/bloque (4 warps), estás a 4/64 = ~6% de ocupación. Letal.

Para subir: usa menos SMEM por bloque con un tiling distinto (tiles más pequeños). O reparte trabajo entre más bloques.

3. Hilos por bloque

El tamaño de bloque que elijas (p.ej., <<<grid, 256>>> = 256 hilos/bloque) debe dividir el cap de hilos del SM. Con un cap de 2048 hilos y bloques de 256 hilos, obtienes 8 bloques/SM (máx.). Con bloques de 1024 hilos, obtienes 2 bloques/SM.

Bloques más grandes = menos bloques/SM pero más hilos/bloque. Bloques más pequeños = más bloques/SM. La elección importa porque cada bloque reserva su propia SMEM y registros — el uso total de recursos escala con el conteo de bloques.

Regla práctica: tamaño de bloque de 128–256 hilos (4–8 warps) suele ser un buen punto de partida. Potencias de dos para tener matemáticas limpias.

La ecuación de ocupación

occupancy = min(
    threads_per_block / (max_threads_per_SM / blocks_per_SM_by_registers),
    threads_per_block / (max_threads_per_SM / blocks_per_SM_by_SMEM),
    threads_per_block / (max_threads_per_SM / blocks_per_SM_by_threads)
)

Esto es lo que la CUDA Occupancy Calculator de NVIDIA (y la salida de nvcc --resource-usage + --ptxas-options=-v) calcula por ti. La Fase 24 las usa. La Fase 23 solo observa el resultado vía cudaOccupancyMaxActiveBlocksPerMultiprocessor().

Ocupación vs rendimiento: no son lo mismo

Un kernel de alta ocupación puede seguir siendo lento:

  • Memory-bound: incluso al 100% de ocupación, estás topado por el ancho de banda de HBM. Añadir más warps no ayuda una vez que la tubería de memoria está llena.
  • Coalescing roto: incluso al 100% de ocupación, si cada warp hace 32 cargas no-coalescidas, tu ancho de banda efectivo es 1/32 del pico. El rendimiento cae 32× independientemente de la ocupación.
  • Divergencia rota: incluso al 100% de ocupación, un warp con divergencia 4-way corre al 25% de eficiencia en la región divergente.

A la inversa, un kernel de baja ocupación puede ser rápido:

  • Compute-bound con alta ILP: un kernel que hace 1000 FLOPs por carga de memoria apenas se detiene; un warp por SM basta para saturar las FPUs (en principio — en realidad quieres 2–3 para llenar el pipeline).
  • Matmul limitado por Tensor Cores: un solo warp emitiendo instrucciones de Tensor Core puede hacer un trabajo enorme; pasar de ~50% de ocupación da rendimientos decrecientes.

Usa la ocupación como diagnóstico, no como objetivo. El objetivo es el rendimiento. Baja ocupación → sospecha (puede ser presión de registros o uso excesivo de SMEM); investiga. Alta ocupación + bajo rendimiento → sospecha de coalescing, divergencia o memory-boundedness.

Cómo la Fase 23 mide la ocupación

No escribirás un kernel en la Fase 23. Pero los benchmarks de experiments/23-device-profile/ llaman a cuBLAS y cudaMemcpy, que tienen características conocidas de ocupación/rendimiento. Puedes:

  1. Consultar al dispositivo por max_threads_per_SM, max_blocks_per_SM, max_warps_per_SM, register_file_size, shared_mem_per_SM.
  2. Calcular el presupuesto de ocupación — es decir, si escribieras un kernel con X registros e Y KiB de SMEM y Z hilos/bloque, ¿qué ocupación obtendrías?
  3. Confirmar que el GEMM de cuBLAS al tamaño que mediste alcanza el pico fp16/bf16 publicado (>80% es realista). Si sí, el tuning de cuBLAS por NVIDIA logra alto rendimiento efectivo; el potencial de la GPU es real.

El experimento real está en lab/02-bandwidth-test.md y lab/03-gpu-roofline.md.

Un ejemplo trabajado de ocupación

Escribes un kernel de fused-softmax con:

  • 128 hilos/bloque.
  • 32 registros/hilo (el compilador lo reporta).
  • 16 KiB de SMEM/bloque (lo usas para alojar una fila de logits para el softmax por filas).

En una A100 (máx 2048 hilos/SM, 64K regs/SM, 164 KiB SMEM/SM, 32 bloques/SM):

  • Límite por hilos: máx 2048 hilos / 128 = 16 bloques/SM. Bajo el cap de 32 bloques.
  • Límite por registros: máx 64K regs / (128 × 32) = 64K / 4096 = 16 bloques/SM. Igual que el límite por hilos.
  • Límite por SMEM: 164 KiB / 16 KiB = 10 bloques/SM. Esta es la restricción dominante.

Bloques/SM máx = min(16, 16, 10) = 10. Hilos máx = 10 × 128 = 1280. Ocupación = 1280 / 2048 = 62.5%.

Para subir la ocupación: reducir SMEM/bloque a ~10 KiB (tile de softmax más pequeño, pero con más recargas globales); o reducir conteo de registros.

Para comprobar si vale la pena subirla: perfila (la Fase 24 usa ncu).

Tensor Cores: un pipeline de cómputo paralelo

Una GPU moderna tiene dos tipos de FPUs por SM: CUDA cores (una operación fp32 por ciclo por core) y Tensor Cores (una operación matrix-multiply-accumulate por ciclo por Tensor Core — pico de rendimiento mucho mayor).

Para operaciones de Tensor Core (matmul de tiles pequeños en fp16/bf16/fp8/int8), el rendimiento es 4–16× mayor que los CUDA cores al mismo dtype. Por eso los números de "fp16 TFLOPS" de una hoja de specs son un orden de magnitud superiores a los de "fp32 TFLOPS" — el número fp16 asume Tensor Cores; el fp32 asume CUDA cores.

La Fase 23 no escribe kernels de Tensor Core. Pero el pico medido en peak_flops.py (vía GEMM de cuBLAS en fp16) golpea Tensor Cores por defecto — así es como cuBLAS llega a ~80% del pico de fabricante. Los labs de la Fase 24 escriben un kernel manual de Tensor Core.

Por qué la ocupación es el último escalón para entender la Fase 23

La ocupación es la abstracción que ata toda la contabilidad de recursos en la GPU. Una vez que puedes:

  1. Calcularla a partir de un esbozo de kernel (usando la ecuación anterior),
  2. Diagnosticar qué la acota (cuál de los tres),
  3. Relacionarla con el rendimiento (alta ocupación es necesaria pero no suficiente),

...tienes el vocabulario para leer cualquier paper de rendimiento de GPU. La Fase 24 será ajustar kernels, y cada perilla de ajuste (conteo de registros, tamaño de bloque, tamaño de tile, split de SMEM) cae en algún punto de este espacio ocupación / rendimiento / coalescing.

Problemas de práctica

  1. En A100 (64K regs/SM, 164 KiB SMEM/SM, cap 2048 hilos, cap 32 bloques): un kernel usa 256 hilos/bloque, 48 regs/hilo, 32 KiB SMEM/bloque. Calcula la restricción dominante y la ocupación resultante.
  2. En H100 (caps similares): el mismo kernel. ¿Misma ocupación? ¿Distinta? ¿Por qué?
  3. Un kernel está "memory-bound al 25% de ocupación y 80% de utilización de ancho de banda." ¿Deberías subir la ocupación? Muestra el razonamiento.
  4. Un kernel está "compute-bound al 100% de ocupación, alcanzando el 70% del pico fp16 de Tensor Core." ¿Cómo intentarías ganar el 30% que falta? (Pista: no subiendo la ocupación.)

Lo que ahora deberías ser capaz de hacer

  1. Indicar los tres límites de recursos sobre la ocupación.
  2. Calcular la ocupación para un kernel dadas su huella de registros/SMEM/hilos.
  3. Explicar por qué la ocupación es necesaria pero no suficiente para alto rendimiento.
  4. Leer la salida de un profile e identificar qué recurso acota la ocupación.

Lo que esta página NO cubre

  • Ajustar un kernel real. El lab de la Fase 24 es donde el ajuste de ocupación se vuelve práctico.
  • Scheduling asíncrono de warps (warp specialization). Característica de vanguardia de Hopper; Fase 27/36.
  • Programación de Tensor Cores. El kernel softmax-over-vocab de la Fase 24 no usa Tensor Cores; los kernels GEMM sí. Tensor Cores es tema de la Fase 27.

Siguiente: theory/04-gpu-roofline.md — re-derivar el plot del roofline para GPU y colocar sobre él los operadores de la Fase 22.