Aprovechamiento de las funciones asíncronas en la reducción del tiempo de ejecución sobre CUDA (Parte I)

APROVECHAMIENTO DE LAS FUNCIONES ASÍNCRONAS EN LA REDUCCIÓN DEL TIEMPO DE EJECUCIÓN SOBRE CUDA.

(Parte I)

Samuel Rodríguez Bernabeu. 
2/Agosto/2012.


*Gracias a la fcsc por hacerme posible el acceso a los recursos computacionales necesarios. Para el ejercicio se han empleado dos tarjetas nvidia Tesla M1060.
Página web de la fundación www.fcsc.es
Aprovechamiento de las funciones asíncronas en la reducción del tiempo de ejecución.
Este pequeño ejercicio pretende sacar partido de las optimizaciones ajenas al kernel, concretamente las que implican a las funciones de carácter asíncrono en la transferencia y la ejecución de cálculos en la GPU.
Las funciones asíncronas, tras iniciarse, devuelven el control al hilo principal de ejecución en el host, pudiéndose de esta manera ejecutar código en la CPU mientras se realizan tareas que involucran a la GPU. Este tipo de técnicas se conocen como “overlap” o “solapamiento”.
La intención de este documento es tomar un kernel y una función en CPU cualesquiera, hacer las mediciones y comparaciones pertinentes y a partir de ello calcular qué volumen de cálculo podemos ocultar mediante estas técnicas y cómo debería ser el planteamiento del programa para ocultarlos.
Para calcular el reparto de volumen de cálculo que asignaremos a la CPU y a la GPU (supondremos que tenemos dentro de un mismo hilo un core de CPU y una GPU) se van a considerar los siguientes parámetros:
    • Tiempo empleado en la transferencia de datos (vía Bus PCI).
    • Tiempo de ejecución del Kernel.
    • Tiempo de ejecución del proceso en la CPU.

Tiempos medios de transferencia a la GPU mediante PCI Express.

Para estimar el porcentaje de volumen de cálculo que debemos repartir entre la GPU y la CPU con el objetivo de aprovechar la ejecución asíncrona del kernel debemos estimar el tiempo empleado en la transferencia de datos a la GPU vía PCI Express.
Con tal fin vamos a medir los tiempos de transferencia para diferentes volúmenes de datos, tomaremos estos tiempos realizando la media aritmética de un número representativo de transferencias, en este caso, 1k transferencias.
El proceso de volcado de datos a la memoria global del device es uno de los puntos clave cuando buscamos minimizar el tiempo de ejecución de nuestra aplicación, y ahora veremos que los tiempos son realmente grandes cuando el volumen de datos comienza a ser significativo.
pastedGraphic.pdf
Como puede apreciarse el tiempo de transferencia sigue un comportamiento bastante lineal, por lo que puede servirnos por el momento la ecuación de la recta que se presenta en la gráfica.
El espacio entre la gráfica y el eje horizontal es tiempo que podemos aprovechar para realizar cálculos en la CPU mientras se realiza la copia, lo que podemos conseguir si realizamos transferencias asíncronas, aunque como veremos más adelante no todo este tiempo es aprovechable.

Si la gráfica sigue la tendencia, todo apunta a ello, cuando estemos explotando la capacidad máxima de la memoria global de la GPU tendremos al procesador ocioso durante un tiempo realmente grande.

Tiempo de ejecución del kernel.

De nuevo tomamos para obtener nuestro tiempo una muestra representativa de datos y obtenemos los resultados que se presentan en la gráfica.
pastedGraphic_1.pdf
En este punto se deben resaltar dos cosas:
  • El tiempo de ejecución de este kernel en concreto es muy bajo comparado con el tiempo de transferencia necesario para llevar el volumen de datos correspondiente la la memoria global de nuestra GPU, lo cual nos lleva a confirmar rápidamente que, tal como indica nvidia, el pci es el punto más susceptible a originar “cuellos de botella”.
  • Al pasar de los 0.3 Gb tenemos muchos más bloques (aún colocando el máximo número de hilos posibles por bloque) de los que la arquitectura permite lanzar. Aprovechando que la memoria global tiene la permanencia del programa he solucionado el problema lanzando el kernel recursivamente, solucionando debidamente el problema de los índices pasando un parámetro adicional.
  • A pesar del “truco” para solucionar el problema del máximo número de bloques, el comportamiento del kernel al aumentar el número de operaciones a realizar es cuasi lineal.

Tiempo de ejecución en la CPU.

Al tratarse de un caso sencillo -suma de elementos contenidos en dos vectores y copia en un tercero- podemos hacer una comparación bastante directa entre los datos, ya que no existen sentencias condicionales u otras que otorguen ventaja a la CPU frente a la GPU y viceversa.
Así pues los tiempos obtenidos son los siguientes:
pastedGraphic_2.pdf

De aquí sólo llamar la atención sobre la diferencia de tiempo que emplean la CPU y la GPU para completar esta tarea en concreto.
Para los fines de este trabajo no importa lo optimizado que sea el kernel o el código de la CPU, ya que sólo se trata de explotar los tiempos muertos de las funciones que tienen que ver con la GPU.


