Computación de altas prestaciones con GPUs

1. Presente y futuro de los sistemas de computación Cursos de verano 2010 Computación de altas prestaciones con GPUs Enrique Arias Antúnez José Luis Sánchez García

2. Índice • Introducción • Arquitectura CUDA • Modelo de ejecución • Programación • Rendimiento

3. Computación de altas prestaciones • Los grandes supercomputadores son la única alternativa para cierto tipo de aplicaciones • Sin embargo, hay aplicaciones con menos exigencias computacionales • GPGPU • Utilizar GPUs para aplicaciones de propósito general • Muy atractivo en términos de rendimiento, consumo y coste • Exclusividad

4. Introducción Rendimiento Fuente: Nvidia

5. Introducción Precios Más de 100 millones de GPUs de Nvidia

6. Introducción Uso de los transistores + Cores - Cache - Control + Cache + Control - Cores

7. Introducción

8. Introducción • GeForce 8800 GTX vs. 2.2GHz Opteron 248 • 10 speedup en un kernel es típico, si hay suficiente paralelismo • 25 a 400 speedup con optimizaciones • Depende mucho del tipo de aplicación

9. Introducción Programabilidad Explosión GPGPU Facilidad de programación OpenCL CUDA Inicios GPGPU Brook++ RapidMind ATI Stream Cg Futuro GPGPU? OpenGL DirectX Ensamblador 2005 2009

10. Índice • Introducción • Arquitectura CUDA • Modelo de ejecución • Programación • Rendimiento

11. SP DP SM Arquitectura CUDA Conjunto de Streaming Multiprocessors (MP) . . . Memoria global 8 Scalar Processors (SP) 1 Unidad de Doble Precisión (DP) 16 KB de Memoria Compartida (SM) 8-16 K Registros

12. Fermi ~1.5TFLOPS (SP) ~800GFLOPS (DP) 230 GB/s DRAM

13. Espacios de memoria en CUDA Fuente: Nvidia

14. Espacios de memoria en CUDA

15. Gestión de memoria con CUDA • CPU y GPU tienen espacios de memoria independientes • Transferencia de datos entre ambos espacios de memoria a través del bus PCIExpress • Reserva, transferencia y liberación explícitas • Las operaciones con memoria realizadas por el host Fuente: Nvidia

16. Gestión de memoria con CUDA • cudaMalloc() • Obtiene espacio en la memoria global • Parámetros: dirección del puntero y el tamaño a reservar • cudaMemset() • Inicializa a un valor dado • Parámetros: dirección, valor y cantidad • cudaFree() • Libera el espacio • Parámetros: dirección Fuente: Nvidia

17. Gestión de memoria con CUDA • cudaMemcpy() • Transfiere datos entre memorias • Requiere 4 parámetros: • Puntero al destino • Puntero al origen • Bytes a copiar • Tipo de transferencia: • Host a host • Host a dispositivo • Dispositivo a host • Dispositivo a dispositivo Fuente: Nvidia

18. Gestión de memoria con CUDA • Ejemplo • Reservar, inicializar y liberar espacio para una matriz 64x64 de elementos float • Enlazar ese espacio a Md #define BLOCK_SIZE 64 float * Md; int size = BLOCK_SIZE*BLOCK_SIZE*sizeof(float); cudaMalloc((void**)&Md,size); cudaMemset(Md,0,size); . . . cudaFree(Md);

19. Gestión de memoria con CUDA • Ejemplo • Transferir una matriz de 64x64 float de la memoria de la GPU a la del host y viceversa • Matriz M está en el host y la matriz Md en el dispositivo cudaMemcpy(M,Md,size,cudaMemcpyDeviceToHost); cudaMemcpy(Md,M,size,cudaMemcpyHostToDevice);

20. Índice • Introducción • Arquitectura CUDA • Modelo de ejecución • Programación • Rendimiento

21. Modelo de ejecución Secuencial Thread Código

22. Modelo de ejecución Paralelo Threads Código

23. Paralelo Miles de threads lanzados a la vez kernel

24. Modelo de ejecución CPU GPU Kernel Threads (instancias del kernel) PCIe Memoria

