Presente y futuro de los sistemas de computacin

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 1

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

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 3

Introducción Rendimiento Fuente: Nvidia 4

Introducción Precios 2 x Xeon Quad. Core + Tesla s 1070 Cores 8 8 + 240*4 = 968 Rendimiento pico (sp) 0, 17 TFlops 4, 14 TFlops Rendimiento pico (dp) 0, 08 TFlops 0, 35 TFlops Precio 2400 euros 9000 euros Consumo 700 W 1500 W Más de 100 millones de GPUs de Nvidia 5

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

Introducción Aplicación Descripción H. 264 SPEC ‘ 06 version, change in guess vector LBM Código Kernel % tiempo 34. 811 194 35% SPEC ‘ 06 version, change to single precision and print fewer reports 1. 481 285 >99% RC 5 -72 Distributed. net RC 5 -72 challenge client code 1. 979 218 >99% FEM Finite element modeling, simulation of 3 D graded materials 1. 874 146 99% RPES Rye Polynomial Equation Solver, quantum chem, 2 -electron repulsion 1. 104 281 99% PNS Petri Net simulation of a distributed system 322 160 >99% SAXPY Single-precision implementation of saxpy, used in Linpack’s Gaussian elim. routine 952 31 >99% TRACF Two Point Angular Correlation Function 536 98 96% FDTD Finite-Difference Time Domain analysis of 2 D electromagnetic wave propagation 1. 365 93 16% MRI-Q Computing a matrix Q, a scanner’s configuration in MRI reconstruction 490 33 >99% 7

Introducción • • Ge. Force 8800 GTX vs. 2. 2 GHz 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 8

Introducción Facilidad de programación Programabilidad Explosión GPGPU Inicios GPGPU Brook++ Rapid. Mind Cg CUDA Open. CL ATI Stream Futuro GPGPU? Open. GL Direct. X Ensamblador 2005 2009 9

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

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

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

Espacios de memoria en CUDA 13 Fuente: Nvidia

Espacios de memoria en CUDA Memoria Localización Cache Acceso Ámbito Vida Registros On-chip N/A R/W Un thread Thread Compartida On-chip N/A R/W Threads en bloque Bloque Global Off-chip No R/W Todos threads y host Aplicación Constantes Off-chip Sí R Todos threads y host Aplicación Texturas Off-chip Sí R Todos threads y host Aplicación 14

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 15

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

Gestión de memoria con CUDA • cuda. Memcpy() – 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 17

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

Gestión de memoria con CUDA • Ejemplo – Transferir una matriz de 64 x 64 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 cuda. Memcpy(M, Md, size, cuda. Memcpy. Device. To. Host); cuda. Memcpy(Md, M, size, cuda. Memcpy. Host. To. Device); 19

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

Modelo de ejecución Secuencial Thread Código 21

Modelo de ejecución Paralelo Threads Código 22

Paralelo kernel Miles de threads lanzados a la vez 23

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

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 GPU • • • Threads pesados Sobrecarga planificación Cambios de contexto lentos CPU • • • Threads ligeros Poca sobrecarga planificación Cambios de contexto rápidos 25

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 Grid 1 Bloque (0, 0) Bloque (1, 0) Bloque (2, 0) Bloque (0, 1) Bloque (1, 1) Bloque (2, 1) Bloque (1, 1) Thread Thread (0, 0) (1, 0) (2, 0) (3, 0) (4, 0) Thread Thread (0, 1) (1, 1) (2, 1) (3, 1) (4, 1) Thread Thread (0, 2) (1, 2) (2, 2) (3, 2) (4, 2) 26

Modelo de ejecución Software Thread Bloque threads Hardware Cada thread se ejecuta en un 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) • Un único kernel concurrente SP MP … Grid • GPU 27

Indentificadores y dimensiones • • block. Idx. x=1 block. Idx. y=1 El tamaño del grid y de los bloques los determina el programador Se usan las variables grid. Dim y block. Dim para referenciar la dimensión de grid y bloque, respectivamente grid. Dim. x grid. Dim. y • 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 thread. Idx y block. Idx para referenciar el identificador del thread dentro del bloque y al bloque dentro del grid, respectivamente Bloque (1, 0) Bloque (2, 0) Bloque (0, 1) Bloque (1, 1) Bloque (2, 1) Thread Thread (0, 0) (1, 0) (2, 0) (3, 0) (4, 0) thread. Idx. x=0 thread. Idx. y=2 z im. k. D loc b Bloque (0, 0) Bloque (1, 1) block. Dim. y • Grid 1 Thread Thread (0, 1) (1, 1) (2, 1) (3, 1) (4, 1) Thread Thread (0, 2) (1, 2) (2, 2) (3, 2) (4, 2) block. Dim. x 28

