Skip to content

Práctica 8: Programación en CUDA

Jump to bottom
Rubén Gran Tejero edited this page Mar 26, 2023 · 1 revision

Objetivos

  • Familiarizarse con la programación en CUDA
  • Modelo de programación CUDA
    • Host side
    • Kernel side
  • Medida de prestaciones
  • Hacer un filtro blurring de una imágen

Requisitos

Para poder realizar esta practica debemos tener credenciales para el acceso a las máquinas del Departamento de Informática e Ingeniería de Sistemas.

Introducción

Desde el año 2010, el éxito de las GPU´s (Graphic Processing Units) como un dispositivo de cálculo de propósito específico ha quedado más que demostrado. Su ámbito de aplicación se ha extendido en todos los ámbitos de los sistemas informáticos. En los centros de supercomputación, por ejemplo, donde el top 10 siempre cuenta con GPUs como parte de capacidad computacional de la instalación (top500). En el ámbito doméstico y de consumo, sobremesa y portátiles, las GPUs han habilitado gráficos fotorealistas en tiempo real para el consumidor medio. Y, finalmente, la utilidad de esta organización masivamente paralela ha demostrado su eficiencia energética ($FLOP/WATT$)lo que le ha convertido en dispositivo calculo estándar en muchas plataformas de sistemas embarcados (tiempo real, edge computing, ...).

El éxito de las GPU se ha debido a que han resultado excepcionalmente eficaces en la ejecución de trabajos masivamente paralelos, y más concretamente, en la ejecución de programas con grandes cantidades de paralelismo de datos. Así pues, las GPUs se pueden entender como un coprocesador de propósito especifico para explotar este paralelismo de datos y, por lo tanto, no encontramos sistemas que sólo dispongan de GPU sino que, además, estos sistemas cuentan con CPU convencionales. Estos sistemas, en los que contamos con diferentes dispositivos de computo se le llama sistemas heterogéneos, y las aplicaciones que ejecutan estos sistemas, si queremos que aprovechen todo el potencial, deben hacer uso de la programación heterogénea.

El modelo de programación heterogénea asigna roles específicos de jefe (master) y trabajador (slave) a cada uno de los dispositivos. Así pues, la CPU tiene el rol de jefa ya que en la misma se ejecutará el hilo de ejecución principal, el correspondiente a la función 'main()' de nuestra aplicación. Por otro lado, la GPU, o aceleradora de regiones con paralelismo de datos, se comportará como trabajadora y unicamente entrará en acción cuando el hilo principal lo ordene. De manera general, las aplicaciones que pueden mejorar en sistemas heterogéneos de este tipo deben contar con regiones de código que exhiban paralelismo, así pues, frente a un paradigma convencional en el que toda la aplicación se ejecuta en una CPU, en el paradigma heterogéneo, las regiones con mucho paralelismo son enviados al dispositivo de calculo paralelo específico (GPU).

Para la programación de sistemas heterogéneos existen distintas alternativas. Algunas de ellas, como OpenCL o SYCL que se postulan como estándares abiertos que pueden ser adoptados por cualquier fabricante de hardware. El problema de estas alternativas es que necesitan interés y soporte por parte de los fabricantes de hardware para que sean modelo de programación de facto en lo chips que venden. Esta circunstancia no siempre ocurre o se extiende a lo largo del tiempo. Por otro lado, otros entornos, surgen directamente a propuesta del fabricante del dispositivo. Así pues AMD y sus RADEON utilizan HIP (Hetergeneous Interface for Portability, HIP) , mientras que, por otro lado, NVIDIA utiliza CUDA (Compute Unified Device Architecture, CUDA)

Entorno de prácticas

