PDF de programación - Kepler

RIO 2014

Río Cuarto, 20 de Febrero, 2014

Contenidos de la charla

Manuel Ujaldón Martínez
Dpto. Arquitectura de Computadores. Univ. Málaga.

NVIDIA CUDA Fellow.

1. Presentación de la arquitectura [7 diapositivas]
2. La memoria y el transporte de datos [9]
3. Los cores SMX [9]
4. Desplegando todo el paralelismo en Kepler [12]
5. Mejoras funcionales [7]

1. Paralelismo dinámico.
2. Hyper-Q.

6. Optimizaciones futuras []

1. Vectorización: El tamaño del warp.
2. Stacked-DRAM: Memoria apilada sobre el procesador.

Kepler, Johannes
(1571-1630)

Autor de las leyes del movimiento planetario.

Primera ley: Las órbitas de los planetas son planas. El sol está en

el plano de la órbita. La trayectoria del planeta respecto del sol es
una elipse en la que el sol ocupa uno de los fotos.

Segunda ley: El radio vector que une al sol y el planeta barre

áreas iguales en tiempos iguales. Un planeta se mueve más
rápidamente en su perihelio que en su afelio, y mientras más
excéntrica sea su órbita, mayor será la diferencia de velocidad entre
sus extremos.

Tercera ley: Los cuadrados de los períodos de revolución en torno

al sol son proporcionales a los cubos de los semiejes mayores de las
órbitas. La velocidad media con que un planeta recorre su órbita
disminuye a medida que el planeta está más lejos del sol. La influen-
cia que el sol ejerce sobre los planetas disminuye con la distancia.

2

4

Kepler1. Presentación de la arquitectura Nuestra Kepler también tiene tres leyes

Y tres innovaciones principales

Consumo

Rendimiento

Programabilidad

SMX:

Un multiprocesador con más
recursos y menos consumo.

Paralelismo
dinámico:

La GPU es más autónoma,
puede lanzar kernels CUDA.

Hyper-Q:

Múltiples kernels pueden
compartir los SMX .

5

Resumen de sus rasgos más sobresalientes

Fabricación: 7100 Mt. integrados a 28 nm. por TSMC.
Arquitectura: Entre 7 y 15 multiprocesadores SMX,

dotados de 192 cores cada uno.

Modelos más populares: K20 (con 13 SMX), K20X (14), K40 (15).

Aritmética: Más de 1 TeraFLOP en punto flotante de doble

precisión (formato IEEE-754 de 64 bits).

Los valores concretos dependen de la frecuencia de cada modelo.
Con sólo 10 racks de servidores, podemos alcanzar 1 PetaFLOP.

Principales innovaciones en el diseño de los cores:

Paralelismo dinámico.
Planificación de hilos (Hyper-Q).

7

Modelos comerciales de Kepler:
GeForce y Tesla frente a frente

GeForce GTX Titan

Diseñada para jugar:

El precio es prioritario (<500€).
Gran disponibilidad/popularidad.
Poca memoria de vídeo (1-2 GB.).
Relojes un poco más rápidos.
Hyper-Q sólo para streams CUDA.
Perfecta para desarrollar código

Orientada a HPC:
Fiabilidad (tres años de garantía).
Pensada para conectar en clusters.
Más memoria de vídeo (6-12 GB.).
Ejecución sin descanso (24/7).
Hyper-Q para procesos MPI.
GPUDirect (RDMA) y otras

que luego pueda disfrutar Tesla.

coberturas para clusters de GPUs.

6

8

Kepler GK110: Disposición física de las UFs
para la Tesla K40 dotada de 15 SMXs

El multiprocesador SMX

Planificación y emisión
de instrucciones en warps

Front-end

Ejecución de instrucciones.
512 unidades funcionales:
- 192 para aritmética entera.
- 192 para aritmética s.p.
- 64 para aritmética d.p.
- 32 para carga/almacen.
- 32 para SFUs (log,sqrt, ...)

Acceso a memoria

9

Back-end

Interfaz

10

La memoria en las Teslas: Fermi vs. Kepler

Tarjeta gráfica Tesla

M2075 M2090

K20