Planificación Bloque 1 • • • 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 2 warp 1 m 2 warp 1 Bloque n m 2 warp 1 warp 8 instrucción 11 warp 1 instrucción 42 tiempo • • • Bloque 2 warp 3 instrucción 95. . . warp 8 instrucción 12 warp 3 instrucción 96 29 m

Planificación • Tres flujos de instrucciones: warp 1, warp 3 y warp 8 Warp Instrucción actual Estado de la Instrucción Warp 1 42 Computando warp 8 instrucción 11 t=k warp 1 instrucción 42 t=k+1 Warp 3 95 Computando warp 3 instrucción 95. . . warp 8 instrucción 12 t=k+2 Warp 8 11 Operandos preparados t=l>k … warp 3 instrucción 96 t=l+1 Planifica en tiempo k 30

Planificación • Tres flujos de instrucciones: warp 1, warp 3 y warp 8 Warp Current Instruction State Warp 1 42 Preparado escribir result. warp 8 instrucción 11 t=k warp 1 instrucción 42 t=k+1 Warp 3 95 Computando warp 3 instrucción 95. . . warp 8 instrucción 12 t=k+2 Warp 8 11 Computando t=l>k … warp 3 instrucción 96 t=l+1 Planifica en tiempo k+1 31

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

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 33

Modificadores de función 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__ 34

Modificadores de variables 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 35

Variables específicas dimensiones Indican las dimensiones que caracterizan a los identificadores de los threads y bloques • • dim 3 grid. Dim; – Dimensiones de los grids en bloques (grid. Dim. z no usado) dim 3 block. Dim; – Dimensiones del bloque en threads dim 3 block. Idx; – Indice del bloque en el grid dim 3 thread. Idx; – Indice del thread en el bloque 36

Directiva ejecución del kernel 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) 37

Ejecución del kernel • Un kernel debe ser ejecutado del siguiente modo: __global__ void Kernel. Function(…); dim 3 Dim. Grid(100, 50); // 5000 bloques dim 3 Dim. Block(4, 8, 8); // 256 threads/bloque Kernel. Function<<< Dimgrid, Dimblock >>>(…); 38

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 = block. Idx. x*block. Dim. x + thread. Idx. 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); 39

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

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 41

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

Solapar computación y comunicación Datos 1 Datos 2 • • • 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 Datos 1 Datos 2 Kernel Cálculos Datos 1 Datos 2 Device→Host Datos 1 Datos 2 Host→Device Kernel Cálculos 43 Device → Host 43

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) 1 transacción 16 transacciones 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) memoria global matriz 1 2 3 4 5 6 7 9 10 13 14 1 memoria global matriz 1 2 3 4 8 3 5 6 7 8 3 11 12 5 9 10 11 12 5 15 16 6 13 14 15 16 6 4 7 2 4 7 8 8 … … 45

Memoria compartida • • • Dividida en módulos (bancos) En las GPUs actuales hay 16 bancos de 1 KB 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) Sin conflictos Con conflictos Thread 0 Banco 0 Thread 1 Banco 1 Thread 2 Banco 2 Thread 3 Banco 3 Thread 4 Banco 4 . . . Thread 8 Banco 5 Thread 14 Banco 14 Thread 9 . . . Thread 15 Banco 15 Direccionamiento lineal Stride = 1 Direccionamiento aleatorio Permutación 1: 1 . . . 46 Direccionamiento lineal Stride = 2

Memoria compartida • • • Dividida en módulos (bancos) En las GPUs actuales hay 16 bancos de 1 KB 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 47

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 48

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 49

Divergencia de caminos Salto Branch Path A Path B • 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 50

Índice • Producto Matriz-Matriz • ¿Cómo explotar más la arquitectura? • Buenas prácticas de programación CUDA • Open. CL 51

Índice • Producto Matriz-Matriz – – – Implementación secuencial Recordar Implementación básica ¿Cómo compilar? Implementación memoria compartida • ¿Cómo explotar más la arquitectura? • Buenas prácticas de programación CUDA • Open. CL 52

Producto Matriz-Matriz • Implementación secuencial 53

