Computación Concurrente, Paralela y Distribuida. PDF

Computación Concurrente, Paralela y Distribuida. Tema 4: Introducción a los aceleradores Hardware. Javier López Fandiño Departamento de Electrónica e Computación Universidade de Santiago de Compostela Índice Bibliografı́a recomendada Objetivos Computación heterogénea: Acelera...

Computación Concurrente, Paralela y Distribuida. Tema 4: Introducción a los aceleradores Hardware. Javier López Fandiño Departamento de Electrónica e Computación Universidade de Santiago de Compostela Índice Bibliografı́a recomendada Objetivos Computación heterogénea: Aceleradores Hardware Programación en GPU: CUDA Repaso 1 Bibliografı́a recomendada Bibliografı́a Zaccone, G. (2019). Python parallel programming cookbook. Packt Publishing Ltd. Capı́tulo 8. Zaccone, G. (2015). Python parallel programming cookbook. Packt Publishing Ltd. Capı́tulo 6. Kirk, D. B., & Wen-mei, W. H. (2013). Programming massively parallel processors: a hands-on approach. Morgan Kaufmann Publishers. Capı́tulos 1-6. 2 Objetivos Objetivos Objetivo principal: Estudiar el paradigma de computación heterogénea mediante el uso de aceleradores Hardware. Se introducirá el paradigma de computación heterogénea. Se revisarán los distintos tipos de aceleradores Hardware disponibles. Se introducirá el modelo de programación en GPU mediante el uso de CUDA. 3 Computación heterogénea: Aceleradores Hardware Computación heterogénea Definición Computación que utiliza más de un tipo de procesador. Algunas posibilidades: Misma familia de procesador con distintas configuraciones o caracterı́sticas: frecuencia, nº de cores, cachés, etc. Co-procesador dedicado a tareas especı́ficas: gráficos, red, sensores. *En muchos contextos, sólo se considera que existe heterogeneidad si los procesadores tienen distintas ISA. 4 Ejemplos de co-procesadores GPUs: Especializadas en cálculos matriciales, lo que las hace particularmente adecuadas para procesado de imagen/vı́deo. Intel xeon Phi: Especializados en cálculo masivamente paralelo (OpenMP). Diseñados para ser una alternativa a las GPUs en sistemas HPC. FPGAs: Procesadores reprogramables, adecuados para acelerar tareas especı́ficas como el procesado digital de señal. NPU: Network Processing Unit, chip diseñado para encargarse de la gestión de tareas asociadas a la red (KillerNIC). 5 Ejemplos de co-procesadores En la actualidad, las CPUs (o las placas base) integran módulos que anteriormente estaban desagregados en co-procesadores dedicados: Gráficos. Sonido. Red/wifi. etc. 6 Aceleradores Hardware Hardware añadido que ha sido diseñado para realizar operaciones especı́ficas de forma muy eficiente. A pesar de su alto rendimiento, no está diseñado para funcionar de forma autónoma ⇒ Necesita una CPU. La CPU lleva el control de las tareas de manera secuencial y asigna al acelerador tareas computacionalmente costosas y que presentan un alto grado de paralelismo. Pueden compartir o no su espacio de memoria. Si no la comparten, existe un sobre coste de comunicación entre la CPU y el acelerador. 7 Aceleradores Hardware: Procesadores Masivamente Paralelos (MPP) El ejemplo más común son las GPU. Figura 1: Arquitectura CPU vs Arquitectura GPU. 8 Evolución de los procesadores La Ley de Moore predijo que el número de transistores por unidad de área se duplicarı́a cada 18 a 24 meses. 9 Evolución de los procesadores Hasta 2005, la frecuencia del procesador siguió la misma tendencia debido a la ganancia en velocidad por la miniaturización de los transistores. 9 Evolución de los procesadores La frecuencia dejó de aumentar debido al power wall: Aumentar más la frecuencia harı́a que el chip se calentara demasiado para enfriarse de manera factible. 9 Evolución de los procesadores El estancamiento en la frecuencia provocó un estancamiento en el rendimiento de un solo subproceso, a excepción de mejoras debidas a los avances en la arquitectura y el compilador. 9 Evolución de los procesadores El estancamiento en el rendimiento de un solo subproceso hizo que la computación paralela se generalizara a medida que se dedicaron transistores a agregar más núcleos a los procesadores. 9 Evolución de los procesadores Se pasó, por tanto, de un modelo orientado a minimizar la latencia a un modelo orientado a maximizar el ancho de banda: 10 Evolución de los procesadores Se pasó, por tanto, de un modelo orientado a minimizar la latencia a un modelo orientado a maximizar el ancho de banda: 10 El modelo CPU vs GPU Se pasó, por tanto, de un modelo orientado a minimizar la latencia a un modelo orientado a maximizar el ancho de banda: Vı́deo demostración Nvidia: Mythbusters Versión corta. Versión larga. 11 El modelo CPU vs GPU CPU GPU Pocas ALUS con frecuencias altas: Muchas ALUS pequeñas: Latencia reducida. Ancho de banda elevado. Cachés grandes: Accesos a caché en lugar de accesos a memoria (con alta latencia). Cachés más pequeñas: Mayor área dedicada a cómputo. Control sofisticado: Predicción de saltos. Control simplificado: Adelantamiento de datos para Mayor área dedicada a cómputo. reducir riesgos de datos. 12 ¿Dónde se usan las GPUs? 13 ¿Dónde se usan las GPUs? GPGPU: General-purpose computing on Graphics Processing Units Uso de GPUs para acelerar el cómputo de tareas de propósito general, acompañando o reemplazando a la CPU. Especialmente adecuado en aplicaciones con muchas operaciones matemáticas. Alto grado de paralelismo: La misma secuencia de operaciones repetida sobre grandes conjuntos de datos. El uso de GPUs está generalizado en los supercomputadores de las listas TOP500 y Green500. 14 ¿Dónde se usan las GPUs? TOP 500 Figura 2: TOP500. Noviembre 2023. Uso de aceleradores Hardware. 15 Casos de uso GPUs Por sus caracterı́sticas, las GPU resultan especialmente adecuadas para resolver problemas, entre otros, en los siguientes campos: Procesamiento gráfico/renderizado. Aprendizaje automático, redes neuronales. Deep learning. Criptografı́a. Modelado computacional. Visualización: médica, sensado remoto, etc. Big Data. 16 Casos de uso GPUs Ejemplo: Conversión de una imagen RGB a escala de grises. 17 Stream processing Stream processing: Paralelismo de datos Paradigma de computación centrado en el procesamiento de flujos de datos. Resulta especialmente adecuado cuando: La carga computacional es elevada. Existe paralelismo de datos: La misma función se aplica a todos los elementos de un mismo input de forma simultánea. Existe localidad en los accesos a datos. 18 Casos de uso GPUs: Paralelismo de datos 19 Casos de uso GPUs: Paralelismo de datos 20 Casos de uso GPUs: Paralelismo de datos 21 Casos de uso GPUs: Paralelismo de datos 22 Casos de uso GPUs: Paralelismo de datos 23 Uso de GPU para optimizar el rendimiento de código Uso de librerı́as que exploten las capacidades de las GPU disponibles: CuPy PyTorch Tensorflow Programación directa de código para ser ejecutado en GPU: CUDA. 24 Programación en GPU: CUDA CUDA Compute Unified Device Architecture 1 Plataforma de computación paralela que incluye 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 (CUDA C) para codificar algoritmos en GPU de Nvidia. *Por medio de wrappers se puede usar desde otros lenguajes, como Python. 1 https://developer.nvidia.com/blog/even-easier-introduction-cuda/ 25 CUDA Ventajas: Explota las ventajas de la arquitectura GPU: permite lanzar un número de hilos muy elevado de manera simultánea. Obtiene grandes speedups en aplicaciones que estén diseñadas para utilizar numerosos hilos realizando tareas independientes. 26 CUDA Limitaciones: Sobrecoste de movimiento de datos: Cuello de botella entre la CPU y la GPU. Debido al ancho de banda de los buses y sus latencias. Necesidad de carga computacional elevada: Los threads o hilos de ejecución deben lanzarse en grupos de al menos 32, llamados warps. 27 Kernels Código que se ejecuta en el dispositivo La función que ejecutarán los diferentes flujos durante la fase paralela. En CUDA un kernel se ejecuta mediante un conjunto de flujos, es decir, es una función la cual al ejecutarse lo hará en N distintos hilos en lugar de en secuencial. Se define añadiendo la palabra clave global en la declaración. 1 //Definición del kernel 2 __global__ void add(int* A, int* B, int* C){ 3 int i = threadIdx.x; 4 C[i] = A[i] + B[i]; 5 } 28 Kernels El modelo de ejecución se basa en el uso de un grid de bloques: Los threads se organizan formando bloques que a su vez se organizan formando un grid. Cada thread tendrá un tid único que indicará que tarea le corresponde ejecutar. Este tid se obtiene a partir del id dentro del bloque y del id del bloque dentro del grid. Tanto los grid como los bloques podrán constar de 1, 2 o 3 dimensiones. Se representan como los ejes x, y, z. Las dimensiones de grid y bloque con las que se debe ejecutar un kernel se definen en la llamada a la función. Dentro de la función se identifica el hilo y bloque correspondiente. 29 Kernels 1 int blockSize = 256; 2 int numBlocks = (N + blockSize - 1) / blockSize; 3 add(N, x, y); 30 Arquitectura de una GPU Cada bloque se asignará a un SM para su ejecución La configuración del grid afectará al nivel de paralelismo alcanzado. Figura 3: Organización en Streaming Multiprocessors de una GPU. 31 Warps Unidad mı́nima de ejecución paralela en un SM (Streaming Multiprocessor) Formada por un conjunto de 32 threads con threadIdx consecutivos. Para maximizar el rendimiento debe: Minimizarse la divergencia entre las instrucciones a ejecutar por los threads de un mismo warp. Minimizarse el número de threads inactivos (idle). Maximizarse la coalescencia en el acceso a los datos (Tarea asignada a cada thread). 32 Movimiento de datos en memoria Existen funciones dedicadas al movimiento de datos entre la memoria del host y el dispostivo: cudaMemcpyHostToDevice() cudaMemcpyDeviceToHost() cudaMemcpyDeviceToDevice() 33 Tipos de memoria Registros: La memoria más rápida y pequeña. Sólo accesible por un thread. Su tiempo de vida es igual al de dicho thread. Memoria local: Memoria privada a la que sólo un thread puede acceder. Más lenta que los registros pero de mayor tamaño. Memoria compartida: Cada bloque tiene su propia memoria compartida entre los threads que le pertenecen. Esta memoria es extremadamente rápida. Requiere sincronización entre hilos. Memoria de constantes: todos los threads de un mismo grid tienen acceso a esta memoria, pero sólo se pueden hacer lecturas. Los datos son persistentes durante la ejecución de la aplicación. Memoria global: Todos los threads del grid tienen acceso a la memoria global. Los datos son persistentes durante la ejecución de la aplicación. 34 Tipos de memoria Figura 4: Memoria de una GPU. 35 Tipos de memoria Factores a tener en cuenta: Para obtener el mejor rendimiento, debe usarse adecuadamente cada uno de los niveles de memoria. En particular, resulta de interés maximizar el uso de la memoria compartida, minimizando ası́ los accesos a la memoria global (más lenta). 36 Sincronización Como en cualquier aplicación concurrente, la sincronización entre hilos puede ser necesaria para el correcto funcionamiento de un código. Desde un kernel, se puede explicitar una barrera mediante syncthreads(), que provocará que todos los hilos esperen a que los demás lleguen a ese punto. 37 Flujo de ejecución 1. Reservar memoria en el dispositivo (GPU). 2. Transferir los datos necesarios desde procesador principal o host (CPU) al dispositivo. 3. Invocar la ejecución del kernel(s). 4. Transferir los resultados del dispositivo al host. 5. Liberar la memoria del dispositivo una vez finalizada la ejecución del kernel. 38 Uso de CUDA desde código Python Existen diferentes librerı́as que permiten la ejecución de código CUDA desde Python: PyCUDA 2. Numba: Compilador Python desde Anaconda que puede compilar código Python para su ejecución en GPUs Nvidia. 2 https://documen.tician.de/pycuda/tutorial.html 39 Programación heterogénea con PyCUDA Pycuda es un wrapper para programación en CUDA desde Python que permite aprovechar las capacidades de las GPU utilizando códigos a más alto nivel. La compilación de un código CUDA consiste en los siguientes pasos: 1. El código del dispositivo es separado del del host. 2. Invocar un compilador (gcc) para compilar el código del host. 3. Generar el código del dispositivo en formato binario o ensamblador. 4. Linkar los códigos resultantes para obtener el ejecutable final. La ejecución de todos estos pasos mediante PyCUDA durante la ejecución de un programa, supone una sobrecarga de tiempo respecto a una aplicación CUDA nativa. 40 Google Colab Colaboratory, o Colab, es un producto de Google Research que permite escribir y ejecutar código Python en el navegador. [...] Desde un punto de vista más técnico, Colab es un servicio de cuaderno alojado de Jupyter que no requiere configuración y que ofrece acceso sin coste adicional a recursos informáticos, como GPUs. 3 3 https://research.google.com/colaboratory/intl/es/faq.html 41 Google Colab Configuración para uso de GPUs: 1. Acceder a https://colab.research.google.com/ 2. Crear nuevo cuaderno. 3. Cambiar tipo de entorno de ejecución: T4 GPU. 4. Instalar pycuda: !pip install pycuda. 5. Comprobar el hardware disponible. 42 Obtener capacidades de GPUs disponibles 1 import pycuda.driver as drv 2 drv.init() 3 print("%d device(s) found." % drv.Device.count()) 4 5 for ordinal in range(drv.Device.count()): 6 dev = drv.Device(ordinal) 7 print("Device #%d: %s" % (ordinal, dev.name())) 8 print("Compute Capability: %d.%d"% dev.compute_capability()) 9 print("Total Memory: %s KB" % (dev.total_memory()//(1024))) 43 Obtener capacidades de GPUs disponibles 1 1 device(s) found. 2 Device #0: Tesla T4 3 Compute Capability: 7.5 4 Total Memory: 15464512 KB 44 Ejemplo: Multiplicación de vectores 1 import pycuda.autoinit 2 import pycuda.driver as drv 3 import numpy 4 import time 5 6 SIZE = 150 7 8 from pycuda.compiler import SourceModule 9 mult = SourceModule(""" 10 __global__ void multiply_naive(float *dest, float *a, float *b) 11 { 12 const int i = threadIdx.x; 13 dest[i] = a[i] * b[i]; 14 } 15 """) 16 17 multiply_naive = mult.get_function("multiply_naive") 45 Ejemplo: Multiplicación de vectores 1 a = numpy.random.randn(SIZE).astype(numpy.float32) 2 b = numpy.random.randn(SIZE).astype(numpy.float32) 3 4 start = time.time() 5 check = [a[i] * b[i] for i in range(SIZE)] 6 end = time.time() 7 print(f"Tiempo CPU {end-start}") 8 9 dest = numpy.zeros_like(a) 10 start = time.time() 11 multiply_naive( 12 drv.Out(dest), drv.In(a), drv.In(b), block=(SIZE,1,1), grid=(1,1)) 13 end = time.time() 14 print(f"Tiempo GPU {end-start}") 15 16 print(numpy.array_equal(dest, check)) 46 Ejemplo: Multiplicación de vectores 1 Tiempo CPU 0.00015854835510253906 2 Tiempo GPU 0.004923820495605469 3 True 47 Ejemplo: Multiplicación de vectores SIZE = 1500 1 --------------------------------------------------------------------------- 2 LogicError Traceback (most recent call last) 3 in () 4 37 dest = numpy.zeros_like(a) 5 38 start = time.time() 6 ---> 39 multiply_naive( 7 40 drv.Out(dest), drv.In(a), drv.In(b), 8 41 block=(SIZE,1,1), grid=(1,1)) 9 10 /usr/local/lib/python3.10/dist-packages/pycuda/driver.py in function_call(func, *args, **kwargs) 11 479 raise ValueError("must specify block size") 12 480 13 --> 481 func._set_block_shape(*block) 14 482 handlers, arg_buf = _build_arg_buf(args) 15 483 16 17 LogicError: cuFuncSetBlockShape failed: invalid argument 48 Ejemplo: Multiplicación de vectores por bloques 1 [...] 2 from pycuda.compiler import SourceModule 3 mult = SourceModule(""" 4 __global__ void multiply_blocks(float *dest, float *a, float *b, int n) 5 { 6 int index = blockIdx.x * blockDim.x + threadIdx.x; 7 if (index < n) 8 dest[index] = a[index] * b[index]; 9 } 10 """) 11 12 multiply_blocks = mult.get_function("multiply_blocks") 13 49 Ejemplo: Multiplicación de vectores por bloques 1 [...] 2 3 dest = numpy.zeros_like(a) 4 start = time.time() 5 multiply_blocks(drv.Out(dest), drv.In(a), drv.In(b), numpy.int32(SIZE), block=(256, 1, ,→ 1), grid=((SIZE + 255) // 256, 1)) 6 end = time.time() 7 print(f"Tiempo GPU {end-start}") 8 9 print(numpy.array_equal(dest, check)) 50 Ejemplo: Multiplicación de vectores por bloques SIZE = 1500, BLOCK SIZE = 256 1 Tiempo CPU: 0.0006 segundos 2 Tiempo GPU 0.0018 segundos 3 Speedup cpu/gpu 0.33x 4 True 51 Ejemplo: Multiplicación de vectores por bloques SIZE = 15000000, BLOCK SIZE = 256 1 Tiempo CPU: 5.7086 segundos 2 Tiempo GPU 0.0509 segundos 3 speedup cpu/gpu 112.02x 4 True 52 Posibles causas de bajo rendimiento en GPU Ausencia de paralelismo: Para operaciones muy simples es posible que no haya suficiente paralelismo para justificar el uso de la GPU. Overhead de transferencia de datos: Coste asociado a la transferencia de datos entre la CPU y la GPU a través del bus PCI Express. Para operaciones pequeñas puede ser significativo en comparación con el tiempo de cómputo real en la GPU. Dimensionalidad: Para tamaños de problema pequeños, el tiempo necesario para lanzar y configurar el kernel en la GPU puede superar el tiempo de cómputo real, lo que resulta en un rendimiento más lento en comparación con la ejecución en la CPU. 53 Ejemplo: Multiplicación de vectores por bloques 1 [...] 2 from pycuda.compiler import SourceModule 3 mult = SourceModule(""" 4 __global__ void multiply_blocks(float *dest, float *a, float *b, int n) 5 { 6 int index = blockIdx.x * blockDim.x + threadIdx.x; 7 if (index < n) 8 dest[index] = a[index] * b[index]; 9 } 10 """) 11 multiply_blocks = mult.get_function("multiply_blocks") 12 [...] 13 multiply_blocks(drv.Out(dest), drv.In(a), drv.In(b), numpy.int32(SIZE), block=(2, 1, ,→ 1), grid=((SIZE + 1) // 2, 1)) 14 [...] 54 Ejemplo: Multiplicación de vectores por bloques La configuración de tamaños y dimensiones del grid impacta directamente en el rendimiento obtenido. El reparto de datos es todavı́a más importante cuando un hilo necesita acceder a información de sus vecinos cercanos. SIZE = 15000000, BLOCK SIZE = 2 1 Tiempo CPU: 5.4351 segundos 2 Tiempo GPU 0.0671 segundos 3 speedup cpu/gpu 80.94x 4 True 55 Matrices: Indexación 1D vs 2D La memoria se organiza como un array unidimensional (row-major vs column-major). El compilador traducirá cualquier indexación 2D o 3D a la representación 1D equivalente. Utilizar directamente indexación 1D puede ayudar a controlar el patrón de accesos empleado (accesos coalescentes). Algunas librerı́as obligan a trabajar con indexación 1D (cuBLAS). Figura 5: Indexación de matriz bidimensional. 56 Repaso Repaso Denominamos computación heterogénea a aquella que combina el uso de distintos tipos de procesadores. Los aceleradores Hardware son co-procesadores muy eficientes para una tarea determinada que se utilizan en conjunto con una CPU de proposito general. La evolución de los procesadores ha dado lugar, por limitaciones fı́sicas, a que para seguir aumentando prestaciones deba usarse un modelo basado en el aumento del ancho de banda. 57 Repaso Las GPUs permiten, gracias a su particular arquitectura, realizar de manera muy eficiente cálculos con alta carga computacional en paralelo. CUDA es una plataforma de computación que permite desarrollar código para ser ejecutado en GPU (Nvidia) con un código muy similar a C. Es posible acceder a sus capacidades desde Python mediante el uso de wrappers. Un conocimiento detallado de la arquitectura sobre la que se proyectará el código resulta crı́tico para maximizar el rendimiento de los programas desarrollados. 58