Í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