Producto Matriz-Matriz • • • cuda. Malloc() cuda. Memcpy() cuda. Thread. Synchronize(); cuda. Free() multiplication<<< grid, threads >>>(d_P, d_M, d_N, M. width, N. width); • __global__ void multiplication(float *P, float *M, float *N, int w. M, int w. N) 54

Producto Matriz-Matriz • Implementación básica – Cada thread calcula un elemento de la matriz • Lee una fila de A • Lee una columna de B 55

Producto Matriz-Matriz • Implementación básica – Ejercicio 1: (En programa principal) • a) Repasar las llamadas y el significado • b) Completar llamada a kernel • c) Cálculo de Giga. Flops – Ejercicio 2: (En kernel) • a) Cálculo de índices • b) Cálculo de elemento – Ejercicio 3: (Pruebas) 56

Producto Matriz-Matriz • Implementación básica – Ejercicio 2. a (Ejemplo) (0, 0) (1, 0) (0, 1) (1, 1) (0, 2) (1, 2) 3 5 (0, 0) (1, 0) (2, 0) (0, 1) (1, 1) (2, 1) (0, 2) (1, 2) (2, 2) (0, 0) (1, 0) (2, 0) (0, 1) (1, 1) (2, 1) (0, 2) (1, 2) (2, 2) 57

Producto Matriz-Matriz • ¿Cómo compilar? gmultiply: gmultiply. cu nvcc gmultiply. cu -keep --ptxas-options=-v -o gmultiply -lcutil -L /usr/local/cuda/sdk/lib -I /usr/local/cuda/sdk/common/inc 58

GPU Ge. Force 285 GTX • 30 multiprocesadores (MPs) – 8 cores (FPUs) por multiprocesador, a 1476 MHz – mad: multiplicación-suma (considera 2 flops a efectos de cálculo de productividad) – 2 SFU (unidades funcionales especiales) por multiprocesador • Total 240 cores • 797 GFLOPS (rendimiento pico) – (30 MPs) x (8 x 2 + 2 Flops/MP) x 1476 MHz = 797 GFLOPS • 1 GB RAM – 1242 MHz (double data rate, 2484 MHz) – 512 bits (ancho bus) – 159 GB/s = (2484 MHz) x (512 / 8 bytes) 59

Producto Matriz-Matriz • Se usan los cores (no las SFUs) – (240 cores) x 1476 MHz = = 354, 24 Goperaciones/s • • De 8 instrucciones sólo 2 son FLOP Máxima productividad – 354, 24 x ¼ = 88, 56 GLOPS • ¿Por qué no se llega a ese valor? • Accesos a memoria $Lt_0_7: ld. global. f 32 mad. f 32 add. s 32 setp. ne. u 32 @%p 1 bra %f 2, [%r 15+0]; %f 3, [%r 17+0]; %f 1, %f 2, %f 3, %f 1; %r 17, 4096; %r 15, 4; %p 1, %r 15, %r 16; $Lt_0_7; – ¼ operaciones son cargas – (240 cores) x (¼ cargas) x (4 bytes/carga) x (1476 MHz) = 354 GBs > 159 GBs El ancho de banda con memoria global no es suficiente 60

Producto Matriz-Matriz • Cada dato de entrada es leído por width threads • • Latencia memoria global ~ 400 ciclos Latencia memoria compartida ~ 10 ciclos • Llevar cada dato de entrada a memoria compartida y que sea leído por varios threads 61

Producto Matriz-Matriz • Cada dato de entrada es leído por width threads • • Latencia memoria global ~ 400 ciclos Latencia memoria compartida ~ 10 ciclos • Llevar cada dato de entrada a memoria compartida y que sea leído por varios threads Producto de sub-matrices • • Cada sub-matriz es calculada por un bloque de threads Los datos de entrada necesarios son llevados a memoria compartida 62

Producto Matriz-Matriz • Implementación memoria compartida – Disminuir accesos a memoria – ¿Cómo lo haríais? 63

Producto Matriz-Matriz B 0, 0 B 1, 0 B 2, 0 B 3, 0 B 0, 1 B 1, 1 B 2, 1 B 3, 1 B 0, 2 B 1, 2 B 2, 2 B 3, 2 B 0, 3 B 1, 3 B 2, 3 B 3, 3 A 0, 0 A 1, 0 A 2, 0 A 3, 0 C 0, 0 C 1, 0 C 2, 0 C 3, 0 A 0, 1 A 1, 1 A 2, 1 A 3, 1 C 0, 1 C 1, 1 C 2, 1 C 3, 1 A 0, 2 A 1, 2 A 2, 2 A 3, 2 C 0, 2 C 1, 2 C 2, 2 C 3, 2 A 0, 3 A 1, 3 A 2, 3 A 3, 3 C 0, 3 C 1, 3 C 2, 3 C 3, 3 64

