English · Español

Break — el kernel que se sale de shared memory; cómo se ve el fallo¶

🇪🇸 Reservamos un buffer en shared memory un poco más pequeño de lo que realmente usamos. En CUDA, este bug puede pasar inadvertido durante meses si el resto de la memoria compartida del bloque queda intacta. Lo causamos a propósito en un kernel pequeño y vemos los síntomas que delatan el problema.

Síntoma que verá Borja¶

Dos kernels CUDA para una matmul simple por bloques con tiling sobre una matriz \(32 \times 32\) usando tiles en shared memory de tamaño 16:

Run A (control): reserva __shared__ float tile[16][16] e indexa dentro de los límites.
Run B (roto): reserva __shared__ float tile[16][15] (off-by-one en la segunda dimensión) pero indexa tile[r][c] para c ∈ [0, 16).

Ambos compilan sin warnings. Ambos se ejecutan sin crashear (no hay CUDA_ERROR_ILLEGAL_ADDRESS en tiempo de ejecución, porque los accesos a shared memory no tienen comprobación de límites por hardware).

Salida:

Run A: coincide con la matmul de referencia en CPU exactamente (dentro del redondeo fp32).
Run B: difiere de la referencia por cantidades variables, con un patrón específico — cada 16º elemento de cada fila está mal, y los valores incorrectos parecen venir de filas vecinas.

Para una salida \(32 \times 32\), Run B tiene ~30/1024 elementos (~3%) que están salvajemente mal, ~50/1024 (~5%) que están sutilmente mal (off por 1-5%), y el resto correctos.

Si ejecutas un test unitario con tolerancia 1e-3 sobre la diferencia máxima elemento a elemento, falla — pero el mensaje de fallo solo da el primer elemento incorrecto, no el patrón. Un learner que arregle el primer elemento malo añadiendo una guarda podría convencerse de que el test ahora pasa, mientras muchos otros elementos siguen corruptos.

El break, mecánicamente¶

// Run A (control)
__global__ void tile_matmul_correct(...) {
    __shared__ float tile_a[16][16];
    __shared__ float tile_b[16][16];
    // ... load, sync, multiply, accumulate, store ...
    tile_a[threadIdx.y][threadIdx.x] = A[...];   // y in [0,16), x in [0,16)
    __syncthreads();
    // ...
}

// Run B (break)
__global__ void tile_matmul_broken(...) {
    __shared__ float tile_a[16][15];   // <-- one column too small
    __shared__ float tile_b[16][16];
    // SAME indexing as above
    tile_a[threadIdx.y][threadIdx.x] = A[...];   // when threadIdx.x = 15, writes tile_a[y][15] — OOB!
    __syncthreads();
    // ...
}

Cuando threadIdx.x = 15, tile_a[threadIdx.y][15] escribe una columna pasada del buffer reservado. En el layout de shared memory de CUDA, esto sobrescribe la siguiente reserva del mismo bloque — que es el primer elemento de tile_b[0][0]. Los 16 threads que escriben la última columna de tile_a corrompen cada uno un elemento de tile_b.

El patrón: cada 16º elemento de tile_b está corrupto (porque cada threadIdx.x=15 de cada fila threadIdx.y golpea la misma región de overflow, pero para distintas filas de tile_b).

Por qué este es el bug paradigmático de GPU¶

En CUDA, la shared memory se reserva por bloque de un pool fijo (96 KiB en Ampere, 64 KiB en arquitecturas más viejas). Múltiples declaraciones __shared__ dentro de un kernel se concatenan en el espacio de direcciones de shared memory. Una escritura fuera de límites en un buffer corrompe silenciosamente el siguiente buffer de shared memory.

No hay comprobación en runtime. El hardware no impone límites en los accesos a shared memory. El compilador sí comprueba límites estáticos para índices constantes en tiempo de compilación, pero cualquier cosa indexada por threadIdx es dinámica — sus límites no se comprueban.

Contrasta con memoria global: las escrituras fuera de límites a memoria global sí se atrapan en runtime (obtienes CUDA_ERROR_ILLEGAL_ADDRESS o, con cuda-memcheck, un diagnóstico preciso). Shared memory no tiene tal guarda. Esto es el trade-off por la latencia a escala de nanosegundos de la shared memory.

El mismo problema existe en Triton. tl.zeros((16, 15)) en lugar de tl.zeros((16, 16)) produce el mismo patrón. Las abstracciones vectoriales de Triton esconden el indexado por thread pero no el error de tamaño de buffer.

Escalera diagnóstica que Borja debe recorrer¶

