USM Numérica

Librerías de Python: PyCUDA

Objetivos

1. Entender la utilidad de aceleradores (como las tarjetas gráficas).
2. Aprender los fundamentos de la programación en tarjetas gráficas.
3. Conocer la librería PyCUDA

0.1 Instrucciones

Las instrucciones de instalación y uso de un ipython notebook se encuentran en el siguiente link.

Después de descargar y abrir el presente notebook, recuerden:

• Desarrollar los problemas de manera secuencial.
• Guardar constantemente con Ctr-S para evitar sorpresas.
• Reemplazar en las celdas de código donde diga FIX_ME por el código correspondiente.
• Ejecutar cada celda de código utilizando Ctr-Enter

0.2 Licenciamiento y Configuración

Ejecutar la siguiente celda mediante Ctr-Enter.

Contenido

1. Computación de alto desempeño (HPC)
   • Aceleradores
   • El rol de Python en HPC
2. Programación en tarjetas gráficas
   • La tarjeta gráfica
   • Lenguaje CUDA
3. PyCUDA
   • Usando la tarjeta gráfica desde Python
   • Creando tus propios "kernels"

1. Computación de alto desempeño

En muchos problemas de ingeniería (en verdad, en la gran mayoría de ellos) es necesario hacer una gran catidad de cálculos para lograr el objetivo. Por ejemplo, la gente que modela dinámica de fluidos con el computador, necesita resolver una ecuación en un dominio con varios millones de nodos. Otro caso es el procesamiento de datos de observatorios astronómicos, que pueden generar varios terabytes de datos por día. Los computadores convencionales se quedan "corto" en estos casos: no son suficientemente rápidos como para realizar todos los cálculos en un tiempo razonable, y la memoria es muy pequeña.

La computación de alto desempeño, o High Performance Computing (HPC), es la disciplina que se hace cargo de este problema. La principal herramienta que se usa en HPC es la computación paralela, donde se usan varios procesadores coordinados haciendo cálculos al mismo tiempo. Otra opción, que se puede usar junto a la programación paralela, es el uso de aceleradores. Este notebook trata del uso de tarjetas gráficas (Graphic Processing Units o GPU) como aceleradores, desde Python.

Como su nombre lo dice, historicamente las tarjetas gráficas han sido usadas para hacer la gráfica en la pantalla, sin embargo, desde mediados de los años 2000, comenzaron desarrollos que utilizaron las GPUs para otros propósitos, y nació el concepto General Purpose GPU.

Python y HPC

A pesar que Python es un lenguaje que prioriza facilidad en uso por sobre eficiencia, existen herramientas que permiten hacer computación de alto desempeño con Python. Por ejemplo, la librería mpi4py contiene funciones que permiten operar sobre más de un procesador y coordinarlos. También, la librería petsc4py da acceso a PETSc, una librería de computación paralela desarrollada en C++ para aplicaciones numéricas.

Por parte de aceleradores, existen librerías que permiten acceder a ellos desde Python. PyCUDA y PyOpenCL, desarrollados por Andreas Kloeckner de University of Illinois Urbana-Champaign, son dos ejemplos de esto. También, la versión pagada de la distribución de Python de Numba puede usar la GPU. En este notebook vamos a revisar las bases de PyCUDA.

2. Programación en tarjetas gráficas

2.1 La tarjeta gráfica

La tarjeta gráfica es un procesador con muchos núcleos. Nosotros estamos acostumbrados a hablar de un procesador (o CPU) de 2 o 4 núcleos, la GPU tiene miles. Eso si, los núcleos de la GPU son menos "inteligentes" que los de un procesador común: la GPU es un procesador de gran capacidad para hacer muchos cómputos simples de forma paralela. Si el problema puede ser descompuesto en muchos cálculos simples, la GPU es una muy buena alternativa, en cambio, si el código es complicado (por ejemplo, contiene muchos if), es probable que un procesador común tenga mejor desempeño.

Figura obtenida del CUDA Programming Guide

2.2 CUDA (Computer Unified Device Architecture)

Existen dos lenguajes para programar tarjetas gráficas: CUDA y OpenCL. Ambos funcionan de manera muy parecida, y en este caso nos focalizaremos en CUDA. CUDA es muy parecido al lenguaje C, con palabras clave cuando queremos que el código corra en la GPU.

CUDA reconoce dos dominios o procesadores: host corresponde a la CPU y device a la GPU. Por defecto, el código corre en el host, sin embargo, hay palabras claves que le indica al computador que una función corra en el device. Esta función de conoce como kernel, y se reconocen porque hay que poner la palabra clave __global__ antes de la definición de la función, por ejemplo:

__global__ int suma_1(in a)
{
    int b = a + 1
    return b
}

El kernel es ejecutado en paralelo por cada núcleo de la GPU. Para una mejor organización, CUDA virtualiza los núcleos en hebras (o threads), que están organizados en bloques (o blocks). El programador puede especificar cómo organizar los threads (con índices en 1, 2 o 3 dimensiones), y los blocks (con índices en 1 o 2 dimensiones)

Figura obtenida del CUDA Programming Guide