Para la presente sesión de prácticas utilizaremos la máquina Berlin (berlin.unizar.es). Esta máquina consta de una CPU multicore de AMD y de una GPU de NVIDIA.

  • AMD EPYC 7313P 16-Core Processor 32 L2 512 KB, RAM: 515140 MB. Cada core tiene 2 hilos. Más información: $promtp> lscpu
  • 2 x NVIDIA A10 Tensor Core 24 GB GDDR6. Con unidades específicas para inteligencia artificial.

Hola mundo en CUDA

Antes de comenzar, y para poder utilizar CUDA en nuestra cuenta de las máquinas del DIIS, tenemos que añadir la palabra cuda al fichero .software que se encuentra en nuestro directorio HOME. De otro modo, no encontraremos los binarios y bibliotecas de CUDA.

Para comenzar compilaremos nuestro primer programa en CUDA, el típico hola mundo. Para ello utilizaremos el fichero hola_mundo.cu, la extensión .cu es característica de CUDA. A continuación se muestra el código:

#include <stdio.h>
#include <stdlib.h>

__global__ void cuda_hello(){
    printf("Hello World from GPU!\n");
}

int main() {
    
    cuda_hello<<<1,1>>>();

    cudaError_t cudaerr = cudaDeviceSynchronize();
    if (cudaerr != cudaSuccess)
        printf("kernel launch failed with error \"%s\".\n",
               cudaGetErrorString(cudaerr));

    return 0;
}

Abrid el fichero fuente y analizad el contenido. Contamos con dos funciones, main hilo principal que se ejecutará en la CPU y ordena acciones a la GPU, y segunda función, cuda_hello, precedida por la palabra clave __global__ y que es lanzada desde el host con la linea cuda_hello<<<1,1>>>();. Otra función que es invocada para controlar a la GPU es cudaDeviceSynchronize().

Pregunta. ¿Para qué sirve la función cudaDeviceSynchronize()? ¿qué ocurre si comentamos dicha linea? ¿por qué?

Para compilar este fuente, utilizad el siguiente comando: $prompt> nvcc hola_mundo.cu -o hola.bin

Escribiendo código CUDA: Suma de Vectores

En este apartado transformaremos un código escrito en secuencial para CPU que realiza la suma de vectores. Realizaremos este proceso en varios pasos, describiendo el modelo de programación heterogénea aplicado a CUDA.

Primer paso: Definición e invocación de la función kernel

Comenzaremos presentando el código secuencial vector_addition.c:

#define N 1024*1024*1024

#include <stdlib.h>
#include <stdio.h>

void vector_add(float *out, float *a, float *b, long int n) {
    for(long int i = 0; i < n; i++){
        out[i] = a[i] + b[i];
    }
}

int main(){
    float *a, *b, *out; 

    printf("%d\n", sizeof(long int));

    // Allocate memory
    a   = (float*)malloc(sizeof(float) * N);
    b   = (float*)malloc(sizeof(float) * N);
    out = (float*)malloc(sizeof(float) * N);

    // Initialize array
    for(long int i = 0; i < N; i++){
        a[i] = 1.0f; b[i] = 2.0f;
    }

    // Main function
    vector_add(out, a, b, N);
}

Tarea. Una vez analizado, añade el código para medir el tiempo de ejecución de la suma. A continuación implementa un versión paralela utilizando openMP y mide tambien tiempo.

Para realizar la conversión del código anterior a cuda, seguiremos los siguientes pasos:

  1. Copia de ese mismo fichero a uno con el mismo nombre y extensión .cu
  2. Cambiamos la declaración de la función vector_add y la invocación desde main.
    • Añadimos la palabra clave __global__ delante de void. De este modo haremos visible al compilador nvcc qué código del fichero fuente debe ser compilado para la GPU. El código que se ejecuta en la GPU también recibe el nombre de kernel de cálculo. En segundo lugar, se explicitan qué datos (parámetros) va a necesitar a través de la declaración de la función. Y finalmente, define un punto de entrada para la ejecución en la GPU.
