Skip to content

English · Español

01 — El modelo de programación CUDA

🇪🇸 CUDA = un lenguaje C/C++ extendido con cualificadores (__global__, __device__, __shared__) y unas pocas variables intrínsecas (blockIdx, threadIdx, blockDim, gridDim). El kernel se lanza con <<<grid, block>>> y cada hilo identifica su trabajo con un índice global que tú calculas a partir de los intrínsecos. Es minimalista; lo difícil no es la sintaxis sino mapear el problema a la jerarquía grid → bloque → warp → hilo.

Esta página da al modelo de programación CUDA el mismo tratamiento compacto que la Fase 23 dio al modelo de ejecución. Tras esto puedes leer cualquier kernel CUDA C e identificar: qué hace cada hilo, qué hace cada bloque, qué memoria toca cada cosa.

Los cinco cualificadores

__global__ void kernel(...)   { ... }   // host calls; runs on device
__device__ int helper(int x)  { ... }   // device-only, callable from kernel
__host__   int main(...)      { ... }   // host-only (default if unqualified)

__shared__ float tile[BLOCK_SIZE];      // SMEM, per-block
__constant__ float lookup[256];         // constant memory, read-only from device

Eso es la mayor parte. CUDA C es C + estos cualificadores + algunas adiciones de sintaxis de lanzamiento + algunos intrínsecos.

El lanzamiento

dim3 grid(num_blocks_x, num_blocks_y, num_blocks_z);   // up to 3D
dim3 block(threads_x, threads_y, threads_z);            // up to 3D, product ≤ 1024
kernel<<<grid, block, shared_mem_bytes, stream>>>(args...);

El <<<...>>> es la única sintaxis no estándar de CUDA C. Le indica a nvcc que emita un lanzamiento de kernel.

  • grid: cuántos bloques lanzar. Cada bloque corre en un SM, no puede migrar.
  • block: cuántos hilos por bloque. Los hilos de un bloque pueden sincronizar (__syncthreads()), compartir SMEM, y se agrupan en warps de 32.
  • shared_mem_bytes: tamaño de SMEM dinámico si el kernel usa extern __shared__. Opcional.
  • stream: en qué stream de CUDA corre este kernel. Múltiples streams habilitan concurrencia entre kernels. Por defecto, el stream por defecto (sincroniza con todo).

Total de hilos lanzados = grid × block (producto de todas las dimensiones). Para un matmul de salida 1024×1024, podrías lanzar <<<dim3(64, 64), dim3(16, 16)>>> — 4096 bloques × 256 hilos = 1M de hilos, uno por elemento de salida (o por grupo cuando hay tiling).

Los intrínsecos

Dentro de un kernel, cada hilo tiene acceso a:

threadIdx.x, threadIdx.y, threadIdx.z    // its position in the block
blockIdx.x,  blockIdx.y,  blockIdx.z     // its block's position in the grid
blockDim.x,  blockDim.y,  blockDim.z     // block shape
gridDim.x,   gridDim.y,   gridDim.z      // grid shape

Más:

warpSize                                  // always 32 today
laneId  =  threadIdx.x % 32              // position within warp (some helpers)
warpId  =  threadIdx.x / 32              // warp index within block

El patrón canónico "¿cuál es mi ID global?" para un lanzamiento 1D:

int tid = blockIdx.x * blockDim.x + threadIdx.x;

Para 2D (trabajo con matrices):

int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;

Escribirás tid o row, col al principio de casi cada kernel.

El kernel mínimo de vector-add

__global__ void vec_add(const float* a, const float* b, float* c, int N) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < N) {
        c[tid] = a[tid] + b[tid];
    }
}

Eso es un kernel completo. Para lanzarlo sobre N = 1M elementos con bloques de 256 hilos:

int N = 1 << 20;
int block_size = 256;
int grid_size = (N + block_size - 1) / block_size;   // ceil(N / block_size)
vec_add<<<grid_size, block_size>>>(d_a, d_b, d_c, N);

El patrón (N + block_size - 1) / block_size es el idiom de "división redondeada hacia arriba" — asegura suficientes bloques para cubrir N incluso cuando N no es múltiplo de block_size, y el guard if (tid < N) dentro del kernel maneja los hilos sobrantes.

Primitivas de sincronización

__syncthreads();              // barrier: all threads in this block must reach this point
__syncwarp();                 // barrier within a warp (cheap, lockstep)
__threadfence();              // memory fence within device
__threadfence_block();        // memory fence within block

cudaDeviceSynchronize();      // host-side: wait for all kernels to finish
cudaStreamSynchronize(stream); // host-side: wait for a stream

Los syncs del lado host ya los conociste en el lab 00 — son los que hacen significativo el cronometraje. El __syncthreads() del lado device es el crucial para la corrección de SMEM: cargar un tile, sync, luego leer el tile. Sin el sync, algunos hilos podrían leer antes de que otros hayan terminado de cargar. Aquí viven los bugs sutiles.

