Universidad Nacional de Colombia

CUDA

CUDA

CUDA son las siglas de Compute Unified Device Architecture (Arquitectura Unificada de Dispositivos de Cómputo) que hace referencia a una plataforma de computación en paralelo incluyendo un compilador y un conjunto de herramientas de desarrollo creadas por nVidia que permiten a los programadores usar una variación del lenguaje de programación C para codificar algoritmos en GPU de nVidia. CUDA intenta explotar las ventajas de las GPU frente a las CPU de propósito general utilizando el paralelismo que ofrecen sus múltiples núcleos, que permiten el lanzamiento de un altísimo número de hilos simultáneos. Por ello, si una aplicación está diseñada utilizando numerosos hilos que realizan tareas independientes (que es lo que hacen las GPU al procesar gráficos, su tarea natural), una GPU podrá ofrecer un gran rendimiento en campos que podrían ir desde la biología computacional a la criptografía, por ejemplo.

Unidad de procesamiento gráfico (GPU)

La computación acelerada por GPU es el uso de una unidad de procesamiento de gráficos (GPU, por sus siglas en inglés) junto a una CPU para acelerar el funcionamiento de las aplicaciones de aprendizaje profundo, análisis e ingeniería. Hoy en día, los aceleradores de GPU, una innovación de NVIDIA allá por el 2007, permiten el funcionamiento de centro de datos con eficiencia energética en laboratorios gubernamentales, universidades, empresas, y pequeñas y medianas empresas en todo el mundo. Tienen un papel muy importante en la aceleración de aplicaciones en plataformas, que abarcan desde la inteligencia artificial hasta automóviles, drones y robots. La computación acelerada por GPU permite asignar a la GPU el trabajo de los aspectos de la aplicación donde la computación es más intensiva, mientras que el resto del código se ejecuta en la CPU. Desde la perspectiva del usuario, las aplicaciones se ejecutan de forma mucho más rápida. Una forma sencilla de comprender la diferencia entre una GPU y una CPU es comparar la forma en que procesan las tareas. Una CPU tiene unos cuantos núcleos optimizados para el procesamiento en serie secuencial, mientras que una GPU cuenta con una arquitectura en paralelo enorme que consiste de mieles de núcleos más pequeños y eficaces, y que se diseñaron para resolver varias tareas al mismo tiempo. Aquí fue necesario del uso de conceptos básicos, tales como reserva de memoria en la GPU (cudaMalloc), paso de información entre memorias (cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost), y liberación de memoria (cudaFree). Se realizó una función almacenada en la memoria global de la GPU (mediante la sentencia __global__).