Producto Matriz-Matriz orden accesos Cada thread calcula un punto de la matriz resultado C 0, 0 thread 0, 0 C 1, 0 thread 1, 0 C 0, 1 thread 0, 1 C 1, 1 thread 1, 1 A 0, 0 x B 0, 0 A 0, 0 x B 1, 0 A 0, 1 x B 0, 0 A 0, 1 x B 1, 0 A 1, 0 x B 0, 1 A 1, 0 x B 1, 1 A 1, 1 x B 0, 1 A 1, 1 x B 1, 1 A 2, 0 x B 0, 2 A 2, 0 x B 1, 2 A 2, 1 x B 0, 2 A 2, 1 x B 1, 2 A 3, 0 x B 0, 3 A 3, 0 x B 1, 3 A 3, 1 x B 0, 3 A 3, 1 x B 1, 3 En una primera fase se llevan a memoria compartida En una segunda fase se llevan a memoria compartida Cada Ai, j y Bi, j se lee dos veces 65

Producto Matriz-Matriz threads (0, 0) (1, 0) calculan C 0, 0 C 1, 0 B 0, 0 B 1, 0 B 2, 0 B 3, 0 C 0, 1 C 1, 1 (0, 1) (1, 1) B 0, 1 B 1, 1 B 2, 1 B 3, 1 B 0, 2 B 1, 2 B 2, 2 B 3, 2 cada thread lleva un dato de A y B As Bs memoria compartida B 0, 3 B 1, 3 B 2, 3 B 3, 3 A 0, 0 A 1, 0 A 2, 0 A 3, 0 C 0, 0 C 1, 0 C 2, 0 C 3, 0 A 0, 1 A 1, 1 A 2, 1 A 3, 1 C 0, 1 C 1, 1 C 2, 1 C 3, 1 A 0, 2 A 1, 2 A 2, 2 A 3, 2 C 0, 2 C 1, 2 C 2, 2 C 3, 2 A 0, 3 A 1, 3 A 2, 3 A 3, 3 C 0, 3 C 1, 3 C 2, 3 C 3, 3 66

Producto Matriz-Matriz threads (0, 0) (1, 0) calculan C 0, 0 C 1, 0 B 0, 0 B 1, 0 B 2, 0 B 3, 0 C 0, 1 C 1, 1 (0, 1) (1, 1) B 0, 1 B 1, 1 B 2, 1 B 3, 1 C 0, 0 = A 0, 0 x B 0, 0 + A 1, 0 x B 0, 1 B 0, 2 B 1, 2 B 2, 2 B 3, 2 C 1, 0 = A 0, 0 x B 1, 0 + A 1, 0 x B 1, 1 B 0, 3 B 1, 3 B 2, 3 B 3, 3 C 0, 1 = A 0, 1 x B 0, 0 + A 1, 1 x B 0, 1 C 1, 1 = A 0, 1 x B 1, 0 + A 1, 1 x B 1, 1 A 0, 0 A 1, 0 A 2, 0 A 3, 0 C 0, 0 C 1, 0 C 2, 0 C 3, 0 A 0, 1 A 1, 1 A 2, 1 A 3, 1 C 0, 1 C 1, 1 C 2, 1 C 3, 1 As Bs A 0, 0 A 1, 0 B 0, 0 B 1, 0 A 0, 2 A 1, 2 A 2, 2 A 3, 2 C 0, 2 C 1, 2 C 2, 2 C 3, 2 A 0, 1 A 1, 1 B 0, 1 B 1, 1 A 0, 3 A 1, 3 A 2, 3 A 3, 3 C 0, 3 C 1, 3 C 2, 3 C 3, 3 memoria compartida 67

