Índice Introducción Arquitectura CUDA Modelo de ejecución Programación Rendimiento 1
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
2
Introducción Rendimiento Fuente: Nvidia 3
Introducción Precios Más de 100 millones de GPUs de Nvidia 4
Introducción Uso de los transistores + Cache + Control – Cores + Cores – Cache – Control 5
Introducción 6
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 7
Introducción Programabilidad 2005 2009 Ensamblador OpenGL DirectX Cg Brook++ RapidMind CUDA OpenCL Facilidad de programación Inicios GPGPU ATI Stream Explosión GPGPU Futuro GPGPU? 8
Índice Introducción Arquitectura CUDA Modelo de ejecución Programación Rendimiento 9
Arquitectura CUDA . . . Memoria global Conjunto de Streaming Multiprocessors (MP) (Gp:) SP (Gp:) DP (Gp:) SM
8 Scalar Processors (SP) 1 Unidad de Doble Precisión (DP) 16 KB de Memoria Compartida (SM) 8-16 K Registros 10
Fermi ~1.5TFLOPS (SP) ~800GFLOPS (DP) 230 GB/s DRAM 11
Espacios de memoria en CUDA 12 Fuente: Nvidia
Espacios de memoria en CUDA 13
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 14 Fuente: Nvidia
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 15 Fuente: Nvidia
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 16 Fuente: Nvidia
Ejemplo Reservar, inicializar y liberar espacio para una matriz 64×64 de elementos float Enlazar ese espacio a Md Gestión de memoria con CUDA #define BLOCK_SIZE 64 float * Md; int size = BLOCK_SIZE*BLOCK_SIZE*sizeof(float); cudaMalloc((void**)&Md,size); cudaMemset(Md,0,size); . . . cudaFree(Md); 17
Ejemplo Transferir una matriz de 64×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 Gestión de memoria con CUDA cudaMemcpy(M,Md,size,cudaMemcpyDeviceToHost);
cudaMemcpy(Md,M,size,cudaMemcpyHostToDevice); 18
Índice Introducción Arquitectura CUDA Modelo de ejecución Programación Rendimiento 19
Modelo de ejecución Secuencial Código Thread 20
Modelo de ejecución Paralelo Código Threads 21
kernel Paralelo Miles de threads lanzados a la vez 22
Modelo de ejecución Memoria Kernel Threads (instancias del kernel) PCIe CPU GPU 23
Modelo de ejecución CPU GPU 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 24
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 (Gp:) Grid 1 (Gp:) Bloque (2,1) (Gp:) Bloque (2,0) (Gp:) Bloque (0,0) (Gp:) Bloque (1,0) (Gp:) Bloque (0,1) (Gp:) Bloque (1,1)
(Gp:) Thread (0,0) (Gp:) Thread (1,0) (Gp:) Thread (2,0) (Gp:) Thread (3,0) (Gp:) Thread (4,0) (Gp:) Thread (0,1) (Gp:) Thread (1,1) (Gp:) Thread (2,1) (Gp:) Thread (3,1) (Gp:) Thread (4,1) (Gp:) Thread (0,2) (Gp:) Thread (1,2) (Gp:) Thread (2,2) (Gp:) Thread (3,2) (Gp:) Thread (4,2) (Gp:) Bloque (1,1)
25
Modelo de ejecución Software Hardware (Gp:) Thread (Gp:) SP
(Gp:) 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)
(Gp:) Bloque threads (Gp:) MP
(Gp:) … (Gp:) Grid (Gp:) GPU (Gp:) Un único kernel concurrente
Cada thread se ejecuta en un SP 26
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 (Gp:) Grid 1 (Gp:) Bloque (2,1) (Gp:) Bloque (2,0) (Gp:) Bloque (0,0) (Gp:) Bloque (1,0) (Gp:) Bloque (0,1) (Gp:) Bloque (1,1)
(Gp:) Thread (0,0) (Gp:) Thread (1,0) (Gp:) Thread (2,0) (Gp:) Thread (3,0) (Gp:) Thread (4,0) (Gp:) Thread (0,1) (Gp:) Thread (1,1) (Gp:) Thread (2,1) (Gp:) Thread (3,1) (Gp:) Thread (4,1) (Gp:) Thread (0,2) (Gp:) Thread (1,2) (Gp:) Thread (2,2) (Gp:) Thread (3,2) (Gp:) Thread (4,2) (Gp:) Bloque (1,1)
(Gp:) gridDim.x (Gp:) gridDim.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 threadIdx y blockIdx para referenciar el identificador del thread dentro del bloque y al bloque dentro del grid, respectivamente (Gp:) blockDim.x (Gp:) blockDim.y (Gp:) blockDim.z
(Gp:) blockIdx.x=1 blockIdx.y=1
(Gp:) threadIdx.x=0 threadIdx.y=2
27
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 (Gp:) tiempo (Gp:) warp 8 instrucción 11 (Gp:) warp 1 instrucción 42 (Gp:) warp 3 instrucción 95 (Gp:) warp 8 instrucción 12 (Gp:) … (Gp:) warp 3 instrucción 96
(Gp:) Bloque 1 (Gp:) Bloque 2 (Gp:) Bloque n (Gp:) warp 1 (Gp:) 2 (Gp:) m (Gp:) warp 1 (Gp:) 2 (Gp:) m (Gp:) warp 1 (Gp:) 2 (Gp:) m
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 28
Planificación Tres flujos de instrucciones: warp1, warp3 y warp8 t=k t=k+1 t=k+2 t=l>k t=l+1 Planifica en tiempo k (Gp:) warp 8 instrucción 11
(Gp:) warp 1 instrucción 42
(Gp:) warp 3 instrucción 95
(Gp:) warp 8 instrucción 12
… (Gp:) warp 3 instrucción 96
29
Planificación (Gp:) warp 8 instrucción 11
(Gp:) warp 1 instrucción 42
(Gp:) warp 3 instrucción 95
(Gp:) warp 8 instrucción 12
… (Gp:) warp 3 instrucción 96
t=k t=k+1 t=k+2 t=l>k t=l+1 Planifica en tiempo k+1 Tres flujos de instrucciones: warp1, warp3 y warp8 30
Índice Introducción Arquitectura CUDA Modelo de ejecución Programación Rendimiento 31
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 32
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__ 33
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 34
Variables específicas dimensiones 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
35
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: 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)
36
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>(…); 37
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(n, 2.0, x, y); 38
Índice Introducción Arquitectura CUDA Modelo de ejecución Programación Rendimiento 39
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
40
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 41
Solapar computación y comunicación 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 42 Datos2 Datos1 Datos2 Datos1 Host?Device Device?Host Kernel Cálculos Datos1 Datos2 Datos2 Host?Device Device ? Host Kernel Cálculos Datos1 Datos2 Datos1 Datos2 Datos1 42
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 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 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 1 2 3 4 5 6 7 8 … memoria global matriz 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 1 2 3 4 5 6 7 8 … memoria global matriz 44
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) (Gp:) Banco 0 (Gp:) Banco 1 (Gp:) Banco 2 (Gp:) Banco 3 (Gp:) Banco 4 (Gp:) Banco 14 (Gp:) Banco 15 (Gp:) Thread 0 (Gp:) Thread 1 (Gp:) Thread 2 (Gp:) Thread 3 (Gp:) Thread 4 (Gp:) Thread 14 (Gp:) Thread 15 (Gp:) . . . (Gp:) . . . (Gp:) Sin conflictos (Gp:) Direccionamiento lineal Stride = 1
(Gp:) Banco 0 (Gp:) Banco 1 (Gp:) Banco 2 (Gp:) Banco 3 (Gp:) Banco 4 (Gp:) Banco 14 (Gp:) Banco 15 (Gp:) Thread 0 (Gp:) Thread 1 (Gp:) Thread 2 (Gp:) Thread 3 (Gp:) Thread 4 (Gp:) Thread 14 (Gp:) Thread 15 (Gp:) . . . (Gp:) . . . (Gp:) Sin conflictos (Gp:) Direccionamiento aleatorio Permutación 1:1
(Gp:) Banco 0 (Gp:) Banco 1 (Gp:) Banco 2 (Gp:) Banco 3 (Gp:) Banco 4 (Gp:) Banco 5 (Gp:) Banco 15 (Gp:) Thread 0 (Gp:) Thread 1 (Gp:) Thread 2 (Gp:) Thread 3 (Gp:) Thread 8 (Gp:) Thread 15 (Gp:) . . . (Gp:) . . . (Gp:) . . . (Gp:) Thread 9 (Gp:) Con conflictos (Gp:) Direccionamiento lineal Stride = 2
45
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 46
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 47
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 48
Divergencia de caminos (Gp:) Branch (Gp:) Path A (Gp:) Path B
Salto 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 49
Índice Producto Matriz-Matriz ¿Cómo explotar más la arquitectura? Buenas prácticas de programación CUDA OpenCL 50
Í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 OpenCL 51
Producto Matriz-Matriz Implementación secuencial 52
Producto Matriz-Matriz cudaMalloc() cudaMemcpy() cudaThreadSynchronize(); cudaFree() multiplication>(d_P, d_M, d_N, M.width, N.width); __global__ void multiplication(float *P, float *M, float *N, int wM, int wN)
53
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
54
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 GigaFlops Ejercicio 2: (En kernel) a) Cálculo de índices b) Cálculo de elemento Ejercicio 3: (Pruebas)
55
Producto Matriz-Matriz Implementación básica Ejercicio 2.a (Ejemplo)
56 (Gp:) (0,0) (Gp:) (1,0) (Gp:) (2,0) (Gp:) (0,1) (Gp:) (1,1) (Gp:) (2,1) (Gp:) (0,2) (Gp:) (1,2) (Gp:) (2,2)
(Gp:) (0,0) (Gp:) (1,0) (Gp:) (2,0) (Gp:) (0,1) (Gp:) (1,1) (Gp:) (2,1) (Gp:) (0,2) (Gp:) (1,2) (Gp:) (2,2)
(Gp:) (0,0) (Gp:) (1,0) (Gp:) (2,0) (Gp:) (0,1) (Gp:) (1,1) (Gp:) (2,1) (Gp:) (0,2) (Gp:) (1,2) (Gp:) (2,2)
(Gp:) (0,0) (Gp:) (1,0) (Gp:) (2,0) (Gp:) (0,1) (Gp:) (1,1) (Gp:) (2,1) (Gp:) (0,2) (Gp:) (1,2) (Gp:) (2,2)
(Gp:) (0,0) (Gp:) (1,0) (Gp:) (2,0) (Gp:) (0,1) (Gp:) (1,1) (Gp:) (2,1) (Gp:) (0,2) (Gp:) (1,2) (Gp:) (2,2)
(Gp:) (0,0) (Gp:) (1,0) (Gp:) (2,0) (Gp:) (0,1) (Gp:) (1,1) (Gp:) (2,1) (Gp:) (0,2) (Gp:) (1,2) (Gp:) (2,2)
5 3 (Gp:) (0,0) (Gp:) (1,0) (Gp:) (0,1) (Gp:) (1,1) (Gp:) (0,2) (Gp:) (1,2)
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 57
GPU GeForce 285GTX 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×2 + 2 Flops/MP) x 1476 MHz = 797 GFLOPS
1GB RAM 1242 MHz (double data rate, 2484 MHz) 512 bits (ancho bus) 159 GB/s = (2484 MHz) x (512 / 8 bytes)
58
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? 59 $Lt_0_7: ld.global.f32 %f2, [%r15+0]; ld.global.f32 %f3, [%r17+0]; mad.f32 %f1, %f2, %f3, %f1; add.s32 %r17, %r17, 4096; add.s32 %r15, %r15, 4; setp.ne.u32 %p1, %r15, %r16; @%p1 bra $Lt_0_7;
Accesos a memoria ¼ operaciones son cargas (240 cores) x (¼ cargas) x (4bytes/carga) x (1476 MHz) = 354 GBs > 159 GBs
El ancho de banda con memoria global no es suficiente
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
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
Producto de sub-matrices
Cada sub-matriz es calculada por un bloque de threads Los datos de entrada necesarios son llevados a memoria compartida 61
Producto Matriz-Matriz Implementación memoria compartida Disminuir accesos a memoria ¿Cómo lo haríais? 62
Producto Matriz-Matriz 63 (Gp:) C1,0 (Gp:) C0,0 (Gp:) C0,1 (Gp:) C2,0 (Gp:) C3,0 (Gp:) C1,1 (Gp:) C0,2 (Gp:) C2,2 (Gp:) C3,2 (Gp:) C1,2 (Gp:) C3,1 (Gp:) C2,1 (Gp:) C0,3 (Gp:) C2,3 (Gp:) C3,3 (Gp:) C1,3
(Gp:) A2,0 (Gp:) A1,1 (Gp:) A1,0 (Gp:) A0,0 (Gp:) A0,1 (Gp:) A3,0 (Gp:) A2,1 (Gp:) A3,1 (Gp:) A2,2 (Gp:) A1,3 (Gp:) A1,2 (Gp:) A0,2 (Gp:) A0,3 (Gp:) A3,2 (Gp:) A2,3 (Gp:) A3,3
(Gp:) B0,3 (Gp:) B1,3 (Gp:) B1,2 (Gp:) B1,1 (Gp:) B1,0 (Gp:) B0,0 (Gp:) B0,1 (Gp:) B0,2 (Gp:) B2,3 (Gp:) B3,3 (Gp:) B3,2 (Gp:) B3,1 (Gp:) B3,0 (Gp:) B2,0 (Gp:) B2,1 (Gp:) B2,2
Producto Matriz-Matriz 64 orden accesos 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 Cada thread calcula un punto de la matriz resultado
Producto Matriz-Matriz 65 C1,0 C0,0 C0,1 C2,0 C3,0 C1,1 C0,2 C2,2 C3,2 C1,2 C3,1 C2,1 C0,3 C2,3 C3,3 C1,3 (Gp:) A2,0 (Gp:) A1,1 (Gp:) A1,0 (Gp:) A0,0 (Gp:) A0,1 (Gp:) A3,0 (Gp:) A2,1 (Gp:) A3,1 (Gp:) A2,2 (Gp:) A1,3 (Gp:) A1,2 (Gp:) A0,2 (Gp:) A0,3 (Gp:) A3,2 (Gp:) A2,3 (Gp:) A3,3
(Gp:) B0,3 (Gp:) B1,3 (Gp:) B1,2 (Gp:) B1,1 (Gp:) B1,0 (Gp:) B0,0 (Gp:) B0,1 (Gp:) B0,2 (Gp:) B2,3 (Gp:) B3,3 (Gp:) B3,2 (Gp:) B3,1 (Gp:) B3,0 (Gp:) B2,0 (Gp:) B2,1 (Gp:) B2,2
(1,1) (1,0) (0,0) (0,1) threads C0,0 C1,0 C1,1 C0,1 calculan memoria compartida As Bs
cada thread lleva un dato de A y B
Producto Matriz-Matriz 66 C1,0 C0,0 C0,1 C2,0 C3,0 C1,1 C0,2 C2,2 C3,2 C1,2 C3,1 C2,1 C0,3 C2,3 C3,3 C1,3 (Gp:) A2,0 (Gp:) A1,1 (Gp:) A1,0 (Gp:) A0,0 (Gp:) A0,1 (Gp:) A3,0 (Gp:) A2,1 (Gp:) A3,1 (Gp:) A2,2 (Gp:) A1,3 (Gp:) A1,2 (Gp:) A0,2 (Gp:) A0,3 (Gp:) A3,2 (Gp:) A2,3 (Gp:) A3,3
(Gp:) B0,3 (Gp:) B1,3 (Gp:) B1,2 (Gp:) B1,1 (Gp:) B1,0 (Gp:) B0,0 (Gp:) B0,1 (Gp:) B0,2 (Gp:) B2,3 (Gp:) B3,3 (Gp:) B3,2 (Gp:) B3,1 (Gp:) B3,0 (Gp:) B2,0 (Gp:) B2,1 (Gp:) B2,2
(1,1) (1,0) (0,0) (0,1) threads C0,0 C1,0 C1,1 C0,1 calculan memoria compartida As Bs A0,0 A1,0 A1,1 A0,1 B0,0 B1,0 B1,1 B0,1 C0,0 = A0,0 x B0,0 + A1,0 x B0,1 C1,0 = A0,0 x B1,0 + A1,0 x B1,1 C0,1 = A0,1 x B0,0 + A1,1 x B0,1 C1,1 = A0,1 x B1,0 + A1,1 x B1,1
Producto Matriz-Matriz 67 C1,0 C0,0 C0,1 C2,0 C3,0 C1,1 C0,2 C2,2 C3,2 C1,2 C3,1 C2,1 C0,3 C2,3 C3,3 C1,3 (Gp:) A2,0 (Gp:) A1,1 (Gp:) A1,0 (Gp:) A0,0 (Gp:) A0,1 (Gp:) A3,0 (Gp:) A2,1 (Gp:) A3,1 (Gp:) A2,2 (Gp:) A1,3 (Gp:) A1,2 (Gp:) A0,2 (Gp:) A0,3 (Gp:) A3,2 (Gp:) A2,3 (Gp:) A3,3
(Gp:) B0,3 (Gp:) B1,3 (Gp:) B1,2 (Gp:) B1,1 (Gp:) B1,0 (Gp:) B0,0 (Gp:) B0,1 (Gp:) B0,2 (Gp:) B2,3 (Gp:) B3,3 (Gp:) B3,2 (Gp:) B3,1 (Gp:) B3,0 (Gp:) B2,0 (Gp:) B2,1 (Gp:) B2,2
(1,1) (1,0) (0,0) (0,1) threads C0,0 C1,0 C1,1 C0,1 calculan memoria compartida As Bs A2,0 A3,0 A3,1 A2,1 B0,2 B1,2 B1,3 B0,3 C0,0 = A0,0 x B0,0 + A1,0 x B0,1 + A2,0 x B0,2 + A3,0 x B0,3 C1,0 = A0,0 x B1,0 + A1,0 x B1,1 + A2,0 x B1,2 + A3,0 x B1,3 C0,1 = A0,1 x B0,0 + A1,1 x B0,1 + A2,1 x B0,2 + A3,1 x B0,3 C1,1 = A0,1 x B1,0 + A1,1 x B1,1 + A2,1 x B1,2 + A3,1 x B1,3
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 GigaFlops Ejercicio 2: (En kernel) a) Cálculo de índices b) Cálculo de elemento Ejercicio 3: (Pruebas)
69
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 70
Índice Producto Matriz-Matriz ¿Cómo explotar más la arquitectura? Buenas prácticas de programación CUDA OpenCL 71
¿Cómo explotar más la arquitectura? Sistemas Distribuidos mediante MPI
Sistemas compartidos mediante OpenMP 72
Índice Producto Matriz-Matriz ¿Cómo explotar más la arquitectura? Buenas prácticas de programación CUDA OpenCL 73
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 74
Índice Producto Matriz-Matriz ¿Cómo explotar más la arquitectura? Buenas prácticas de programación CUDA OpenCL ¿Qué es OpenCL? ¿Qué diferencia OpenCL de CUDA? Ejemplo: Suma de vectores 75
¿Qué es OpenCL? OpenCL: Open Computing Language
Propuesto por a
¿Quién está involucrado?
¿Qué diferencia OpenCL de CUDA? (I) Punteros CUDA struct Node {Node* next} n=n->next Opencl struct Node {unsigned int next;} next=bufBase+n
¿Qué diferencia OpenCL de CUDA? (II) Kernels CUDA Programa compilado en formato binario OpenCL Se compila en tiempo de ejecución
Palabras Clave y lenguaje utilizado para los kernel
Ejemplo: Suma de vectores CUDA _global_ void SumaVec(const float *a, const float *b, float *c) // Índice al elemento del vector int indice=blockIdx.x*blockDim.x+threadIdx.x
c[indice]=a[indice]+b[indice] } OpenCL _kernel void SumaVec(_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 OpenCL 80