Primer chequeo: el test unitario falla. Mira el error: "max diff at element [3, 7] is 8.2". Un elemento. El kernel produjo el número equivocado.
Segundo chequeo: compara todos los elementos con la referencia. Patrón: 3% están salvajemente mal, 5% sutilmente mal. El cluster "salvajemente mal".
Tercer chequeo: los elementos salvajemente-mal comparten una estructura — están en los índices de columna 0, 15, 16, 31 de la salida. O filas 0, 1, 16, 17. El patrón sugiere "cada 16º".
Cuarto chequeo: cuda-memcheck (o compute-sanitizer en toolchains modernos) reporta la escritura OOB. Esta es la pistola humeante. Salida: "Invalid shared write of size 4 at ...".
Diagnóstico: el buffer de shared memory es una columna demasiado pequeño.

Reproductor¶

# Compile both
just phase-24-build-cuda

# Run with the broken version; observe failure
./phase24_matmul_broken 32 > /tmp/output_broken.txt
diff /tmp/output_broken.txt /tmp/output_reference.txt | head -20

# Run with compute-sanitizer
compute-sanitizer ./phase24_matmul_broken 32
# Look for "Invalid __shared__ write"

O en Triton:

# Triton version of the same bug — replace BLOCK_SIZE_K = 16 with BLOCK_SIZE_K = 15 in the kernel
just phase-24-triton-matmul broken

Cascada de pistas¶

(Suave) "El test unitario reporta un único elemento malo. Dibuja el diff completo elemento a elemento. ¿Cuál es el patrón?"
(Media) "Ejecuta compute-sanitizer (o cuda-memcheck) sobre el kernel. ¿Qué reporta?"
(Directa) "El tamaño de la reserva en shared memory es uno menos que el límite del bucle. Iguálalos."

Arreglo¶

Restaura __shared__ float tile_a[16][16]. O, defensivamente, usa constantes con nombre BLOCK_SIZE y las dimensiones del buffer: __shared__ float tile_a[BLOCK_SIZE][BLOCK_SIZE] donde BLOCK_SIZE = 16 es constexpr.

Mejor: escribe una aserción de tamaño en el kernel mediante un static_assert cuando sea posible, o una comprobación en runtime cuando no. Para Triton, declara BLOCK_M, BLOCK_K, BLOCK_N como parámetros tl.constexpr y úsalos de forma consistente.

Lo que hace este break educativo¶

Este bug demuestra el trade-off que define la programación de GPU: el hardware se salta las comprobaciones de límites que el OS / runtime de la CPU haría, porque comprobar límites en cada acceso a shared memory costaría la mitad del throughput. El coste: puedes corromper tu propio programa silenciosamente.

La defensa es tooling: compute-sanitizer lo atrapa en runtime; nvcc -G (modo debug) ayuda; las comprobaciones de límites en tiempo de compilación de Triton ayudan (cuando los índices son constantes en tiempo de compilación); tu suite de tests debe incluir comprobaciones de equivalencia de salida completa, no chequeos puntuales de un único elemento.

Este es el análogo GPU del bug de strcpy en C — buffer overflow silencioso con fallo observable diferido. La Fase 24 introduce tanto la clase de bug como las herramientas (compute-sanitizer, tests de equivalencia exhaustivos, límites en tiempo de compilación) que defienden contra él.

Fallback solo en CPU¶

Si Borja no tiene acceso a una GPU CUDA (el i5-8250U no tiene hardware NVIDIA), este break puede simularse en código CPU: escribe un array C plano float tile_a[16*15], indexa tile_a[r*16 + c] para c ∈ [0, 16). El array C tampoco tiene comprobación de límites; la escritura OOB corrompe lo que toque estar a continuación en la memoria de pila. El patrón es más difícil de reproducir de forma fiable en CPU porque el layout de pila varía, pero el concepto es idéntico.

El lab 01-device-query.md de la Fase 23 y 00-hello-cuda.md de la Fase 24 ya filtran los pasos que requieren GPU detrás de un chequeo de CUDA. Adapta este break a la ruta filtrada.

Lo que este break NO es¶

No es un bug de corrección del algoritmo de matmul.
No es un bug de precisión numérica.
No es un bug de reserva de memoria a nivel del host (sin fallo de malloc).

Es un overflow silencioso de shared memory — la clase más insidiosa de bug en GPU, porque el hardware se niega a ayudarte a encontrarlo. La defensa es disciplina de tooling, no habilidad algorítmica.

Cross-refs¶

theory/01-cuda-programming-model.md — el modelo de ejecución / memoria que lo hace posible.
theory/05-triton-vs-cuda-build-before-abstract.md — Triton hereda el mismo riesgo.
Fase 25 theory/01-dispatcher-and-aten.md — cómo los kernels de PyTorch lo evitan usando constantes en tiempo de compilación y parámetros de plantilla.