Skip to content

English · Español

05 — Triton vs CUDA crudo: qué te da cada uno, y por qué esperamos hasta la Fase 24

🇪🇸 La regla §0.1 de CLAUDE.md dice: construye antes de abstraer. PyTorch llega en la Fase 24. Pero dentro de la Fase 24 hay otra elección — ¿escribimos CUDA directo o usamos Triton? Esta página justifica la respuesta, y explica por qué escribir CUDA primero aunque sea brevemente es no negociable.

Este archivo es el complemento en profundidad de theory/03-triton.md. Comparamos CUDA C++ y Triton frente a frente en un único kernel (vector add), mostramos qué oculta Triton y qué no, y explicamos por qué la disciplina construir-antes-de-abstraer de CLAUDE.md §0.1 dice CUDA primero, Triton después, kernels internos del framework al final.

Qué te da CUDA

CUDA C++ expone el modelo de ejecución de la GPU directamente:

  • Hilos organizados en bloques, bloques en una grid.
  • Shared memory declarada e indexada explícitamente.
  • Primitivas a nivel de warp (shuffles, vote, ballot).
  • Jerarquía de memoria controlada a mano (global, shared, registros, L1 / L2, texture).
  • Sincronización vía __syncthreads(), atomics, barreras.

Escribes __global__ void kernel(...) { int idx = blockIdx.x * blockDim.x + threadIdx.x; ... }. Cada concepto mapea directamente a hardware. El compilador (nvcc) hace casi ninguna transformación de alto nivel — lo que escribes es lo que corre.

Coste: verboso. Un matmul "bueno" tuneado a mano en CUDA son 200-400 líneas para un tile 2-D con double buffering. Una versión "decente" son 50-80 líneas. La versión de juguete de 5 líneas rinde al 1-5% del pico.

Qué te da Triton

Triton (Tillet et al. 2019) es un DSL embebido en Python que compila a PTX. Diferencias clave:

  • Abstracciones a nivel de bloque. En lugar de indexación a nivel de hilo, operas sobre tiles (p. ej., un bloque \(128 \times 128\) de una matriz). El compilador decide cómo mapear tiles a hilos / warps.
  • Auto-tuning. Los tamaños de bloque son parámetros declarativos; tú das pistas, el compilador enumera configuraciones.
  • Pipelining implícito. El pipelining software (solapar loads de memoria con cómputo) se genera automáticamente; en CUDA lo escribes a mano con intrínsecos __pipeline_*.
  • Gestión de shared memory. Asignada, indexada y resuelta de conflictos de bancos automáticamente.

Coste: menos control. Para el 95% de los kernels, "el compilador escoge buenos defaults" es más rápido que "Borja tunea a mano". Para el 5% restante — habitualmente el más crítico en rendimiento — las abstracciones de Triton ocultan cosas que podrías necesitar controlar (warp specialization, async copies, indexación custom).

Comparación concreta: vector add

CUDA C++

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

// Launch
int threads = 256;
int blocks = (N + threads - 1) / threads;
vec_add<<<blocks, threads>>>(a, b, c, N);

8 líneas de kernel + 3 de lanzamiento. Cada concepto (hilo, bloque, bounds check) es explícito.

Triton

import triton
import triton.language as tl

@triton.jit
def vec_add(a_ptr, b_ptr, c_ptr, N, BLOCK_SIZE: tl.constexpr):
    pid = tl.program_id(0)
    offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
    mask = offsets < N
    a = tl.load(a_ptr + offsets, mask=mask)
    b = tl.load(b_ptr + offsets, mask=mask)
    tl.store(c_ptr + offsets, a + b, mask=mask)