Dentro del kernel tenemos acceso al número del thread y block que está ejecutando el kernel con threadIdx.x y blockIdx.x, respectivamente. De esta forma, podemos hacer que ls threads operen sobre diferentes partes de la data. Por ejemplo, si queremos sumarle 1 a un arreglo de tamaño N, en vez de hacer un loop for que corra por todo el arreglo y sume 1, podemos ejecutar un kernel en la GPU donde el thread 0 opere sobre el primer elemento del arreglo, el tread 1 sobre el segundo, y así sucesivamente. En C/C++ una función que sume 1 a cada elemento de un arreglo A y guarda el resultado en el arreglo B sería

void suma_1(float *A, float *B, int N)
{
    for (int i=0; i<N; i++)
        B[i] = A[i] + 1; 
}

en cambio, de correr en la GPU no necesitamos el loop y quedaría

__global__ void suma_1(float *A, float *C, int N)
{
    int i = threadIdx.x + blockDim.x*blockIdx.x;

    if (i<N)
        B[i] = A[i] + 1; 
}

El cambio parece muy fácil, sin embargo hay un pequeño detalle: la memoria de la GPU y la CPU no es la misma. Por lo tanto, antes de correr el kernel, debemos mover la data que necesitaremos a la memoria de la GPU (en el ejemplo, A y B):

float *A, *B;
cudaMalloc( (void**) &A, N*sizeof(float));
cudaMalloc( (void**) &B, N*sizeof(float));
cudaMemcpy(A, A_h, N*sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(B, B_h, N*sizeof(float), cudaMemcpyHostToDevice);

donde A_h y B_h son los arreglos que están en la memoria de la CPU.

y una vez finalizada la ejecución del kernel, debemos traer el resultado a la CPU:

cudaMemcpy(B, B_d, N*sizeof(float), cudaMemcpyDeviceToHost);

La llamada al kernel también tiene una pequeña diferencia con C:

int blocks = int N/256;
suma_1<<<blocks, 256>>> (A, B, N);

donde N es el número de datos en el arreglo A, 256 es el número de threads que contiene cada block y blocks es el número de bloques necesarios para poder operar sobre toda la data. Pareciera ser que elegimos 256 de forma aleatoria, y de cierta manera, es verdad, sin embargo, generalmente funciona bien pues es un múltiplo del número de núcleos en la GPU.

¿Más detalles? Vean la guía de programación en CUDA

3. Accediendo a la tarjeta gráfica desde Python: PyCUDA

Ya vimos que CUDA nos permite acceder a la GPU desde C; PyCUDA nos permite hacerlo desde Python. Antes de cualquier cosa, nos tenemos que hacer dos preguntas:

• ¿Mi GPU tiene soporte para CUDA? Hay una lista de los modelos de GPU que soportan CUDA. Ve si la tarjeta gráfica en tu computador está en esa lista!
• ¿Cómo instalo PyCUDA? Revisen las instrucciones en página de PyCUDA.

Existen dos formas de acceder a la GPU usando PyCUDA: usando funciones prehechas que usan la GPU, o generando sus propios kernels. Veamos estas dos por separado.

3.1 Funciones de PyCUDA

PyCUDA viene con una interfaz tipo numpy, cuyas funciones corren en la GPU. Esta es la forma más fácil de usar la GPU, ya que no es necesario hacer un kernel. Existe una lista con las funciones disponibles en PyCUDA usando esta interfaz.

Para decirle a Python que vamos a usar la GPU, debemos importar autoinit

y las principales funciones de PyCUDA están en el driver

Además, así como numpy, PyCUDA tiene su propia estructura de datos. Por ejemplo, lo equivalente a un arreglo de numpy es un gpuarray

Cuando un arreglo es un gpuarray, se puede mover la data a la GPU con la función to_gpu. Por ejemplo, si creamos un arreglo de numpy lleno de 1 y queremos mover esa data a la GPU hacemos

Podemos operar sobre a_gpu como si fuera un arreglo de numpy. Por ejemplo, para sumar 1 a su valor

y luego podemos traerlo a la CPU con get() para imprimir su valor

Desafío 1

Hagan un arreglo con $1000$ números equidistantes entre $0$ y $2\pi$, y calculen la función $\sin(x)$ en la GPU usando las funciones de PyCUDA

3.2 Generando nuestros propios kernels

La cantidad de funciones disponibles en PyCUDA es limitada, y es muy probable que, de necesitar algo específico, sea necesario hacer un kernel especial. En PyCUDA el kernel es exactamente igual que en CUDA, y se escribe dentro de un string. PyCUDA usa el compilador de CUDA, llamado nvcc, para compilar lo que está dentro del string cuando se ejecuta el código Python. PyCUDA guarda el código compilado, así es que solamente recompila el string cuando ve que hay una modificación.

Usemos el mismo ejemplo que antes, y sumemos 1 a los valores de un arreglo, esta vez, haciendo el kernel a mano. La función que compila el string con el kernel se llama SourceModule, y hacemos a esa función accesible por Python con get_function

y podemos ejecutar el kernel de forma similar a CUDA, especificando el número de threads y blocks

La palabra clave block cuando invocamos el kernel reemplaza al <<<blocks,256>>> que usamos en CUDA.

Desafío 2

Hagan el mismo cálculo que en el Desafío 1, esta vez, haciendo ustedes mismos el kernel