25. Modelo de ejecución • Cada thread tiene un identificador • Todos los threads ejecutan el mismo código • Cada thread opera sobre distintos datos • Modelo SIMD CPU • Threads pesados • Sobrecarga planificación • Cambios de contexto lentos CPU • Threads ligeros • Poca sobrecarga planificación • Cambios de contexto rápidos CPU GPU

26. Grid 1 Bloque (0,0) Bloque (1,0) Bloque (2,0) Bloque (0,1) Bloque (1,1) Bloque (2,1) Thread (2,1) Thread (3,1) Thread (4,0) Thread (2,0) Thread (4,2) Thread (2,2) Thread (3,2) Thread (3,0) Thread (4,1) Bloque (1,1) Thread (0,0) Thread (1,0) Thread (0,1) Thread (0,2) Thread (1,2) Thread (1,1) Estructura jerárquica de threads • Los threads se agrupan en bloques de threads • Los bloques de threads se agrupan en Grids • Los threads de un bloque se comunican mediante la memoria compartida • Todos los threads del Grid se comunican a través de la memoria global • Un thread se ejecuta en un procesador escalar (SP) • Un bloque de threads se lanza en un multiprocesador (MP) • Un MP puede desarrollar varios bloques

27. Software Cada thread se ejecuta en un SP Thread SP • Un bloque se ejecuta en un MP • Los bloques no migran entre MPs • Varios bloques a la vez en un MP (según registros y memoria compartida) Bloque threads MP … • Un único kernel concurrente Grid GPU Modelo de ejecución Hardware

28. blockIdx.x=1 blockIdx.y=1 gridDim.x gridDim.y Grid 1 Bloque (0,0) Bloque (1,0) Bloque (2,0) Bloque (0,1) Bloque (1,1) Bloque (2,1) Thread (2,1) Thread (2,0) Thread (4,0) Thread (3,1) Thread (3,0) Thread (2,2) Thread (3,2) Thread (4,2) Thread (4,1) Bloque (1,1) Thread (0,0) Thread (1,0) blockDim.y Thread (0,2) Thread (0,1) Thread (1,2) Thread (1,1) threadIdx.x=0 threadIdx.y=2 blockDim.x blockDim.z Indentificadores y dimensiones • El tamaño del grid y de los bloques los determina el programador • Se usan las variables gridDim y blockDim para referenciar la dimensión de grid y bloque, respectivamente • Un thread queda indentificado por: • Un identificador propio dentro del bloque al que pertenece • El identificador del bloque al que pertenece • Se usan las variables threadIdx y blockIdx para referenciar el identificador del thread dentro del bloque y al bloque dentro del grid, respectivamente

29. Bloque 1 Bloque 2 Bloque n m m m 2 2 2 warp 1 warp 1 warp 1 warp 8 instrucción 11 warp 1 instrucción 42 warp 3 instrucción 95 tiempo ... warp 8 instrucción 12 warp 3 instrucción 96 Planificación • Se agrupan los threads en bloques • Se asignan identificadores a bloques y threads • Se distribuyen los bloques de threads entre los multiprocesadores • Los threads de un bloque se ejecutan concurrentemente en un multiprocesador • Los threads de un bloque son agrupados en warps • Un warp es la unidad mínima de planificación y está formada por 32 threads • Varios warps en cada multiprocesador, pero sólo uno está en ejecución en cada momento • Los warps cuyos operandos están listos son seleccionados para ejecución • Todos los threads en un warp ejecutan la misma instrucción

30. warp 8 instrucción 11 warp 1 instrucción42 warp 3 instrucción95 warp 8 instrucción12 warp 3 instrucción96 Planificación • Tres flujos de instrucciones: warp1, warp3 y warp8 t=k t=k+1 t=k+2 Planifica en tiempo k ... t=l>k t=l+1

31. warp 8 instrucción 11 warp 1 instrucción42 warp 3 instrucción95 warp 8 instrucción12 warp 3 instrucción96 Planificación • Tres flujos de instrucciones: warp1, warp3 y warp8 Planifica en tiempo k+1 t=k t=k+1 t=k+2 ... t=l>k t=l+1

32. Índice • Introducción • Arquitectura CUDA • Modelo de ejecución • Programación • Rendimiento

33. Programación • Computación heterogénea • Parte de código se ejecuta en la CPU • Parte de código se ejecuta en la GPU (kernel) • La API de CUDA es una extensión al lenguaje ANSI C • Curva de aprendizaje suave • Extensiones básicas • Modificadores de función • Modificadores de variables • Variables específicas de dimensionado • Directiva para ejecución del kernel

