Transcript proyectos:gopac:presentacion_gpus.ppt

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 QuadCore 2 x Xeon QuadCore + Tesla s1070
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%
RC5-72
Distributed.net RC5-72 challenge client
code
1.979
218
>99%
FEM
Finite element modeling, simulation of 3D
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
2D 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
457 316
431 263
210
79
GPU Speedup
Relative to CPU
60
50
40
Ke rn e l
Ap p lic a tio n
30
20
10
0
H .2 6 4
LBM
R C 5 -7 2
F EM
R PES
PN S
SA XPY T PA C F
FDTD
M R I-Q
M R IFHD
•
•
•
•
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
8
Introducción
Facilidad de programación
Programabilidad
Explosión GPGPU
Inicios GPGPU
Brook++
RapidMind
Cg
CUDA
OpenCL
ATI Stream
Futuro GPGPU?
OpenGL
DirectX
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.5TFLOPS (SP)
~800GFLOPS (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
• 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
16
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
17
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);
18
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);
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
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 Thread Thread Thread
(0,0)
(1,0)
(2,0)
(3,0)
(4,0)
Thread Thread Thread Thread Thread
(0,1)
(1,1)
(2,1)
(3,1)
(4,1)
Thread Thread Thread 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
•
Un thread queda indentificado por:
–
–
•
gridDim.x
gridDim.y
•
blockIdx.x=1
blockIdx.y=1
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 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
threadIdx.x=0
threadIdx.y=2
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 Thread Thread Thread
(0,0)
(1,0)
(2,0)
(3,0)
(4,0)
blockDim.y
•
Thread Thread Thread Thread Thread
(0,1)
(1,1)
(2,1)
(3,1)
(4,1)
Thread Thread Thread Thread Thread
(0,2)
(1,2)
(2,2)
(3,2)
(4,2)
blockDim.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 29
96
m
Planificación
• Tres flujos de instrucciones:
warp1, warp3 y warp8
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:
warp1, warp3 y warp8
Warp
Current
Instruction
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
•
•
•
•
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
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 KernelFunction(…);
dim3 DimGrid(100,50); // 5000 bloques
dim3 DimBlock(4,8,8); // 256 threads/bloque
KernelFunction<<< 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 = 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);
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
• 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
42
Solapar computación y comunicación
Datos1 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
Datos2
Kernel Cálculos
Datos1 Datos2
Device→Host
Datos1 Datos2
Datos1
Datos2
Datos1
Datos2
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
memoria global
matriz
1
1
1
2
3
4
2
1
2
3
4
2
5
6
7
8
3
5
6
7
8
3
4
4
9
10
11
12
5
9
10
11
12
5
13
14
15
16
6
13
14
15
16
6
7
7
8
8
…
…
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)
Sin conflictos
Sin conflictos
Con conflictos
Thread 0
Banco 0
Thread 0
Banco 0
Thread 0
Banco 0
Thread 1
Banco 1
Thread 1
Banco 1
Thread 1
Banco 1
Thread 2
Banco 2
Thread 2
Banco 2
Thread 2
Banco 2
Thread 3
Banco 3
Thread 3
Banco 3
Thread 3
Banco 3
Thread 4
Banco 4
Thread 4
Banco 4
...
Banco 4
...
...
...
...
Thread 8
Banco 5
...
Thread 14
Banco 14
Thread 14
Banco 14
Thread 9
Thread 15
Banco 15
Thread 15
Banco 15
Thread 15
Direccionamiento lineal
Stride = 1
Direccionamiento aleatorio
Permutación 1:1
...
Banco 15
46
Direccionamiento lineal
Stride = 2
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
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
Branch
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
50
Índice
• Producto Matriz-Matriz
• ¿Cómo explotar más la arquitectura?
• Buenas prácticas de programación
CUDA
• OpenCL
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
• OpenCL
52
Producto Matriz-Matriz
• Implementación secuencial
53
Producto Matriz-Matriz
•
•
•
•
•
cudaMalloc()
cudaMemcpy()
cudaThreadSynchronize();
cudaFree()
multiplication<<< grid, threads >>>(d_P, d_M, d_N,
M.width, N.width);
• __global__ void multiplication(float *P, float *M, float *N,
int wM, int wN)
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 GigaFlops
– 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,0)
(1,0)
(2,0)
(0,1)
(1,1)
(2,1)
(0,1)
(1,1)
(2,1)
(0,2)
(1,2)
(2,2)
(0,2)
(1,2)
(2,2)
(0,0)
(1,0)
(2,0)
(0,0)
(1,0)
(2,0)
(0,1)
(1,1)
(2,1)
(0,1)
(1,1)
(2,1)
(0,2)
(1,2)
(2,2)
(0,2)
(1,2)
(2,2)
(0,0)
(1,0)
(2,0)
(0,0)
(1,0)
(2,0)
(0,1)
(1,1)
(2,1)
(0,1)
(1,1)
(2,1)
(0,2)
(1,2)
(2,2)
(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 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 (8x2 + 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)
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.f32
ld.global.f32
mad.f32
add.s32
add.s32
setp.ne.u32
@%p1 bra
%f2, [%r15+0];
%f3, [%r17+0];
%f1, %f2, %f3, %f1;
%r17, %r17, 4096;
%r15, %r15, 4;
%p1, %r15, %r16;
$Lt_0_7;
– ¼ 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
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
B0,0 B1,0 B2,0 B3,0
B0,1 B1,1 B2,1 B3,1
B0,2 B1,2 B2,2 B3,2
B0,3 B1,3 B2,3 B3,3
A0,0 A1,0 A2,0 A3,0
C0,0 C1,0 C2,0 C3,0
A0,1 A1,1 A2,1 A3,1
C0,1 C1,1 C2,1 C3,1
A0,2 A1,2 A2,2 A3,2
C0,2 C1,2 C2,2 C3,2
A0,3 A1,3 A2,3 A3,3
C0,3 C1,3 C2,3 C3,3
64
Producto Matriz-Matriz
orden accesos
Cada thread calcula un punto de la matriz resultado
C0,0
thread0,0
C1,0
thread1,0
C0,1
thread0,1
C1,1
thread1,1
A0,0 x B0,0
A0,0 x B1,0
A0,1 x B0,0
A0,1 x B1,0
A1,0 x B0,1
A1,0 x B1,1
A1,1 x B0,1
A1,1 x B1,1
A2,0 x B0,2
A2,0 x B1,2
A2,1 x B0,2
A2,1 x B1,2
A3,0 x B0,3
A3,0 x B1,3
A3,1 x B0,3
A3,1 x B1,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
C0,0 C1,0
B0,0 B1,0 B2,0 B3,0
C0,1 C1,1
(0,1) (1,1)
B0,1 B1,1 B2,1 B3,1
B0,2 B1,2 B2,2 B3,2
cada thread lleva un dato de A y B
As
Bs
memoria compartida
B0,3 B1,3 B2,3 B3,3
A0,0 A1,0 A2,0 A3,0
C0,0 C1,0 C2,0 C3,0
A0,1 A1,1 A2,1 A3,1
C0,1 C1,1 C2,1 C3,1
A0,2 A1,2 A2,2 A3,2
C0,2 C1,2 C2,2 C3,2
A0,3 A1,3 A2,3 A3,3
C0,3 C1,3 C2,3 C3,3
66
Producto Matriz-Matriz
threads
(0,0) (1,0)
calculan
C0,0 C1,0
B0,0 B1,0 B2,0 B3,0
C0,1 C1,1
(0,1) (1,1)
B0,1 B1,1 B2,1 B3,1
C0,0 = A0,0 x B0,0 + A1,0 x B0,1
B0,2 B1,2 B2,2 B3,2
C1,0 = A0,0 x B1,0 + A1,0 x B1,1
B0,3 B1,3 B2,3 B3,3
C0,1 = A0,1 x B0,0 + A1,1 x B0,1
C1,1 = A0,1 x B1,0 + A1,1 x B1,1
A0,0 A1,0 A2,0 A3,0
C0,0 C1,0 C2,0 C3,0
A0,1 A1,1 A2,1 A3,1
C0,1 C1,1 C2,1 C3,1
As
Bs
A0,0 A1,0
B0,0 B1,0
A0,2 A1,2 A2,2 A3,2
C0,2 C1,2 C2,2 C3,2
A0,1 A1,1
B0,1 B1,1
A0,3 A1,3 A2,3 A3,3
C0,3 C1,3 C2,3 C3,3
memoria compartida
67
Producto Matriz-Matriz
threads
(0,0) (1,0)
calculan
C0,0 C1,0
B0,0 B1,0 B2,0 B3,0
C0,1 C1,1
(0,1) (1,1)
B0,1 B1,1 B2,1 B3,1
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
B0,2 B1,2 B2,2 B3,2
B0,3 B1,3 B2,3 B3,3
C1,1 = A0,1 x B1,0 + A1,1 x B1,1 + A2,1 x B1,2 + A3,1 x B1,3
A0,0 A1,0 A2,0 A3,0
C0,0 C1,0 C2,0 C3,0
A0,1 A1,1 A2,1 A3,1
C0,1 C1,1 C2,1 C3,1
As
Bs
A2,0 A3,0
B0,2 B1,2
A0,2 A1,2 A2,2 A3,2
C0,2 C1,2 C2,2 C3,2
A2,1 A3,1
B0,3 B1,3
A0,3 A1,3 A2,3 A3,3
C0,3 C1,3 C2,3 C3,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 GigaFlops
– 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
• OpenCL
72
¿Cómo explotar más la
arquitectura?
• Sistemas Distribuidos
mediante MPI
• Sistemas compartidos
mediante OpenMP
73
Índice
 Producto Matriz-Matriz
 ¿Cómo explotar más la
arquitectura?
• Buenas prácticas de programación
CUDA
• OpenCL
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
• OpenCL
– ¿Qué es OpenCL?
– ¿Qué diferencia OpenCL de CUDA?
– Ejemplo: Suma de vectores
76
¿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
81