Tres cosas para recordar:

  1. __syncthreads() debe llamarlo todos los hilos del bloque. Llamarlo dentro de una rama divergente es comportamiento indefinido (deadlock o resultado erróneo).
  2. __syncthreads() no atraviesa bloques. Los bloques son independientes. La sincronización entre bloques requiere salir del kernel y lanzar otro.
  3. __syncwarp() es más barato que __syncthreads() pero solo sincroniza dentro de un warp. Útil para reducciones a nivel de warp.

Espacios de memoria, como cualificadores cuda

Cualificador Dónde Acceso Alcance
float* x = cudaMalloc(...); Global (HBM) r/w Todo el dispositivo
__shared__ float t[N]; SMEM r/w Por bloque
int local = ...; Registro (o memoria local si hay spill) r/w Por hilo
__constant__ float c[N]; Constant cache solo lectura desde device Todo el dispositivo
texture<...> Texture cache acceso especial Todo el dispositivo

SMEM y registros se cubrieron en el theory/02 de la Fase 23. La memoria constante es un nicho infrautilizado: muy rápida para valores leídos uniformemente (todos los hilos leen la misma dirección). La memoria de textura es mayormente legado gráfico; los kernels modernos de IA (AI) raramente la usan.

Un patrón canónico: reducción con tile en SMEM

Este es el patrón que escribirás una docena de veces:

__global__ void block_sum(const float* x, float* out, int N) {
    __shared__ float tile[BLOCK];

    int tid = threadIdx.x;
    int gid = blockIdx.x * blockDim.x + tid;

    // Load coalesced into SMEM.
    tile[tid] = (gid < N) ? x[gid] : 0.0f;
    __syncthreads();

    // Reduce in SMEM (tree reduction).
    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
        if (tid < s) tile[tid] += tile[tid + s];
        __syncthreads();
    }

    // Thread 0 writes the block's sum.
    if (tid == 0) out[blockIdx.x] = tile[0];
}

Léelo hasta que sea obvio. Cada línea mapea a un concepto de la Fase 23:

  • tile[tid] = x[gid]: load global coalesced (hilos adyacentes → direcciones adyacentes).
  • __syncthreads() tras el load: requerido antes de que cualquier hilo lea valores del tile escritos por otro.
  • La reducción en árbol: log₂(BLOCK) rondas de sync, cada una halvando los hilos activos.
  • El if-guard if (tid < s): la mitad de los hilos se quedan ociosos cada ronda. ¿Óptimo? No, pero simple. Las reducciones reales usan intrínsecos a nivel de warp (__shfl_down_sync) para saltarse las últimas rondas.

Una reducción real también maneja N > blocks × block_size haciendo que cada bloque reduzca varios chunks de entrada, pero el principio es el mismo.

Overhead del lanzamiento de kernel

Un lanzamiento de kernel desde el host cuesta ~5–10 microsegundos. Suena minúsculo — pero para un kernel de 50 microsegundos, eso es un overhead del 10–20%. Para un kernel de 5 microsegundos, el overhead domina.

Implicaciones:

  • Lanzar muchos kernels diminutos es malo. Fusiónalos.
  • torch.compile / Inductor en la Fase 25 fusionarán por ti.
  • Los kernels escritos a mano a menudo combinan varios operadores (p. ej., GEMM + bias + ReLU) en un solo lanzamiento.
  • Los CUDA Graphs (Fase 33) agrupan lanzamientos de kernel en un solo envío desde el host, eliminando el overhead por lanzamiento.

Para el kernel de la Fase 24, esto es mayormente irrelevante (el kernel es el trabajo, no el lanzamiento). Pero merece la pena conocerlo.

La asincronía en un párrafo

Los lanzamientos de kernel devuelven inmediatamente en el host — se añaden a un stream y se ejecutan asincrónicamente. El host continúa; el device corre en paralelo. cudaDeviceSynchronize() bloquea el host hasta que el device se ponga al día.

cudaMemcpyAsync es asíncrono; cudaMemcpy es síncrono. Los streams se intercalan: puedes transferir H2D un batch mientras computas el kernel del batch anterior, solapando memoria y cómputo. La Fase 33 usará esto; la Fase 24 no.

Lo que esto te da para los laboratorios

Al final de esta página puedes leer los laboratorios sin consultar:

  • __global__, <<<grid, block>>>, threadIdx, blockIdx, __shared__, __syncthreads() son todos vocabulario primitivo.
  • El patrón "tamaño de grid redondeado hacia arriba + guard con if (tid < N)" es estándar.
  • El patrón "cargar tile en SMEM → sync → reducir → sync → escribir de vuelta" es estándar.

Eso basta para leer el kernel ingenuo de softmax en lab/01 y el tuneado en lab/02.

Lo que esta página NO cubre

  • Templates de CUDA C++, uso avanzado de __launch_bounds__, CUTLASS. La Fase 24 se queda en CUDA con sabor a C; los templates no son necesarios para el kernel elegido.
  • Cooperative groups (cg::thread_block_tile). Una alternativa moderna al __syncthreads crudo; se menciona en la Fase 27 si hace falta.
  • Ensamblador PTX inline. No lo necesitarás para la Fase 24.
  • Sintaxis de CPU/Triton/PyTorch. Las cubren las teorías 02, 03 y 04.

Siguiente: theory/02-from-naive-to-tiled.md — la ruta canónica de optimización para un solo kernel.