Publicado el 20 de Mayo del 2018
583 visualizaciones desde el 20 de Mayo del 2018
3,0 MB
7 paginas
RIO 2014
Río Cuarto (Argentina), 20 de Febrero, 2014
Contenidos de la charla [18 diapositivas]
1. Balanceo dinámico de la carga. [2]
2. Mejorando el paralelismo con streams. [5]
3. Ejecución dependiente de los datos. [2]
4. Algoritmos paralelos recursivos. [4]
5. Llamadas a librerías desde los kernels. [3]
6. Simplificar la división CPU/GPU. [2]
2
Manuel Ujaldón Martínez
Dpto. Arquitectura de Computadores. Univ. Málaga.
NVIDIA CUDA Fellow.
Seis formas de mejorar
nuestros códigos CUDA con Kepler
Balanceo dinámico de la carga
Desdoble de las colas de los SMX
para ocupar más y mejor la GPU
Ejecución dependiente de los datos
Algoritmos paralelos recursivos
Llamadas a librerías desde los kernels
Simplificar la división CPU/GPU
Ocupación
Ejecución
d
a
b ili d
a m a
r
g
o
P r
Paralelismo
dinámico y
planificación de
hilos en Kepler
3
Ejemplos de optimización para Kepler1. Balanceo dinámico de la carga Generación dinámica de la carga
Despliega paralelismo según el nivel de detalle
Asigna los recursos dinámicamente según se vaya
requiriendo precisión en tiempo real, lo que facilita la
computación de aplicaciones irregulares en GPU.
Malla gruesa
Malla fina
Malla dinámica
Rendimiento elevado,
precisión baja
Rendimiento bajo,
precisión elevada
Sacrifica rendimiento sólo
donde se requiere precisión
5
La potencia computacional
se asocia a las regiones
según su interés
CUDA hasta 2012:
• La CPU lanza kernels
de forma regular.
• Todos los píxeles
se procesan igual.
CUDA sobre Kepler:
• La GPU lanza un número
diferente de kernels/bloques
para cada región
computacional.
El gestor de kernels/mallas: Fermi vs. Kepler
Fermi
Colas de streams
(colas ordenadas de mallas)
Stream 1
Kernel C
Kernel B
Kernel A
Stream 2
Kernel R
Kernel Q
Kernel P
Stream 3
Kernel Z
Kernel Y
Kernel X
Una sola cola hardware
multiplexa los streams
Distribuidor de carga
para los bloques lanzados desde las mallas
16 mallas activas
Kepler GK110
Cola de streams
C
B
A
R
Q
P
Z
Y
X
Hardware paralelo de streams
Gestor de kernels/mallas
Mallas pendientes y suspendidas
Miles de mallas pendientes
Permite suspender mallas
Distribuidor de carga
Se encarga de las mallas activas
32 mallas activas
A
D
U
C
r
o
p
a
d
a
r
e
n
e
g
o
j
a
b
a
r
t
e
d
a
g
r
a
C
SM
SM
SM
SM
SMX
SMX
SMX
SMX
6
8
2. Mejorando el paralelismo con streams Relación entre las colas software y hardware
Relación entre las colas software y hardware
...pero los streams se multiplexan en una cola única
...pero los streams se multiplexan en una cola única
Fermi:
El hardware de
la GPU puede
albergar hasta
16 mallas en
ejecución...
A--B--C P--Q--R X--Y--Z
Oportunidad para solapar:
Sólo en las fronteras entre streams
A -- B -- C
Stream 1
P -- Q -- R
Stream 2
X -- Y -- Z
Stream 3
Fermi:
El hardware de
la GPU puede
albergar hasta
16 mallas en
ejecución...
A--B--C P--Q--R X--Y--Z
Oportunidad para solapar:
Sólo en las fronteras entre streams
Kepler:
Desaparecen las dependencias entre streams
El número
de mallas
en ejecución
crece hasta 32
A--B--C
P--Q--R
X--Y--Z
9
Concurrencia plena entre streams
Caso estudio para explotar la concurrencia de
la GPU en Fermi (15 SMs) y Kepler (15 SMXs)
mykernel <<< 100, 128, ... >>> [Aquí tenemos un déficit en warps]
Lanza 100 bloques de 128 hilos (4 warps), esto es, 400 warps.
Hay 26.66 warps para cada multiprocesador, ya sea SM o SMX.
En Fermi: Hasta 48 warps activos (21 bajo el límite), que no puede aprovecharse.
En Kepler: Hasta 64 warps activos (37 bajo el límite), que pueden activarse desde
hasta un máx. de 32 llamadas a kernels desde: MPI, threads POSIX, streams CUDA.
mykernel <<< 100, 384, ... >>>
Lanza 100 bloques de 384 hilos (12 warps), esto es, 1200 warps.
Hay 80 warps para cada multiprocesador. Hemos alcanzado el máx.
de 64 warps activos, así que 16 warps * 15 SMX = 240 warps esperan
en colas de Kepler para ser activados con posterioridad.
mykernel <<< 1000, 32, ... >>>[Aquí tenemos un exceso de bloques]
66.66 bloques para cada SMX, pero el máx. es 16. Mejor <100, 320>
11
Lecciones a aprender (y conflictos asociados)
Bloques suficientemente grandes para evitar el límite de 16
por cada SMX.
Pero los bloques consumen memoria compartida, y alojar más
memoria compartida significa menos bloques y más hilos por bloque.
Suficientes hilos por bloque como para saturar el límite de
64 warps activos por cada SMX.
Pero los hilos consumen registros, y utilizar muchos registros nos
lleva a tener menos hilos por bloque y más bloques.
Sugerencias:
Al menos 3-4 bloques activos, cada uno con al menos 128 hilos.
Menos bloques cuando la memoria compartida es crítica, pero...
... abusar de ella penaliza la concurrencia y la ocultación de latencia.
A -- B -- C
Stream 1
P -- Q -- R
Stream 2
X -- Y -- Z
Stream 3
A -- B -- C
Stream 1
P -- Q -- R
Stream 2
X -- Y -- Z
Stream 3
10
12
Paralelismo dependiente
del volumen de datos (2)
El programa paralelo más elemental:
Los bucles son paralelizables.
Conocemos la carga de trabajo a priori.
for i = 1 to N
for j = 1 to M
convolution (i, j);
El programa imposible más elemental:
Desconocemos la carga de trabajo.
El reto es la partición de los datos.
M
max(x[i])
N
N
for i = 1 to N
for j = 1 to x[i]
convolution (i, j);
Una solución mala: Superconjunto.
Una solución peor: Serialización.
Lo que hace posible el paralelismo dinámico:
Los dos lazos se ejecutan en paralelo
El programa CUDA para Kepler:
__global__ void convolution(int x[])
{
for j = 1 to x[blockIdx] // Cada bloque lanza x[blockIdx] ...
kernel <<< ... >>> (blockIdx, j) // ... kernels desde la GPU
}
convolution <<< N, 1 >>> (x); // Lanza N bloques de un hilo en GPU
// (las filas comienzan en paralelo)
s
e
u
q
o
b
l
N
x[blockIdx] llamadas a kernels
Intercambiando estos dos parámetros,
el programa es más rápido, pero no sirve para
más de 1024 filas (máximo tamaño del bloque).
En CUDA 5.0 se permiten hasta 24 lazos anidados.
15
3. Ejecución dependiente de los datos4. Algoritmos paralelos recursivos Un ejemplo sencillo de recursividad paralela:
Quicksort
Es el típico algoritmo divide y vencerás que cuesta a Fermi
La ejecución depende de los datos.
Los datos se particionan y ordenan recursivamente.
Algoritmos paralelos recursivos
antes de Kepler
Los primeros modelos de programación CUDA no
soportaban recursividad de ningún tipo.
CUDA comenzó a soportar funciones recursivas en la
versión 3.1, pero podían fallar perfectamente si el tamaño de
los argumentos era considerable.
En su lugar, puede utilizarse una pila definida por el
usuario en memoria global, pero a costa de una considerable
merma en rendimiento.
Gracias al paralelismo dinámico, podemos aspirar a una
solución eficiente para GPU.
17
El código CUDA para quicksort
Resultados experimentales para Quicksort
El número de líneas de código se reduce a la mitad.
El rendimiento se mejora en un factor 2x.
Versión ineficiente
Versión más eficiente en Kepler
_global_ void qsort(int *data, int l, int r)
{
int pivot = data[0];
int *lptr = data+l, *rptr = data+r;
// Particiona datos en torno al pivote
partition(data, l, r, lptr, rptr, pivot);
_global_ void qsort(int *data, int l, int r)
{
int pivot = data[0];
int *lptr = data+l, *rptr = data+r;
// Particiona datos en torno al pivote
partition(data, l, r, lptr, rptr, pivot);
// Lanza la siguiente etapa recursivamente
int rx = rptr-data; lx = lptr-data;
if (l < rx)
qsort<<<...>>>(data,l,rx);
if (r > lx)
qsort<<<...>>>(data,lx,r);
}
Las ordenaciones de la parte derecha
e izquierda se serializan
// Utiliza streams para la recursividad
cudaStream_t s1, s2;
cudaStreamCreateWithFlags(&s1, ...);
cudaStreamCreateWithFlags(&s2, ...);
int rx = rptr-data; lx = lptr-data;
if (l < rx)
qsort<<<...,0,s1>>>(data,l,rx);
if (r > lx)
qsort<<<...,0,s2>>>(data,lx,r);
}
Utiliza "streams" separados
para lograr concurrencia
19
18
20
Conceptos básicos del modelo CUDA:
Sintaxis y semántica en tiempo de ejecución
__device__ float buf[1024];
__global__ void dynamic(float *data)
{
int tid = threadIdx.x;
if (tid % 2)
buf[tid/2] = data[tid]+data[tid+1];
__syncthreads();
if (tid == 0) {
launchkernel<<<128,256>>>(buf);
cudaDeviceSynchronize();
}
__syncthreads();
if (tid == 0) {
cudaMemCpyAsync(data, buf, 1024);
cudaDeviceSynchronize();
}
}
Este lanzamiento se produce para cada hilo
CUDA 5.0: Espera a que concluyan todos los
lanzamientos y llamadas que el bloque haya
efectuado anteriormente.
Los hilos sin trabajo esperan al resto aquí
CUDA 5.0: Sólo se permiten lanzamientos
asíncronos para la recogida de datos
22
Un ejemplo de llamada sencilla a una librería
utilizando cuBLAS (ya disponible para CUDA 5.0)
La relación padre-hijo en bloques CUDA
La CPU lanza
el kernel
__global__ void libraryCall(float *a,
float *b,
float *c)
{
// Todos los hilos generan datos
createData(a, b);
__syncthreads();
// El primer hilo llama a librería
if (threadIdx.x == 0) {
cublasDgemm(a, b, c);
cudaDeviceSynchronize();
}
// Todos los hilos esperan los resultados
__syncthreads();
consumeData(c);
}
Generación
de datos por
cada bloque
Llamadas
a la librería
externa
Uso del
resultado
en paralelo
__global__ void libraryCall(float *a,
float *b,
float *c)
{
// Todos los hilos generan datos
createData(a, b);
__syncthreads();
// El primer hilo llama a la librería
if (threadIdx.x == 0) {
cublasDgemm(a, b, c);
cudaDeviceSynchronize();
}
Se ejecuta
la función
externa
// Todos los hilos esperan los resultados
__syncthreads();
consumeData(c);
}
Ejecución por cada hilo
Una solo llamada a la función de
Comentarios de: Ejemplos de optimización para Kepler (0)
No hay comentarios