A medida que se han ido incorporando en la tarjeta de vídeo funciones más avanzadas, influidas por la necesidad de satisfacer requisitos más exigentes en la generación de gráficos y también las nuevas API de programación (OpenGL y DirectX), ha surgido un nuevo concepto: el de GPU (Graphics Processor Unit), como analogía de las CPU (Central Processing Unit) o microprocesadores clásicos. En realidad es algo que lleva casi dos décadas evolucionando, aunque para los programadores ajenos al campo de los vídeojuegos pueda considerarse algo nuevo.
Si bien la denominación GPU hace referencia a una arquitectura especializada, dirigida específicamente al tratamiento gráfico, cada vez es mayor el número de aplicaciones que aprovechan la potencia de estos circuitos integrados para otro tipo de propósitos. Es aquí, precisamente, donde cobran protagonismo soluciones como Cg, ATI Stream, OpenCL y CUDA, infraestructuras compuestas de bibliotecas, compiladores y lenguajes que dan a luz a una nueva filosofía de desarrollo: GPGPU (General Purpose computing on Graphics Processing Units), la computación de propósito general usando GPUs en lugar de CPUs.
¿Cuál es la razón de que los programadores se interesen en desarrollar aplicaciones que se ejecuten en una GPU? La respuesta surge por sí sola simplemente aportando algunos datos: las actuales CPU cuentan con cuatro, seis u ocho núcleos y son capaces de ejecutar hasta doce hilos de manera simultánea (dos por núcleo en algunos casos), las GPU más avanzadas disponen de hasta 1024 núcleos de procesamiento y tienen capacidad para ejecutar hasta 128 hilos por procesador, lo que ofrece un total de hilos muy superior. Sistemas compuestos únicamente de ocho tarjetas de vídeo de este tipo están superando, en cuanto a rendimiento se refiere, a clústeres de ordenadores como Blue Gene, formados por 512 nodos con CPUs clásicas.
Programación de GPU
Para aprovechar la potencia de las GPU es necesario contar con herramientas de desarrollo adecuadas, capaces de explotar el alto nivel de paralelismo que ofrecen estos dispositivos. Hasta no hace mucho dichas herramientas eran bastante primitivas, ya que su objetivo era facilitar exclusivamente la programación de shaders que (véase el Curso de shaders en el margen derecho) son pequeños bloques de código que aplican un cierto procesamiento a los vértices de la geometría de una escena y los fragmentos resultantes de la rasterización. Ese código se ejecuta paralelamente en cada núcleo, lo cual permite aplicar un cierto algoritmo masivamente a miles o millones de vértices y píxeles.
Estos bloques de código tienen una longitud generalmente muy limitada y se programan en una suerte de lenguaje ensamblador a medida, por lo que difícilmente pueden aplicarse más que a la función para la que están pensados desde un principio. Existen diferentes versiones, denominadas Shader Models, que han ido evolucionando en paralelo a Microsoft DirectX y OpenGL y que tanto ATI como nVidia han ido implementando en su hardware. Al desarrollar una aplicación gráfica se utiliza una API, como las citadas DirectX u OpenGL, para escribir el código que se ejecutará en la CPU, usando el ensamblador del shader model correspondiente para escribir el código a ejecutar en la GPU. Tanto el tipo de operaciones que puede llevar a cabo ese código como la memoria a la que tiene acceso están limitados.
El desarrollo de Cg por parte de nVidia, hace prácticamente una década, fue un primer avance al facilitar la codificación de funciones a ejecutar en la GPU. En lugar de escribir el código en ensamblador se usa un lenguaje de más alto nivel, similar al C. Una función como la mostrada en el siguiente fragmento se ejecutaría una vez para cada vértice, pero no de manera secuencial sino paralelamente.
void trataVertice(
in float4 posicion : POSITION,
out float4 posicionD : POSITION,
out float4 colorD : COLOR0,
const uniform float4x4 ModelViewMatrix,
const uniform float4 color)
{
posicionD = mul(posicion, ModelViewMatrix);
colorD = color;
}A diferencia de CUDA, no obstante, Cg se dirige específicamente a la generación de gráficos. A medida que el número de núcleos de proceso en una GPU se fue incrementando, y ganando en rendimiento al poder operar con datos en coma flotante, se hizo cada vez más patente la necesidad de aprovechar esa potencia bruta de cálculo para propósitos alternativos, aparte de la evidente aplicación en videojuegos de última generación. Solamente se precisaban herramientas de trabajo de corte más general, con un espectro de aplicación más amplio.
¿Qué es CUDA?
La mayoría de los lenguajes de programación no cuentan con estructuras nativas que faciliten la paralelización de procesos, entendiendo como tales partes de un algoritmo que pueden ser ejecutados de manera simultánea y no como lo que se entiende por procesos en el contexto de un sistema operativo.
Es cierto que existen API y bibliotecas de funciones que facilitan, hasta cierto punto, la programación paralela, pero prácticamente ninguna de ellas está pensada para ejecutar el código explotando una GPU. En la mayoría de los casos lo único que hacen es iniciar varios hilos de ejecución dejando en manos del sistema operativo el reparto de tiempo de proceso entre las unidades con que cuente la CPU. Para trasladar la aplicación a otro tipo de procesador, así como para ampliar o reducir el número de hilos en ejecución, es corriente tener que alterar, o incluso rescribir por completo, el código fuente.
La solución que ofrece CUDA (Compute Unified Device Architecture) es mucho más flexible y potente y, además, se basa en estándares existentes. Los programas se escriben en lenguaje C, no en el ensamblador de un cierto procesador o en un lenguaje especializado como es el caso de Cg. Esto facilita el acceso a un grupo mucho mayor de programadores.
Al desarrollar una aplicación CUDA el programador escribe su código como si fuese a ejecutarse en un único hilo, sin preocuparse de crear y lanzar threads, controlar la sincronización, etc. Ese código será ejecutado posteriormente en un número arbitrario de hilos, asignado cada uno de ellos a un núcleo de proceso, de manera totalmente transparente para el programador. Éste no tendrá que modificar el código fuente, ni siquiera recompilarlo, dependiendo de la arquitectura del hardware donde vaya a ejecutarse.
Incluso existe la posibilidad de recompilar el código fuente, dirigido originalmente a ejecutarse sobre una GPU, para que funcione sobre una CPU clásica, asociando los hilos CUDA a hilos de CPU en lugar de a núcleos de ejecución de GPU. Obviamente el rendimiento será muy inferior ya que el paralelismo al nivel de CPU no es, actualmente, tan masivo como en una GPU.
Componentes de CUDA
Los objetivos planteados en el desarrollo de CUDA han dado como fruto un conjunto de tres componentes, disponibles gratuitamente en Developer Zone de nVidia para versiones de 32 y 64 bits de Windows XP, Windows Vista, Windows 7, múltiples distribuciones de GNU/Linux y Mac OS X.
El controlador CUDA es el componente básico, ya que es el encargado de facilitar la ejecución de los programas y la comunicación entre CPU y GPU. Este controlador se aplica a prácticamente toda la gama GeForce 8XX, 9XX y GTX 2XX y posteriores, así como a la línea de adaptadores Quadro y los procesadores Tesla. En cualquier caso se requiere una cantidad mínima de 256 MB de memoria gráfica para poder funcionar, por lo que en adaptadores con menos memoria no es posible aprovechar CUDA.
Instalado el controlador, el siguiente componente fundamental para el desarrollo de aplicaciones es el toolkit CUDA, compuesto a su vez de un compilador de C llamado nvcc, un depurador específico para GPU, un perfilador de código y una serie de bibliotecas con funciones de utilidad ya predefinidas, entre ellas la implementación de la Transformada rápida de Fourier (FFT) y unas subrutinas básicas de álgebra lineal (BLAS). También en la misma web se encuentran múltiples documentos de introducción a la programación de GPUs con CUDA, manuales de referencia, etc.
El tercer componente de interés es el CUDA Developer SDK, un paquete formado básicamente por código de ejemplo y documentación. Se ofrece más de medio centenar de proyectos en los que se muestra cómo integrar CUDA con DirectX y OpenGL, cómo paralelizar la ejecución de un algoritmo y cómo utilizar las bibliotecas FFT y BLAS para realizar diversos trabajos: generación de un histograma, aplicación de convolución a una señal, operaciones con matrices, etc.
Conjuntamente estos tres componentes ponen al alcance del programador todo lo que necesita para aprender a programar una GPU con CUDA y comenzar a desarrollar sus propias soluciones, apoyándose en código probado como el de los ejemplos facilitados o el de las bibliotecas FFT y BLAS.
Estructura de una aplicación CUDA
El código de un programa escrito para CUDA siempre estará compuesto de dos partes: una cuya ejecución quedará en manos de la CPU y otra que se ejecutará en la GPU. Al código de la primera parte se le denomina código para el host y al de la segunda código para el dispositivo.
Al ser procesado por el compilador nvcc, el programa generará por una parte código objeto para la GPU y por otra código fuente u objeto para la CPU. El código objeto específico para la GPU se denomina cubin. El código fuente para la CPU será procesado por un compilador de C/C++ corriente, enlazando el código cubin como si de un recurso se tratase.
La finalidad del código host es inicializar la aplicación, transfiriendo el código cubin a la GPU, reservando la memoria necesaria en el dispositivo y llevando a la GPU los datos de partida con los que se va a trabajar. Esta parte del código puede escribirse en C o en C++, lo cual permite aprovechar el paradigma de orientación a objetos si se quiere.
El código a ejecutar en el dispositivo debe seguir estrictamente la sintaxis de C, contemplándose algunas extensiones de C++. Normalmente se estructurará en funciones llamadas kernels, cuyas sentencias se ejecutarán en paralelo según la configuración hardware del dispositivo final en el que se ponga en funcionamiento la aplicación.
Lo que hace el entorno de ejecución de CUDA, a grandes rasgos, es aprovechar el conocido como paralelismo de datos o SIMD (Simple Instruction Multiple Data), consistente en dividir la información de entrada, por ejemplo una gran matriz de valores, en tantos bloques como núcleos de procesamiento existan en la GPU. Cada núcleo ejecuta el mismo código, pero recibe unos parámetros distintivos que le permiten saber la parte de los datos sobre los que ha de trabajar.
El listado siguiente corresponde a una función kernel muy sencilla, cuyo objetivo es hallar el producto escalar de una matriz por una constante. La función solamente opera sobre un elemento de la matriz, el que le indica la variable threadIdx.x que identifica el hilo en que está ejecutándose el código. Esta función se ejecutaría paralelamente en todos los núcleos de la GPU, por lo que en un ciclo se obtendría el producto de una gran porción de la matriz o incluso de ésta completa, dependiendo de su tamaño y el número de núcleos disponibles.
__global__ void ProdEscalar(
float* M1,
float* M2,
float Constante)
{
// Se obtiene el número de thread
int i = threadIdx.x;
// y se procesa el dato que corresponde
M2[i] = M1[i] * Constante;
}</textarea>En una CPU moderna, como los Athlon Phenom o Core i7, es posible dividir los datos de entrada en cuatro o seis partes pero sin ninguna garantía de que se procesarán en paralelo, salvo que se programe explícitamente el reparto trabajando a bajo nivel. En una GPU y usando CUDA, por el contrario, esos datos se dividirán en bloques mucho más pequeños, al existir 240, 512, 1024 o más núcleos de procesamiento, garantizándose la ejecución en paralelo si necesidad de recurrir a la programación en ensamblador.