# Launch
grid = lambda meta: ((N + meta["BLOCK_SIZE"] - 1) // meta["BLOCK_SIZE"],)
vec_add[grid](a, b, c, N, BLOCK_SIZE=1024)

8 líneas de kernel + 2 de lanzamiento. Superficie similar, pero:

  • tl.arange(0, BLOCK_SIZE) te da un vector de índices, no un índice escalar de hilo.
  • tl.load(..., mask=mask) maneja el bounds check internamente.
  • BLOCK_SIZE es una constante de tiempo de compilación — el compilador especializa el kernel para ese valor.

Para un kernel tan simple, Triton no es dramáticamente más conciso. La ventaja está al nivel de matmul / attention, donde 200 líneas de CUDA → 30 líneas de Triton.

Cuándo gana CUDA

  • Warp specialization. Tratar warps dentro de un bloque como si tuvieran roles distintos (p. ej., producer-consumer en attention). Triton tiene soporte limitado.
  • Trasteo del layout de Tensor Cores. Cuando el layout de datos (row-major vs col-major vs swizzled) interactúa con la instrucción MMA de Tensor Core, el control a mano importa.
  • Indexación no uniforme. Si tu kernel necesita que los hilos de un warp hagan cosas distintas según los datos, CUDA te da el control de divergencia; la abstracción vectorial de Triton asume uniformidad.

Cuándo gana Triton

  • Matmul / attention en bloques. Los kernels donde el 80% de la velocidad viene de acertar el tamaño de tile; Triton auto-tunea esto.
  • Prototipado rápido. Tener un kernel corriendo en minutos, luego iterar.
  • Mantenibilidad. Un kernel Triton es más fácil de leer 6 meses después que un kernel CUDA.

Por qué esperamos hasta la Fase 24

Según CLAUDE.md §0.1 ("construye antes de abstraer"):

  1. Fases 0-8: NumPy. Todo a mano. Ver cómo un matmul, una softmax, una layer-norm se descomponen en bucles.
  2. Fases 9-17: Componer módulos en Python/NumPy crudo. El mecanismo de attention, la MLP, el lookup de embedding — todo visible.
  3. Fases 18-22: Entrenamiento, evaluación, inferencia. Aún NumPy. El optimizer, el harness de eval, el KV cache — todo visible.
  4. Fase 23: Modelo mental de GPU. Aún sin código GPU. Solo por qué las GPU difieren de las CPU.
  5. Fase 24: Primer kernel GPU. CUDA primero (un kernel, en el lab), luego Triton para el resto.
  6. Fase 25: Autograd / dispatcher de PyTorch. Ahora podemos leer el código de PyTorch porque ya lo escribimos todo a mano.

El orden importa. Si hubiéramos empezado con PyTorch en la Fase 9, cada ejercicio de depuración posterior tendría un misterio "lo hace el framework". Si hubiéramos empezado con Triton en la Fase 24, las optimizaciones generadas por el compilador serían invisibles. Al escribir el vector-add de CUDA a mano primero, las abstracciones que aporta Triton se vuelven visibles — sabes qué están ocultando.

Un diff en coste mental:

Etapa Concepto que debes sostener Nueva abstracción
Matmul NumPy Fases 0-8 indexación de array, broadcast ninguna
Modelo mental Fase 23 warp, occupancy, jerarquía de memoria vocabulario de hardware
CUDA Fase 24 (un kernel) índice de hilo/bloque, shared mem lenguaje: __global__, <<<>>>
Triton Fase 24 (resto) tile, máscara, tamaño de bloque DSL sobre CUDA
Internals de PyTorch Fase 25 dispatcher, autograd framework sobre CUDA / Triton

Cada paso es una capa nueva. Si te saltas una capa, la siguiente capa se vuelve mucho más difícil de depurar.

Lo que esto no significa

No significa "usa siempre CUDA en producción". Significa:

  • Aprende CUDA primero escribiendo uno o dos kernels a mano.
  • Luego usa Triton para el resto de labs de la Fase 24.
  • Luego usa los built-ins de PyTorch (torch.nn.functional, torch.compile) a partir de la Fase 25.

En producción, la estratificación es la misma: kernels Triton para rutas calientes, built-ins PyTorch para el resto, CUDA crudo solo para los 1-2 sitios donde las abstracciones de Triton son limitantes. Pero la secuencia de aprendizaje es bottom-up, aunque la secuencia de despliegue sea top-down.

La salvedad de §A13

El mini-GPT de §A13 no se beneficia de aceleración GPU a esta escala (ver quiz de la Fase 23 q-23-05). Los kernels de la Fase 24 son ejercicios de escribir código GPU, no herramientas de producción para nuestro modelo. El artefacto pedagógico es:

  • Un kernel CUDA (vector-add o matmul escalar) en lab/00-hello-cuda.md.
  • Un port a Triton del mismo kernel en lab/01-naive-kernel.md.
  • Un matmul Triton con tiling en lab/02-tuned-kernel.md.
  • Cableado en PyTorch en lab/03-triton-and-pytorch.md.

Ese es todo el alcance de la Fase 24. No escribimos FlashAttention desde cero; eso es un EDUCATIONAL_STUB según CLAUDE.md §0.1.8.

Cita

Tillet, P., Kung, H. T., & Cox, D. (2019). Triton: An Intermediate Language and Compiler for Tiled Neural Network Computations. MAPL'19. https://www.eecs.harvard.edu/~htk/publication/2019-mapl-tillet-kung-cox.pdf — el paper original de Triton, la sección 3 sobre la IR a nivel de tile es la fuente del encuadre de "qué te da Triton" aquí.

Recapitulación en un párrafo

CUDA expone el modelo de hilo/bloque/shared-memory de la GPU directamente; Triton abstrae sobre tiles y auto-tunea tamaños de bloque. Para kernels simples, el número de líneas es comparable; para matmul/attention, Triton es 5-10× más corto. Escribimos un kernel CUDA a mano en la Fase 24 para asentar las abstracciones que aporta Triton, luego usamos Triton para el resto. La disciplina construir-antes-de-abstraer de CLAUDE.md §0.1 produce esta estratificación: NumPy → CUDA → Triton → PyTorch → helpers del framework. Cada capa tiene sentido solo si la inferior se escribió primero a mano.

Cross-refs: theory/01-cuda-programming-model.md, theory/03-triton.md, theory/04-pytorch-as-substrate.md, CLAUDE.md §0.1 (la regla construir-antes-de-abstraer), Fase 23 theory/05-cpu-only-roofline-i5-8250u.md (el contexto de hardware).