El moderno conjunto de herramientas CUDA en la práctica
CUDA sigue siendo la base de la computación acelerada por GPU y sostiene desde simulaciones científicas hasta el entrenamiento de modelos de inteligencia artificial a gran escala.
Pero escribir código CUDA correcto, mantenible y de buen rendimiento puede ser difícil: los errores de memoria se esconden a plena vista, los cuellos de botella de rendimiento resultan invisibles sin la instrumentación adecuada, y los algoritmos de GPU escritos a mano rara vez igualan la eficiencia de las bibliotecas optimizadas. Por suerte, la cadena de herramientas moderna de CUDA maduró de manera significativa y muchos de esos problemas hoy tienen soluciones directas.
¿Qué cubre el recorrido?
En seis pasos incrementales, y con cambios de pocas líneas en cada uno, el artículo hace el código de ejemplo más seguro, más fácil de mantener y más rápido:
- Cómo encontrar fácilmente errores de indexación adoptando la API moderna de CCCL y Compute Sanitizer
- Cómo mejorar las mediciones de Nsight Systems con NVTX
- Cómo usar los algoritmos optimizados de CUB a nivel de bloque y de dispositivo
- Cómo administrar la memoria de la GPU mediante contenedores agrupados
- Cómo acelerar las transferencias desde el anfitrión hacia el dispositivo con contenedores anclados
- Cómo paralelizar el trabajo en GPU dándole a cada hilo su propio flujo y transferencias asincrónicas
Como complemento, NVIDIA publica el código y la opción de ejecutarlo en Google Colab.
Punto de partida: una tubería de procesamiento de imágenes
A partir de un flujo de entrada de imágenes en rojo, verde y azul, el ejercicio comienza transfiriendo los datos desde la CPU hacia la GPU. Después convierte esas imágenes de RGB a escala de grises.
Luego, para cada baldosa de 32 por 32 píxeles de la imagen, calcula la mediana ordenando los píxeles y seleccionando el valor del medio. Por último, copia la mediana de cada baldosa de vuelta a la CPU.
El código base define dos kernels. El primero, computeRGBToGray, carga los valores de las imágenes de entrada roja, verde y azul, los convierte y los escribe en la imagen de salida en escala de grises con los coeficientes clásicos de luminancia: 0,299 para el rojo, 0,587 para el verde y 0,114 para el azul. El segundo, computeMedian, calcula la mediana de cada baldosa de la imagen en escala de grises. Cada bloque de hilos carga su baldosa desde memoria global hacia memoria compartida y después un único hilo ordena el arreglo y escribe el valor del índice central en el arreglo global de medianas.
En la función principal, tras definir las constantes del ejemplo, se reserva la memoria de CPU para cada imagen y para las medianas. La tubería corre después para cada una de las tres imágenes, en paralelo, usando OpenMP. Empieza reservando la memoria necesaria en la GPU antes de transferir los datos desde la CPU, lanza los dos kernels y finalmente copia los resultados de vuelta al anfitrión antes de liberar la memoria. Con 250 baldosas por lado y baldosas de 32 píxeles, cada imagen del ejemplo mide 8.000 por 8.000 píxeles.
Ese código tiene varios defectos, que el recorrido va corrigiendo uno por uno.
¿Cómo se caza un error de memoria compartida?
Al ejecutar el programa aparece un mensaje poco informativo.
code_steps$ ./build/0_base_error_example
CUDA error in 0_base_error_example.cu at line 105: an illegal memory access was encounteredAunque el código incluye verificación de errores, frente a un mensaje de acceso ilegal a memoria conviene partir por Compute Sanitizer, la suite de verificación funcional de NVIDIA, que identifica de inmediato un error difícil de ver a simple vista:
$ compute-sanitizer ./build/0_base_error_example
========= COMPUTE-SANITIZER
========= Invalid __shared__ write of size 1 bytes
========= at void computeMedian<(int)32, (int)256>(unsigned char *, unsigned char *, int, int)+0x170 in 0_base_error_example.cu:55
========= by thread (0,3,0) in block (20,0,0)
========= Access at 0x6440 is out of boundsLa salida muestra directamente una escritura fuera de rango en memoria compartida, en la línea 55 del archivo:
tile[index] = d_image_gray[index];La línea carga mal los datos en memoria compartida, porque usa un índice global. Como la memoria compartida está definida a nivel de bloque de hilos, hay que cambiar la indexación. Para evitar ese tipo de equivocaciones se introdujo en CCCL una API nueva que distingue el indexado global del indexado a nivel de bloque. Para usarla, primero hay que lanzar el kernel con la nueva API cuda::launch:
auto config = cuda::make_config(cuda::block_dims(...), cuda::grid_dims(...));
cuda::launch(stream, config, kernel_name<decltype(config)>, input)Y después, dentro del kernel, aplicar la nueva API de indexación. El recorrido completo con los cinco pasos restantes, incluidos NVTX, CUB, los contenedores agrupados y anclados y los flujos por hilo, está disponible en la publicación original junto con el repositorio de código.
¿Qué se lleva de aquí quien no programa GPU a diario?
El patrón que ordena todo el ejercicio no es específico de CUDA: primero corregir, después medir, y solo entonces optimizar. El error de la línea 55 es del tipo que puede sobrevivir meses en producción porque escribe fuera de rango sin necesariamente hacer caer el programa, y ninguna cantidad de ajuste de rendimiento habría servido antes de resolverlo. Que la ordenación de la baldosa la haga un solo hilo de los 1.024 disponibles en el bloque, mediante un ordenamiento de burbuja, es el segundo problema evidente y es exactamente lo que resuelven los algoritmos de CUB en los pasos siguientes.
Vale destacar además que el ejemplo completo se puede correr en Google Colab, sin GPU propia. Para quien aprende computación acelerada en una universidad de la región, sin acceso a un laboratorio con tarjetas dedicadas, esa es la diferencia entre leer el recorrido y ejecutarlo.
