CUDA Tile llega a C++: kernels GPU sin gestionar hilos

Hasta ahora, la programación por tiles en GPUs NVIDIA solo era accesible desde Python. Con la versión CUDA 13.3, la compañía abrió ese mismo modelo a C++, una pieza relevante para los equipos que mantienen codebases extensos en C++ y no pueden migrar a Python sin un costo enorme.

CUDA Tile C++ es la expresión del modelo de tiles sobre el lenguaje C++, construida encima de la especificación CUDA Tile IR. Permite escribir kernels expresando las operaciones sobre porciones (tiles) de arreglos multidimensionales, en lugar del clásico esquema SIMT (single instruction, multiple threads) donde el programador define el comportamiento de cada hilo por separado.

¿Qué cambia respecto al modelo SIMT tradicional?

En el modelo SIMT el desarrollador especifica qué hace cada hilo, calcula índices manualmente y lanza el kernel con un número explícito de bloques e hilos. El ejemplo canónico es la suma de vectores:

__global__ void vecAdd(float* A, float* B, float* C, int vectorLength)
{
    int workIndex = threadIdx.x + blockIdx.x * blockDim.x;
    if (workIndex < vectorLength) {
        C[workIndex] = A[workIndex] + B[workIndex];
    }
}

En el modelo por tiles, el código equivalente no menciona hilos. El desarrollador divide los datos en tiles y describe la operación matemática sobre esos tiles. El resto, incluida la paralelización dentro del bloque y el movimiento de memoria, lo resuelve el compilador.

#include "cuda_tile.h"
__tile_global__ void vectorAdd(float* a, float* b, float* out, size_t n) {
    namespace ct = cuda::tiles;
    using namespace ct::literals;

    auto aSpan = ct::tensor_span{a,   ct::extents{n}};
    auto bSpan = ct::tensor_span{b,   ct::extents{n}};
    auto oSpan = ct::tensor_span{out, ct::extents{n}};

    auto aView = ct::partition_view{aSpan, ct::shape{8_ic}};
    auto bView = ct::partition_view{bSpan, ct::shape{8_ic}};
    auto oView = ct::partition_view{oSpan, ct::shape{8_ic}};

    int bx = ct::bid().x;
    auto aTile = aView.load(bx);
    auto bTile = bView.load(bx);

    auto oTile = aTile + bTile;
    oView.store(oTile, bx);
}

El decorador __tile_global__ marca el kernel como un kernel de tiles. ct::tensor_span es esencialmente un puntero a un arreglo multidimensional con información de forma y layout, equivalente al std::mdspan de C++23. ct::partition_view envuelve ese tensor span y lo presenta como una serie de particiones de tamaño fijo, no superpuestas. Aquí el tamaño de tile es 8 elementos, declarado con la constante entera 8_ic.

Optimizaciones que sí debe especificar el desarrollador

CUDA Tile C++ automatiza mucho, pero hay tres ajustes que el compilador no puede deducir solo y que pesan en rendimiento real:

• __restrict__ en los punteros: indica al compilador que no hay aliasing entre los arreglos. Es el mismo recurso que existe en CUDA C++ tradicional.
• Alineamiento de 16 bytes con ct::assume_aligned<16>: avisa al compilador que las bases de los punteros están alineadas, lo que habilita patrones de acceso a memoria más eficientes. Los punteros devueltos por cudaMalloc siempre cumplen esta condición (de hecho tienen alineamiento de 256 bytes).
• Tile size razonable: el ejemplo usa 8 para mostrar el flujo, pero en código real conviene usar tiles mucho más grandes (1024 en el ejemplo final del blog).

Para datos cuyo largo no es divisible por el tile size, CUDA Tile C++ provee load_masked y store_masked, que manejan los bordes sin que el desarrollador tenga que escribir condicionales.

Multiplicación de matrices y tensor cores

El caso de uso donde CUDA Tile C++ aporta más es en multiplicación de matrices, donde los kernels pueden apoyarse en operaciones de multiplicación-acumulación matricial (mma) que aprovechan directamente los tensor cores del hardware. El blog muestra cómo combinar particiones de tiles con esas operaciones para acumular resultados parciales de manera eficiente, sin escribir manualmente la sincronización entre hilos del bloque.

Profiling y compatibilidad

Los kernels escritos en CUDA Tile C++ se perfilan con NVIDIA Nsight Compute y entregan estadísticas específicas por tile, además de las métricas de rendimiento a nivel de fuente que ya están disponibles en CUDA C++. La integración no obliga a cambiar de toolchain.

Los requisitos para usar el modelo completo son tres:

• GPUs con compute capability 8.x o superior (Ampere, Ada, Hopper y posteriores).
• Driver NVIDIA R580 o posterior.
• CUDA Toolkit 13.3 o posterior.

¿Por qué importa este cambio?

La extensión a C++ apunta a un público muy concreto: los equipos de HPC, simulación y backend de IA que tienen millones de líneas en C++ y no pueden moverse a Python solo para aprovechar abstracciones modernas. Hasta ahora, esos equipos seguían escribiendo kernels SIMT explícitos. Con CUDA Tile C++ pueden expresar sus cargas como operaciones sobre tiles dentro del mismo lenguaje, con la promesa de portabilidad automática entre arquitecturas NVIDIA.

Como referencia, kernels que en CUDA C++ requieren cientos de líneas de gestión manual de threads y memoria compartida quedan en pocas decenas de líneas con CUDA Tile C++, sin sacrificar rendimiento según los ejemplos publicados por NVIDIA. El precio: el código solo corre en hardware NVIDIA y queda atado al CUDA Toolkit 13.3.