Procedimiento

  1. Extracción de información de los píxeles de la imagen( pasada en los parámetros de ejecución).
    • La librería libpng ofrece varias herramientas para poder leer la imagen y manejar los pixeles de la imagen a nuestro gusto, en este link libpng docj podrá ver algunos ejemplos de la implementacion de esta librería. La función void read_png_file(char* file_name) recibe el nombre de la imagen a procesar. Aca lo que se hace es sacar la informacion util de la imagen tal como el numero de pixeles que tiene la imagen a lo ancho ( width = png_get_image_width(png, info); ) , el numero de pixeles que tiene la imagen a lo largo ( height = png_get_image_height(png, info);, entre otros. Se reserva el espacio de memoria para la matriz (cada bit) y se guarda la dirección de memoria del primer bit en un apuntador de tipo png_bytep : png_bytep * rowPointer; de la siguiente manera : png_read_image(png, rowPointer); . Luego en el main se hará así la llamada a esta funcion: read_png_file(argv[1]); .
    • La imagen es representada en una matriz de tamaño n x m,en donde cada posición tiene un arreglo de tamaño 3 en donde se almacenan los valores determinados de RGB. Una vez se tenga el apuntador rowPointer se pasaran los valores de R, G y B de cada pixel en los arreglos h_rI[x], h_gI[x], h_bI[x] respectivamente
  2. Reservar memoria de los vectores que se necesiten en la memoria de la GPU , mediante cudaMalloc.
    3.  cudaMalloc( ( void **) & d_bI, ( height*width ) );
  3. Copiar los datos de la memoria de la CPU a la memoria de la GPU , mediante cudaMemcpy. Por ejemplo: 
    4. cudaMemcpy(d_gI, h_gI, ( height*width ) , cudaMemcpyHostToDevice);
  4. Para el balanceo de cargas y distribución de datos entre hilos se implementó blockwise. Este método pretende de manera equitativa dividir discretamente la carga y asignar dicha cantidad a los hilos correspondientes para que operen bajo cierto conjunto el conjunto de datos. Como se comentó anteriormente, se divide el numero total de pixeles ( (height) * (width ) ), entre el numero de hilos a lanzar:
    5. divi = (height*width/numthreads) ;

    Así ya tendremos almacenado en la variable divi cuantos pixeles debe procesar cada hilo. Luego lo que nos importa es saber cuales van a ser los rangos de pixeles que procesaran cada uno de los hilos, este calculo lo realizara cada hilo deprendiendo de su id. Esto se explicara en el paso 6.



  5. Lanzamiento del kernel:
    6. myBlur<<< blocksPerGrid,  threadsPerBlock >>>(d_rI, d_gI, d_bI, d_rO, d_gO, d_bO, height,numthreads,begin2,end2,p,inputKernel,width);
    Este llamado ejecutara el kernel: 
    __global__ void myBlur3;


  6. Definición de los rangos de pixeles a procesar por cada hilo, esto lo ejecuta cada hilo ( es decir este codigo va dentro del kernel ).
    • Lo que hace este código es multiplicar el numero de pixeles que le toca a cada hilo multiplicado por el id del hilo y así cada hilo sabrá en que pixel comenzar y en cual terminar. En el caso del ultimo hilo, este tomara los pixeles que queden.
  7. Después de tener claro los rangos, cada hilo itera sobre los arreglos con la información de los pixeles, en este caso se hallará el promedio de todas los pixeles para cada valor RGB y se le asignará el resultado al píxel posición.
  8. Copiar los datos de la memoria de la GPU a la memoria de la CPU , mediante cudaMemcpy, por ejemplo:
    9. err = cudaMemcpy(h_rI, d_rI, size, cudaMemcpyDeviceToHost);
  9. Reescribir la nueva imagen con el nuevo nombre recibido en los parametros : 
    10. Write_png_file(argv[2]);
  10. Liberar la memoria de GPU
    11. cudaFree(d_rI);

Unidad de procesamiento gráfico o GPU (Graphics Processing Unit) es un coprocesador dedicado al procesamiento de gráficos u operaciones de coma flotante, para aligerar la carga de trabajo del procesador central en aplicaciones como los videojuegos o aplicaciones 3D interactivas. De esta forma, mientras gran parte de lo relacionado con los gráficos se procesa en la GPU, la unidad central de procesamiento (CPU) puede dedicarse a otro tipo de cálculos (como la inteligencia artificial o los cálculos mecánicos en el caso de los videojuegos).









Llenando los vectores RGB con la informacion de la imagen
x =0;
for(int c = 0; c < height; c++){
    row = rowPointer[c];
    for(int d=0; d < width ; d++ ){
        wrow = &(row[d*4]);
        desrow = wrow[0];
        desrow2 = wrow[1];
        desrow3 = wrow[2];
        h_rI[x] = desrow;
        h_gI[x] = desrow2;
        h_bI[x] = desrow3;
        x++;
    }
}




Parametros que se envian en el kernel.

  • blocksPerGrid: Es el numero de bloques a lanzar.
  • threadsPerBlock: Es el numero de hilos por bloque a lanzar
  • d_rI, d_gI, d_bI : Son los 3 Vectores anteriormente descritos con la información RGB de los pixeles de la imagen.
  • d_rO, d_gO, d_bO: Son los 3 vectores resultantes ( luego de aplicar el efecto borroso a operar).
  • height: Es el numero de pixeles que tiene la imagen a lo ancho.
  • width: Es el numero de pixeles que tiene la imagen a lo Alto.
  • numthreads: Es el numero de hilos a lanzar.
  • begin2: Es el primer pixel a modificar por el primer hilo.
  • end2: Es el ultimo pixel a modificar el primer hilo.
  • inputKernel: Es el tamaño del kernel.
  • p: Es el radio de suma que tendrá el pixel (Depende de cada Kernel).
Balanceo de carga realizdo por cada hilo
int idx = threadIdx.x;
int istart,iend;
if ( idx!= 0 ) {
    istart = idx * height*width / nthrds;
    iend = (idx +1) * height *width/ nthrds;
    if (iend + p >= height*width){
        iend = (height*width)-p ;
    }
}else{
    istart =begin2;
    iend = end2;
}




Ejemplo del calculo del promedio de los pixeles vecinos para kernel 3
for( i = istart; i < iend; i++){
    if ( i < width ){
        aux=0;
    }
    if ( i+width >= (height*width) ){
        aux=0;
    }
    if(  nrows == 3){
        r_I[i] = (double)(r_I[i]+r_I[i-1]+r_I[i+1]+r_I[i-aux]+r_I[i-1-aux]+r_I[i+1-aux]+r_I[i+aux]+r_I[i-1+aux]+r_I[i+1+aux])/9;
        g_I[i] = (double)(g_I[i]+g_I[i-1]+g_I[i+1]+g_I[i-aux]+g_I[i-1-aux]+g_I[i+1-aux]+g_I[i+aux]+g_I[i-1+aux]+g_I[i+1+aux])/9;
        b_I[i] = (double)(b_I[i]+b_I[i-1]+b_I[i+1]+b_I[i-aux]+b_I[i-1-aux]+b_I[i+1-aux]+b_I[i+aux]+b_I[i-1+aux]+b_I[i+1+aux])/9;
    }
}