__global__ void vector_add(float *out, float *a, float *b, long int n) {
    for(long int i = 0; i < n; i++){
        out[i] = a[i] + b[i];
    }
}
* Invocación a función. Desde el programa principal, función ```main()```, debemos indicar cuando queremos iniciar la ejecución de la función kernel (vector_add) en la GPU. Debemos explicitar los argumentos de la función para que identifique de dónde debe leer los datos de entrada y dónde debe escribir los resultados, así como parámetros para configurar la ejecución del propio kernel. El lanzamiento o inicio del kernel es una orden o *command* que será enviado a una cola de comandos de la GPU. Allí, las ordenes deben esperar en orden FIFO (modo por defecto) a que la GPU pueda atenderlas. Para evitar penalizar el progreso del hilo de ejecución en el *host*, la invocación del kernel es una orden asíncrona. Es decir, el hilo host puede haber rebasado la sentencia que encola el comando de invocación de kernel, pero eso no implica que el kernel haya finalizado, se este ejecutando, o que ni siquiera la ejecución del mismo se haya iniciado. A continuación se muestra cómo se invoca el kernel desde la función main:

    vector_add<<<1,1>>>(out, a, b, N);

Segundo Paso: Espacio de memoria en la GPU y transferencias de datos entre el Host y Device

  1. Esta función trabaja con vectores de floats. Recordemos que host y device tienen espacios de direccionamiento distintos [^1], con lo cual, se hace necesario reservar memoria en la GPU para pode utilizar los vectores a, b y c. Los malloc() que hay actualmente en el código sólo reservan memoria en el espacio de memoria de la CPU. Para reservar memoria en la GPU utilizaremos cudaMalloc().

[^1]: En el caso de aceleradores discretos, a través de bus PCIe, los módulos de memoria son físicamente independientes. Por el contrario, en aceleradores integrados en el mismo chip, los módulos de memoria son compartidos pero utilizan espacios de direccionamiento distintos.

    float *d_a, *d_b, *d_out;
    
    cudaMalloc((void **) &d_a, sizeof(float)*N);
    cudaMalloc((void **) &d_b, sizeof(float)*N);
    cudaMalloc((void **) &d_out, sizeof(float)*N);

    . . .

    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_out);

El primer parámetro de cudaMalloc es de salida, es decir, el runtime de CUDA devolverá un puntero a una estructura interna que identifica la región de memoria creada en la memoria del dispositivo. Por ese motivo ese primer argumento requiere una referencia a un puntero. Lo cual en lenguaje C se representa como un doble puntero (void **). El segundo parámetro de la función es el tamaño en bytes.

  1. Mover datos entre los distintos espacios de direccionamiento. Los datos de entrada a la función kernel deben ser copiados a memoria del dispositivo acelerador antes de lanzar a ejecutar el kernel. Así mismo, antes de poder leer los resultados en el host thread, debe realizarse una copia de los resultados a un buffer visible por el host. Para realizar estas copias de información utilizaremos la función cudaMemcpy.
    // Copy source data from host to device
    cudaerr = cudaMemcpy (d_a, a, sizeof(float)*N, cudaMemcpyHostToDevice);
    cudaerr = cudaMemcpy (d_b, b, sizeof(float)*N, cudaMemcpyHostToDevice);

    // Main function
    vector_add<<<1,1>>>(d_out, d_a, d_b, N);

    // Copy result data from device to host 
    cudaerr = cudaMemcpy (out, d_out, sizeof(float)*N, cudaMemcpyDeviceToHost);

Para finalizar este apartado, mediremos el rendimiento del kernel ejecutado. Debemos recordar que el lanzamiento del kernel es una tarea asíncrona, y por lo tanto, debemos asegurar que el kernel ha terminado de ejecutarse para poder medir el tiempo que realmente ha requerido el kernel para ejecutarse. La primera alternativa es la siguiente:

    cudaDeviceSynchronize();  // Con esto detenemos la ejecución del host thread hasta que se procese el
                              // ultimo comando. Por si acaso hubiese comandos pendientes en la cola.
    t1 = myCPUTimer();
    vector_add<<<1,1>>>(d_out, d_a, d_b, N);
    cudaDeviceSynchronize();  // Con esto detenemos la ejecución del host thread hasta que se procese el
                              // ultimo comando
    t2 = myCPUTimer();

