CUDA 13.1 ha introducido CUDA Tile, un nuevo paradigma de programación GPU basado en tiles que busca facilitar el paralelismo fino. Uno de sus principales beneficios es la apertura del lenguaje: cualquier lenguaje de programación puede dirigirse a CUDA Tile, lo que permite a los desarrolladores incorporar aceleración GPU basada en tiles en una amplia variedad de ecosistemas.
En respuesta a la fuerte demanda de desarrolladores experimentados, se lanza cuTile BASIC para GPUs, que lleva la programación CUDA Tile a este lenguaje que había sido desatendido.
¿Qué es cuTile BASIC?
cuTile BASIC es una expresión del modelo de programación CUDA Tile en BASIC, basado en la especificación CUDA Tile IR. Permite escribir kernels en BASIC utilizando un modelo basado en tiles, que se ajusta de forma natural a un lenguaje como BASIC, que precede a la programación multihilo.
cuTile BASIC combina la potencia de las GPUs con la simplicidad y el encanto anacrónico del lenguaje BASIC, que evoca una era más pixelada. ¡Numerar manualmente tus líneas de código nunca había sido tan atractivo y rápido!
¿Para quién es cuTile BASIC?
BASIC es uno de los lenguajes de programación más antiguos y es venerado por una generación de desarrolladores que recuerdan con nostalgia el sonido de un módem de 300 baudios. Para muchos de ellos, BASIC fue su primera introducción a la programación.
Ahora, los desarrolladores que aún tienen el BASIC grabado en su memoria pueden llevar aplicaciones heredadas a la computación acelerada por GPU de NVIDIA por primera vez. Esto desbloquea un rendimiento y funcionalidad que el lenguaje BASIC nunca había imaginado anteriormente, permitiendo que tu Lunar Lander se desplace por la superficie lunar más rápido que una misión Artemis.
Cómo instalar cuTile BASIC
Para comenzar, instala cuTile BASIC usando PIP:
pip install git+https://github.com/nvidia/cuda-tile.git@basic-experimental
Los requisitos de hardware y software para ejecutar cuTile BASIC se detallan al final de esta publicación (se recomiendan 64k de RAM o más).
Ejemplo de cuTile BASIC
Si has aprendido CUDA C++, probablemente te has encontrado con el kernel canónico de suma de vectores. Un kernel de suma de vectores en CUDA C++ se ve así:
__global__ void vecAdd(float* A, float* B, float* C, int vectorLength) { /* calcular el índice de hilo */ int workIndex = threadIdx.x + blockIdx.x*blockDim.x; if(workIndex < vectorLength) { /* realizar la suma de vectores */ C[workIndex] = A[workIndex] + B[workIndex]; } }
En este kernel, el trabajo de cada hilo está explícitamente especificado, y el programador, al lanzar este kernel, deberá especificar el número de bloques e hilos a lanzar.
Ejemplo de suma de vectores en cuTile BASIC
Ahora, veamos el código equivalente en cuTile BASIC. No necesitamos especificar qué hace cada hilo. Solo tenemos que dividir los datos en tiles y especificar qué operaciones matemáticas deben realizarse en esos tiles. Todo lo demás se maneja automáticamente.
El kernel de suma de vectores en cuTile BASIC es el siguiente:
10 REM Suma de Vectores: C = A + B 20 INPUT N, A(), B() 30 DIM A(N), B(N), C(N) 40 TILE A(128), B(128), C(128) 50 LET C(BID) = A(BID) + B(BID) 60 OUTPUT C 70 END
Este ejemplo es básico y utiliza la sintaxis estándar de BASIC, con algunas consideraciones importantes:
- Indexar un arreglo devuelve un tile, que es un subconjunto del arreglo.
- BID es una variable incorporada que especifica el índice del bloque de tiles.
- TILE especifica el tamaño de los tiles en los que se deben particionar los arreglos.
Observa que no tuvimos que especificar nada más que la operación de suma. Todo lo demás es manejado por cuTile BASIC.
Todo junto
Ahora mostraremos cómo ejecutar este kernel de suma de vectores en BASIC. El flujo de trabajo es sencillo: primero se compila la función BASIC a un cubin y luego se lanza en la GPU. Por razones de brevedad, hemos omitido el código aburrido de host y wrapper en Python, pero puedes encontrarlo en nuestro repositorio de GitHub.
Si tienes las versiones adecuadas de CUDA Toolkit y Python instaladas y has descargado el repositorio cuTile BASIC desde GitHub, puedes ejecutar el siguiente comando:
$ python examples/vector_add.py [1/2] Compilando a cubin ... Arreglos: ['A', 'B', 'C'], tile_shapes={'A': [128], 'B': [128], 'C': [128]}, grid_size=8 [2/2] Lanzando kernel en GPU ... Resultados (mostrando 5 muestras de 1024): C[ 0] = 0.0 (esperado 0.0) C[ 1] = 3.0 (esperado 3.0) C[ 511] = 1533.0 (esperado 1533.0) C[ 512] = 1536.0 (esperado 1536.0) C[1023] = 3069.0 (esperado 3069.0) VERIFICACIÓN APROBADA (max_diff=0.000000, 1024 elementos)
Si tu salida se ve igual, ¡felicitaciones! Has ejecutado tu primer programa en cuTile BASIC, y posiblemente tu primer programa de cualquier tipo escrito en BASIC. Max Headroom estaría orgulloso.
Multiplicación de matrices en BASIC
BASIC es un lenguaje simple donde pocas líneas de código pueden expresar algoritmos comunes. Considera un kernel de multiplicación de matrices (GEMM) en BASIC, que se muestra a continuación:
10 REM GEMM: C(M,N) = A(M,K) * B(K,N) 15 INPUT M, N, K, A(), B() 20 DIM A(M, K), B(K, N), C(M, N) 30 TILE A(128, 32), B(32, 128), C(128, 128), ACC(128, 128) 40 LET TILEM = INT(BID / INT(N / 128)) 50 LET TILEN = BID MOD INT(N / 128) 60 LET ACC = 0.0 70 FOR KI = 0 TO INT(K / 32) - 1 80 LET ACC = MMA(A(TILEM, KI), B(KI, TILEN), ACC) 90 NEXT KI 100 LET C(TILEM, TILEN) = ACC 110 OUTPUT C 120 END
En este kernel, además de la sintaxis estándar de BASIC, TILE especifica cómo A, B y C deben ser tileados y el tamaño del tile acumulador, ACC. MMA es la llamada a la función para multiplicar y acumular matrices. Nota lo simple que es este código. Se especifica cómo se debe subdividir la información en tiles y se indica el algoritmo a un alto nivel, mientras que CUDA Tile maneja todo lo demás.
Este ejemplo también está disponible en la carpeta de ejemplos del repositorio de GitHub. Al ejecutarlo, produce la siguiente salida:
$ python examples/gemm.py [1/2] Compilando a cubin ... M=512, N=512, K=512, tile_shapes={'A': [128, 32], 'B': [32, 128], 'C': [128, 128]}, grid_size=16 [2/2] Lanzando kernel en GPU ... Resultados (mostrando 5 muestras de 512x512 = 262144 elementos): C[0,0] = -0.1199 (esperado -0.1199) C[0,1] = -14.4456 (esperado -14.4456) C[256,0] = -15.8891 (esperado -15.8891) C[256,1] = -2.8646 (esperado -2.8646) C[511,511] = 11.4724 (esperado 11.4724) VERIFICACIÓN APROBADA (max_diff=0.000012, tol=0.005120)
La multiplicación de matrices, como se muestra arriba, es fundamental en herramientas de inteligencia artificial como los modelos de lenguaje de gran tamaño. Con cuTile BASIC, los desarrolladores ahora pueden explorar los límites de la inteligencia artificial en modelos con billones de parámetros desde un lenguaje que apenas podía imaginar un megabyte de memoria del sistema.
Cómo obtener cuTile
Para ejecutar programas cuTile BASIC, necesitas lo siguiente:
- Una GPU con capacidad de cómputo 8.x, 10.x, 11.x o 12.x (en futuras versiones de CUDA se añadirá soporte para arquitecturas de GPU adicionales)
- Controlador NVIDIA R580 o posterior (se requiere R590 para soporte de herramientas de desarrollador específicas de tiles)
- CUDA Toolkit 13.1 o posterior
- Python versión 3.10 o superior
- paquete cuTile BASIC
Comienza a programar
Una vez que tengas todo el software, consulta la documentación completa de cuTile BASIC, prueba todos los programas de ejemplo disponibles en GitHub, y comienza a programar en cuTile BASIC hoy. Aprovecha la oportunidad de portar tu código moderno de IA o computación científica a un lenguaje históricamente importante, manteniendo la capacidad de ejecutarse en el hardware más potente disponible. ¡Solo no le digas a tu Commodore 64!
CUDA Tile en cualquier lenguaje
Aunque BASIC puede no ser el primer lenguaje que los desarrolladores consideran para computación paralela de alto rendimiento, es una demostración instructiva de que, gracias al diseño de la pila de software CUDA, CUDA Tile podría utilizarse desde casi cualquier lenguaje de programación. Al compilar al formato CUDA Tile IR, CUDA Tile puede ser integrado en prácticamente cualquier lenguaje… ¡incluso BASIC!
Nota del editor: En retrospectiva, los desarrolladores que pidieron un amplio soporte del modelo de programación CUDA Tile deberían haber sido más específicos. Espera el cuTile COBOL en abril de 2027.