Registros de 32 bits / multiprocesador

Tamaño caché L1 + mem. compartida

Anchura de los 32 bancos de m.c.

32768

64 KB.

32 bits

32768

64 KB.

32 bits

65536

64 KB.

64 bits

K20X

65536

64 KB.

64 bits

K40

65536

64 KB.

64 bits

Frecuencia de SRAM (la de la GPU)

575 MHz

650 MHz

706 MHz

732 MHz 745, 810, 875

Ancho de banda L1 y mem. compartida

73.6 GB/s. 83.2 GB/s. 180.7 GB/s 187.3 GB/s 216.2 GB/s.

Tamaño de la caché L2

768 KB.

768 KB.

1.25 MB.

1.5 MB.

1.5 MB.

Ancho de banda L2 (bytes por ciclo)

384

384

1024

1024

1024

L2 en ops. atómicas (dir. compartida)

1/9 por clk 1/9 por clk

1 por clk

1 por clk

1 per clk

L2 en ops. atómicas (dirs. separadas)

24 por clk 24 por clk 64 por clk 64 por clk

64 per clk

Anchura del bus de memoria DRAM

384 bits

384 bits

320 bits

384 bits

384 bits

Frecuencia de la memoria (MHz)

2x 1500

2x 1850

2x 2600

2x 2600

2 x 3000

Ancho de banda DRAM (no ECC)

144 GB/s. 177 GB/s. 208 GB/s. 250 GB/s.

288 GB/s.

Tamaño DRAM (todas GDDR5)

6 GB.

6 GB.

5 GB.

6 GB.

12 GB.

Bus externo de conexión con la CPU

PCI-e 2.0 PCI-e 2.0

PCI-e 3.0

PCI-e 3.0

PCI-e 3.0

12

2. La memoria y el transporte de datos Diferencias en la jerarquía de memoria:
Fermi vs. Kepler

Motivación para usar la nueva caché de datos

48 Kbytes extra para expandir el tamaño de la caché L1.
Posee el mayor ancho de banda en caso de fallo a caché.
Usa la caché de texturas, pero de forma transparente al
programador, y elimina el tiempo de configuración de ésta.

Permite que una dirección global pueda buscarse y
ubicarse en esta caché, utilizando para ello un camino
separado del que accede a caché L1 y memoria compartida.

Es flexible, no requiere que los accesos estén alineados.
Gestionada automáticamente por el compilador.

Cómo utilizar la nueva caché de datos

Declarar los punteros con el prefijo "const __restrict__".
El compilador automáticamente mapeará la carga de esos

valores en la caché para usar el nuevo camino de datos a
través de la memoria de texturas.

__global__ void saxpy(float x, float y,
const float * __restrict__ input,
float * output)
{
size_t offset = threadIdx.x +
(blockIdx.x * blockDim.x);

// El compilador utilizará la nueva caché para "input"
output[offset] = (input[offset] * x) + y;
}

13

15

Comparativa con la memoria de constantes

A comparar

Memoria de constantes

Caché de datos de sólo lectura

Disponibilidad

Desde CUDA Compute Capability 1.0

A partir de CCC 3.5

(aunque desde CCC 1.0 se podía usar
la memoria de texturas manualmente)

Tamaño

64 Kbytes

48 Kbytes

Implementación
hardware

Una partición de

la memoria global (DRAM)

Caché de texturas

que expande la L1 (SRAM)

Acceso

Mejor rasgo

Peor rasgo

Mejor
escenario
de uso

A través de una caché de 8 Kbytes
que posee cada multiprocesador

SM(X)

Latencia muy baja

Menor ancho de banda

Mediante un camino aparte en
el cauce de segmentación gráfico

Gran ancho de banda

Mayor latencia

Acceso con el mismo coeficiente

Cuando el kernel es memory-bound,

(sin usar threadIdx) a un pequeño

aún después de haber saturado el ancho

conjunto de datos de sólo lectura

de banda con memoria compartida

14

16

Comunicación entre las memorias de las GPU

Kepler + CUDA 5 soportan GPUDirect-RDMA
[Remote Direct Memory Access]

En Fermi se puso en marcha GPU Direct 1.0 para permitir