La segunda opción es utilizar los eventos de cuda. Estos eventos pueden encolarse en la cola de comandos, y cada vez que uno alcanza la cabeza de la cola se le agrega al evento el una marca temporal. Un posible código sería el siguiente:

    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);

    cudaEventRecord(start);
    vector_add<<<1,1>>>(d_out, d_a, d_b, N);
    cudaEventRecord(stop);

    cudaEventSynchronize(stop);
    float milliseconds = 0;
    cudaEventElapsedTime(&milliseconds, start, stop);

Tarea. Comparar estos resultados con los obtenidos en el resto del versiones: secuencial y paralela.

También hay otras herramientas a las que le podéis echar un vistazo como time, nvidia-smio ncu.

Tercer paso: Paralelizar el código en la GPU

Si nos fijamos en el lanzamiento que hace el host del kernel:

    vector_add<<<1,1>>>(d_out, d_a, d_b, N);

Se puede ver que los argumentos que planifican la creación de hilos es: <<<1,1>>>. Esto quiere decir que únicamente se creará un hilo, es decir, sólo habrá un hilo ejecutando el código de la función kernel:

__global__ void vector_add(float *out, float *a, float *b, long int n) {
    for(long int i = 0; i < n; i++){
        out[i] = a[i] + b[i];
    }
}

Esto quiere decir que, hasta ahora, se ha estado ejecutando un sólo hilo que se ha encargado de ejecutar todo el trabajo.

La estrategia de paralelización en CUDA consiste en crear un conjunto de threads (muchos) y que cada uno de ellos ejecute la misma función kernel. En el caso anterior, <<<1,1>>>, sólo se crea un hilo en la GPU que se encarga de ejecutar todo el trabajo del bucle. Este parámetro de invocación <<<M,T>>> ayuda a estructurar la cantidad masiva de threads que se crean por el runtime de CUDA. El parámetro M indica en cuantos paquetes de threads se van a crear, por otro lado, el parámetro T indica cuantos threads hay en cada uno de esos paquetes. Así pues, el total de threads planificados es MxT.

Cada uno de los múltiples hilos que ejecuta la función kernel puede identificarse para adaptar la ejecución de cada hilo a una región del espacio de trabajo que requiere el problema. Las variables que podemos utilizar dentro del kernel para idientifir cada hilo son los siguientes:

  • threadIdx.x contiene el índice del hilo dentro del bloque
  • blockDim.x contiene el número de hilos que hay dentro de cada bloque
  • blockIdx.x contiene el índice del bloque al que pertenece el hilo
  • gridDim.x contiene el número de bloques en total

La estrategia de paralelización para el problema va a consistir en crear un hilo por cada una de las componentes del los vectores, es decir, un total de N threads. Vamos a tomar seleccionar un tamaño de 512 hilos en cada uno de los bloques. Así pues, podemos utilizar la siguiente invocación para lanzar el kernel:

    // Asumiremos que N es división entera de 512
    int block_size = 512;
    int grid_size = (N / block_size);
    vector_add<<<grid_size,block_size>>>(d_out, d_a, d_b, N);

Por otro lado, en la función del kernel tenemos que calcular qué posición del vector debe calcular cada uno de los hilos creados. Una alternativa de código sería la siguiente:

__global__ void vector_add(float *out, float *a, float *b, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    
    out[tid] = a[tid] + b[tid];
}

Podéis explorar los valores de M y T en <<<M,T>>>.

Ejercicio

Implementar ...

Clone this wiki locally