Producto Matriz-Matriz threads (0, 0) (1, 0) calculan C 0, 0 C 1, 0 B 0, 0 B 1, 0 B 2, 0 B 3, 0 C 0, 1 C 1, 1 (0, 1) (1, 1) B 0, 1 B 1, 1 B 2, 1 B 3, 1 C 0, 0 = A 0, 0 x B 0, 0 + A 1, 0 x B 0, 1 + A 2, 0 x B 0, 2 + A 3, 0 x B 0, 3 C 1, 0 = A 0, 0 x B 1, 0 + A 1, 0 x B 1, 1 + A 2, 0 x B 1, 2 + A 3, 0 x B 1, 3 C 0, 1 = A 0, 1 x B 0, 0 + A 1, 1 x B 0, 1 + A 2, 1 x B 0, 2 + A 3, 1 x B 0, 3 B 0, 2 B 1, 2 B 2, 2 B 3, 2 B 0, 3 B 1, 3 B 2, 3 B 3, 3 C 1, 1 = A 0, 1 x B 1, 0 + A 1, 1 x B 1, 1 + A 2, 1 x B 1, 2 + A 3, 1 x B 1, 3 A 0, 0 A 1, 0 A 2, 0 A 3, 0 C 0, 0 C 1, 0 C 2, 0 C 3, 0 A 0, 1 A 1, 1 A 2, 1 A 3, 1 C 0, 1 C 1, 1 C 2, 1 C 3, 1 As Bs A 2, 0 A 3, 0 B 0, 2 B 1, 2 A 0, 2 A 1, 2 A 2, 2 A 3, 2 C 0, 2 C 1, 2 C 2, 2 C 3, 2 A 2, 1 A 3, 1 B 0, 3 B 1, 3 A 0, 3 A 1, 3 A 2, 3 A 3, 3 C 0, 3 C 1, 3 C 2, 3 C 3, 3 memoria compartida 68

Producto Matriz-Matriz • • • Bloque de threads 16 x 16 = 256 threads Width = 4096 256 x 256 = 65536 bloques de threads Cada dato en una sub-matriz es leído por 16 threads Los accesos a memoria global se reducen en un factor 16 al usar memoria compartida • Se requiere ahora 354 / 16 ~ 22 GBs Ahora la memoria no es motivo para no alcanzar la productividad deseada

Producto Matriz-Matriz • Implementación memoria compartida – Ejercicio 1: (En programa principal) • a) Repasar las llamadas y el significado • b) Completar llamada a kernel • c) Cálculo de Giga. Flops – Ejercicio 2: (En kernel) • a) Cálculo de índices • b) Cálculo de elemento – Ejercicio 3: (Pruebas) 70

Producto Matriz-Matriz • Implementación memoria compartida – Disminuir accesos a memoria – ¿Cómo lo haríais? – Veamos el código ahora e identifiquemos los diferentes elementos 71

Índice ü Producto Matriz-Matriz • ¿Cómo explotar más la arquitectura? • Buenas prácticas de programación CUDA • Open. CL 72

¿Cómo explotar más la arquitectura? • Sistemas Distribuidos mediante MPI • Sistemas compartidos mediante Open. MP 73

Índice ü Producto Matriz-Matriz ü ¿Cómo explotar más la arquitectura? • Buenas prácticas de programación CUDA • Open. CL 74

Guía de buenas prácticas de programación CUDA • Maximizar ejecución paralela • Optimizar el uso de la memoria de cara obtener un mayor ancho de banda en el acceso a memoria • Optimizar el uso de las instrucciones para conseguir una mayor productividad 75

Índice ü Producto Matriz-Matriz ü ¿Cómo explotar más la arquitectura? ü Buenas prácticas de programación CUDA • Open. CL – ¿Qué es Open. CL? – ¿Qué diferencia Open. CL de CUDA? – Ejemplo: Suma de vectores 76

¿Qué es Open. CL? • Open. CL: Open Computing Language • Propuesto por a • ¿Quién está involucrado?

¿Qué diferencia Open. CL de CUDA? (I) • Punteros – CUDA struct Node {Node* next} n=n->next – Opencl struct Node {unsigned int next; } next=buf. Base+n

¿Qué diferencia Open. CL de CUDA? (II) • Kernels – CUDA • Programa compilado en formato binario – Open. CL • Se compila en tiempo de ejecución – Palabras Clave y lenguaje utilizado para los kernel

Ejemplo: Suma de vectores • CUDA _global_ void Suma. Vec(const float *a, const float *b, float *c) // Índice al elemento del vector int indice=block. Idx. x*block. Dim. x+thread. Idx. x c[indice]=a[indice]+b[indice] } • Open. CL _kernel void Suma. Vec(_global const float *a, _global const float *b, _global float *c) // Índice al elemento del vector int indice=get_global_id(0) c[indice]=a[indice]+b[indice] }

Índice ü Producto Matriz-Matriz ü ¿Cómo explotar más la arquitectura? ü Buenas prácticas de programación CUDA ü Open. CL 81