Tiempo de ejecución del cuerpo del programa de cálculo.

Se ha medido aquí la estructura del programa, el tiempo que invierte el código en declarar las variables, reservar espacio para los arrays, liberarlos, comprobar errores… etc. No se realiza en este punto cálculo alguno.
pastedGraphic_3.pdf

Comparativa CPU vs GPU.

Con el fin de comprender mejor la ventaja del paralelismo en esta aplicación, se presenta la siguiente gráfica. En este caso, se ha sumado al tiempo del kernel el tiempo de transferencia asociado a su volumen de datos.
pastedGraphic_4.pdfpastedGraphic_5.pdf
Podríamos decir en este punto que el tiempo que toma la ejecución del kernel no es significativo cuando el volumen de datos es grande (a penas 50 ms frente a unos 2000 ms que emplea la CPU o 1000 ms de la transferencia en el caso de los 2.5 Gb ).


Estrategia de optimización.

Una vez disponemos de los datos, podemos calcular el tiempo final de la aplicación tras ocultar la latencia de la transmisión de datos y de la ejecución del kernel. Para puede procederse como sigue; teniendo un volumen de datos fijado queremos calcular la fracción que tenemos que asignar en cada ciclo a la CPU y a la GPU
  • Calcular el volumen de datos que estamos manejando, para ello hemos planteado las gráficas de esa forma.
  • Plantear un sistema de ecuaciones como el siguiente:
    • y1 = y2
    • x1 + x2 = volumen de datos a particional
    • Ecuación de aproximación del tiempo de ejecución en la CPU.
    • Ecuación de aproximación del tiempo de ejecución en la GPU.
  • Tras resolver el sistema tendremos las fracciones de datos con las que debe trabajar cada uno de los dispositivos.
El solapamiento entre transferencia y cálculo es realmente efectivo cuando planteamos una estrategia secuencial de transferencia-cálculo-transferencia valiéndonos de diferentes streams, pero ahora vamos a emplear solamente transferencias en un mismo stream ya que las otras requieren de implementaciones más complejas.
El módulo de construcción básico es el siguiente:
cudaMemcpyAsync (destino, fuente, tamaño, cudaMemcpyHostToDevice, id_stream);
kernel_suma <<<grid, block>>> (argumentos_kernel);
calculo_en_CPU();

Los datos de nuestro experimento:
– Volumen total de datos a procesar en cada ciclo: 1.8 Gb.
  • Transferencia de 1.2 Gb del Host al Device: Tiempo 482 ms.
  • Tiempo empleado por la GPU en realizar los cálculos para 1.2 Gb de datos; 24 ms.
  • En 482 ms nuestro programa en CPU puede efectuar cálculos sobre un volumen de datos de 0.6 Gb. 

Resultados obtenidos:

En el mejor de los casos podríamos procesar todo el volumen de datos en la GPU (ya que ha demostrado ser más rápida en este problema) y deberíamos hacer una transferencia de 1.8 Gb, tomando el proceso total (sin contar la copia de vuelta de los datos) 1460 ms.
Si decidimos usar tanto la CPU como la GPU sin emplear solapamiento tendríamos en el mejor de los casos 989’37 ms. Un tiempo significativamente menor.
Si empleamos tanto la CPU como la GPU pero empleando técnicas de solapamiento el programa encargado de medir el tiempo para el programa completo ha registrado una mejora de 198,34 ms respecto al caso anterior, esto es un tiempo total de 791,03 ms.
Llama la atención el hecho de que no se ha ocultado el tiempo completo de ejecución de la función en la CPU, no tengo muy claro a qué se debe este hecho. En algunos diagramas de nvidia aparece solapado y en otros, como el de esta página no. Quizás el empleo de streams juegue un papel importante, pero en nuestro código no los he implementado.

Conclusiones:

Aún empleando técnicas simples se han obtenido muy buenos resultados empleando técnicas de solapamiento que hacen posible las transferencias asíncronas de datos entre la CPU y la GPU y que permiten ejecutar código tanto en el host como en el device mientras se completan.
Sin duda, un estudio con más tiempo hubiese permitido depurar el modo en el que se realizan estas transferencias y encontrar un balance más adecuado entre las variables.
Este tipo de optimizaciones son ajenas al kernel, no tienen que ver con la memoria compartida o la coalescendia de datos, ni las reducciones más eficientes, pero ayudan a comprender cómo ese mismo kernel puede ejecutarse de una forma más eficiente. 
pastedGraphic_6.pdf
pastedGraphic_7.pdf
Esta entrada fue publicada en C/C++, CUDA, Supercomputador. Guarda el enlace permanente.

Responder

Introduce tus datos o haz clic en un icono para iniciar sesión:

Logo de WordPress.com

Estás comentando usando tu cuenta de WordPress.com. Cerrar sesión / Cambiar )

Imagen de Twitter

Estás comentando usando tu cuenta de Twitter. Cerrar sesión / Cambiar )

Foto de Facebook

Estás comentando usando tu cuenta de Facebook. Cerrar sesión / Cambiar )

Google+ photo

Estás comentando usando tu cuenta de Google+. Cerrar sesión / Cambiar )

Conectando a %s