34. Indica dónde se ejecuta la función: GPU (device) o CPU (host) __device__La función debe ejecutarse en el dispositivo Sólo puede ser llamada por el propio dispositivo Recursividad no soportada No pueden declararse variables estáticas dentro de la función La función no puede tener un número variable de argumentos __global__La función es un kernel que debe ejecutarse en el dispositivo Sólo puede ser llamada por el host Recursividad no soportada No pueden declararse variables estáticas dentro de la función La función no puede tener un número variable de argumentos La función debe devolver siempre void __host__La función debe ejecutarse en el host Sólo puede ser llamada por el host No puede utilizarse junto con __global__ Modificadores de función

35. Indica en qué parte de la memoria se localiza la variable __device__La variable reside en el dispositivo Requiere que se indique uno de los otros dos modificadores de variables para indicar dónde exactamente reside la variable en el dispositivo __constant__La variable reside en el espacio de memoria constante del dispositivo Está viva durante todo el tiempo de ejecución de la aplicación Accesible por todos los threads del grid, así como desde el host __shared__La variable reside en el espacio de memoria compartida del bloque de threads en el dispositivo Está viva mientras el bloque está vivo Accesible sólo por los threads del bloque Modificadores de variables

36. Indican las dimensiones que caracterizan a los identificadores de los threads y bloques dim3 gridDim; Dimensiones de los grids en bloques (gridDim.z no usado) dim3 blockDim; Dimensiones del bloque en threads dim3 blockIdx; Indice del bloque en el grid dim3 threadIdx; Indice del thread en el bloque Variables específicas dimensiones

37. Indica cómo debe ejecutarse el kernel en el dispositivo Cada vez que se llama a __global__, también debe especificarse la configuración de ejecución del kernel Se inserta entre el nombre de la función y el paréntesis del argumento una expresión de la forma: <<< Dg, Db, Ns, S >>> Dg: indica las dimensiones de la malla, es decir, la cantidad de bloques Db: indica las dimensiones de cada bloque, es decir, el número de threads por bloque Ns: indica el número de bytes de memoria compartida asignada dinámicamente por bloque (argumento opcional, 0 por defecto) S: especifica un stream (argumento opcional, 0 por defecto) Las invocaciones de kernels son asíncronas El control vuelve al programa principal (CPU) Directiva ejecución del kernel

38. Ejecución del kernel • Un kernel debe ser ejecutado del siguiente modo: __global__ void KernelFunction(…); dim3 DimGrid(100,50); // 5000 bloques dim3 DimBlock(4,8,8); // 256 threads/bloque KernelFunction<<< Dimgrid,Dimblock >>>(…);

39. Ejemplo: SAXPY // Definición de la función void saxpy_serial (int n, float a, float *x, float *y) { for (int i = 0; i < n; ++i) y[i] = a*x[i] + y[i]; } // Llamada a la función saxpy(n,2.0,x,y) // Definición del kernel __global__ void saxpy_parallel (int n, float a, float *x, float *y) { int i = blockIdx.x*blockDim.x + threadIdx.x; if (i < n) y[i] = a*x[i] + y[i]; } // Llamada al kernel int nblocks = (n + 255) / 256; saxpy_parallel<<<nblocks, 256>>>(n, 2.0, x, y);

40. Índice • Introducción • Arquitectura CUDA • Modelo de ejecución • Programación • Rendimiento

41. Optimizaciones • Reducir sobrecargas en el acceso a memoria • Realizar transferencias asíncronas • Mejorar accesos a memoria global • Usar memoria compartida • Reducir conflictos de bancos en memoria compartida • Configurar la ejecución • Seleccionar el número de bloques y threads por bloque • Optimizar a nivel de instrucción • Instrucciones aritméticas • Instrucciones para acceso a memoria • Saltos y sentencias condicionales

42. Transferencias asíncronas • cudaMemcpy es bloqueante • Se devuelve el control al host una vez finalizada la transferencia • cudaMemcpyAsync es no bloqueante • Se devuelve el control inmediatamente • Las transferencias no bloqueantes permiten solapar computación y comunicación