la comunicación entre GPUs dentro de clusters de CPUs.

Esto permite una transferencia más directa entre GPUs.
Normalmente, el enlace es PCI-express o InfiniBand.

Receptor

Transmisor

17

GPUDirect-RDMA en Maxwell

La situación será más compleja en la próxima generación

de GPUs, pues tendrán un espacio de memoria unificado
con la CPU.

Resultados preliminares de GPUDirect-RDMA
(mejoran con CUDA 6.0 & OpenMPI)

)
s
o
d
n
u
g
e
s
o
r
c
m

i

(

U
P
G
-
U
P
G

a
c
n
e
t
a
L

i

)
s
o
d
n
u
g
e
s
(

l
a
t
o
t

n
ó
c
u
c
e
j
e

e
d

o
p
m
e
T

i

i

Tamaño del mensaje (bytes)

Número lateral

Latencia inter-nodo usando:

GPUs Tesla K40m (no GeForces).
Librería MPI MVAPICH2.
ConnectX-3, IVB 3GHz.

Mejor escalado en MPI:
Código: HSG (bioinformática).
2 nodos de GPU.
4 procesos MPI por nodo.

19

18

20

Un breve recordatorio de CUDA

GPU

Multiprocesador N

Multiprocesador 2

Multiprocesador 1

Memoria compartida

Registros

Registros

Registros

Procesador 1

Procesador 2

…

Procesador M

Memoria global

Unidad

de

Control
SIMD

Caché para
constantes

Caché para
texturas

Memoria
integrada
en la GPU

Memoria
externa
a la GPU
(incluida
dentro de
la tarjeta
gráfica)

Hilo

···

Bloque de hilos

· · ·

· · ·

· · ·

Malla 0

· · ·

· · ·

· · ·

(grid)

···

···

···

···

···

···

Malla 1

···

···

···

22

... y de cómo va escalando la arquitectura

Kepler en perspectiva:
Recursos hardware y rendimiento pico

Arquitectura

Tesla
G80

Tesla
GT200

Fermi
GF100

Fermi
GF104

Kepler
GK104

Kepler
GK110

Marco temporal

2006-07 2008-09 2010

2011

2012

2013-14

CUDA Compute
Capability (CCC)

1.0

1.2

2.0

2.1

3.0

3.5

N (multiprocs.)

M (cores/multip.)

16

8

30

8

16

32

7

48

8

15

192

192

Número de cores

128

240

512

336

1536

2880

Tarjeta Tesla (modelo comercial)

M2075 M2090

K20

K20X

Modelo GeForce similar en cores

GTX 470 GTX 580

-

GTX Titan

K40

-

Generación de GPU (y CCC)

Fermi GF100

mi GF100 (2.0)

KepKepler GK110 (

er GK110 (3.5)

Multiprocesadores x (cores/multipr.)

14 x 32

16 x 32

13 x 192

14 x 192

15 x 192

Número total de cores

Tipo de multiprocesador

448

512

2496

2688

2880

SMSM

SMX con p
SMX con paralelismo diná

lismo dinámico y HyperQ

Distancia de integración de los trans.

40 nm.

40 nm.

28 nm.

28 nm.

28 nm.

Frecuencia de la GPU (para gráficos)

575 MHz 650 MHz 706 MHz 732 MHz 745,810,875 MHz

Frecuencia de los cores (para GPGPU)

1150 MHz 1300 MHz 706 MHz 732 MHz 745,810,875 MHz

Número de cores (simple precisión)

GFLOPS (pico en simple precisión)

Número de cores (doble precisión)

GFLOPS (pico en doble precisión)

448

1030

224

515

512

1331

256

665

2496

3520

832

1170

2688

3950

896

1310

2880

4290

960

1680

23

24

3. Los cores SMX GPU Boost

Cada aplicación tiene un comportamiento
distinto en relación al consumo

Permite acelerar hasta un 17% el reloj de la GPU si el

Aquí vemos el consumo medio (vatios) en la Tesla K20X

consumo de una aplicación es bajo.

de aplicaciones muy populares en el ámbito HPC:

Se retornará al reloj base si se sobrepasan los 235 W.
Se