43. Solapar computación y comunicación Datos1 Datos2 Datos1 Datos2 Datos2 Datos1 • Un stream es un conjunto de operaciones que se completan secuencialmente • Se puede solapar la comunicación y computación de diferentes streams • Para ello se debe permitir que la memoria del host esté mapeada en el espacio de direcciones del dispositivo Host→Device Kernel Cálculos Device→Host Datos1 Datos2 Datos1 Datos2 Datos1 Datos2 Host→Device Kernel Cálculos 43 Device → Host

44. Memoria global • Latencia alta • Usarla lo menos posible • Accesos por half-warp (16 threads) • Intentar completarlos en el menor número de transacciones posible (coalescing) 16 transacciones 1 transacción

45. Memoria global • Latencia alta • Usarla lo menos posible • Accesos por half-warp (16 threads) • Intentar completarlos en el menor número de transacciones posible (coalescing) memoria global memoria global matriz matriz 1 1 1 2 3 4 1 2 3 4 2 2 3 3 5 6 7 8 5 6 7 8 4 4 9 10 11 12 9 10 11 12 5 5 6 6 13 14 15 16 13 14 15 16 7 7 8 8 … …

46. Sin conflictos Sin conflictos Con conflictos Banco 0 Banco 0 Banco 0 Thread 0 Thread 0 Thread 0 Banco 1 Banco 1 Banco 1 Thread 1 Thread 1 Thread 1 Banco 2 Banco 2 Banco 2 Thread 2 Thread 2 Thread 2 Banco 3 Banco 3 Banco 3 Thread 3 Thread 3 Thread 3 . . . Banco 4 Banco 4 Banco 4 Thread 4 Thread 4 Thread 8 . . . . . . . . . . . . Banco 5 . . . Thread 9 Banco 14 Banco 14 Thread 14 Thread 14 . . . Banco 15 Banco 15 Banco 15 Thread 15 Thread 15 Thread 15 Direccionamiento aleatorio Permutación 1:1 Direccionamiento lineal Stride = 2 Direccionamiento lineal Stride = 1 Memoria compartida • Dividida en módulos (bancos) • En las GPUs actuales hay 16 bancos de 1KB • Acceso simultáneo a los bancos • Baja latencia (similar registros) • Problema: conflictos en los bancos (acceso al mismo banco por dos o más threads)

47. Memoria compartida • Dividida en módulos (bancos) • En las GPUs actuales hay 16 bancos de 1KB • Acceso simultáneo a los bancos • Baja latencia (similar registros) • Problema: conflictos en los bancos (acceso al mismo banco por dos o más threads) • Cargar los datos desde la memoria global a la memoria compartida • Sincronizar threads • Procesar usando sólo memoria compartida • Sincronizar • Llevar resultados a memoria global

48. Bloques y threads • Ocupación: • Número de warps activos por MP con respecto al máximo posible de warps activos • El número total de bloques y de threads por bloque son factores importantes • Número de bloques • Dependiente de los recursos disponibles • Al menos tantos como MPs • Más de uno por MP para que unos mantengan los recursos ocupados mientras otros se sincronizan • Suficientemente alto para escalar en futuras versiones • Número de threads por bloque • Múltiplo del tamaño del warp para facilitar coalescing • Múltiplos de 64, para evitar conflictos en acceso a registros • Entre 128 y 256 buena elección • Más threads no implica necesariamente mayor ocupación • Experimentación

49. Rendimiento a nivel de instrucción • La productividad a nivel de instrucción depende de • Número de operaciones • Latencia de la memoria • Ancho de banda de la memoria • Maximizar ancho de banda efectivo • Maximizar el uso de la memoria compartida • Minimizar los accesos a memoria global • Maximizar coalescing en los accesos a memoria global • Solapar comunicación y computación • Aumentar el rendimiento de la ejecución de las instrucciones • Instrucciones aritméticas más rápidas, si se prefiere velocidad en lugar de precisión • Instrucciones para accesos a memoria con menos latencia • Evitar sentencias condicionales que generen diferentes caminos en el mismo warp, pues éstos son serializados

50. Branch Path A Path B Divergencia de caminos • Los caminos son serializados • Incremento del número de instrucciones ejecutadas por el warp • Los threads vuelven a converger cuando todos los caminos se completan • 50% de pérdida de rendimiento Salto Path A Path B