Introducción
distances
vectorsum
sgemm
Más cosas
Resumen
Optimizaciones
Carlos Bederián, Nicolás Wolovick
FaMAF, Universidad Nacional de Córdoba, Argentina
1 de Agosto de 2012
[email protected]
Revisión 3742, 2012-08-03
Introducción
distances
vectorsum
sgemm
Más cosas
Resumen
Introducción
Ejemplo: distancias de un vector 4d
Ejemplo: suma de vector de a bloques
Ejemplo: explorando tamaños de bloque
Más cosas
Resumen
Introducción
distances
vectorsum
sgemm
Más cosas
Resumen
Introducción
Vamos a ver algunos ejemplos de optimizaciones típicas
• AoS vs. SoA.
• Atomics sobre globales y sobre shared.
• Exploración de tamaños de bloque.
Quedan muchos temas por ver.
En el práctico y el exámen aumentarán bastante su base de
conocimiento.
Introducción
distances
vectorsum
sgemm
Más cosas
Resumen
Introducción
Ejemplo: distancias de un vector 4d
Ejemplo: suma de vector de a bloques
Ejemplo: explorando tamaños de bloque
Más cosas
Resumen
Introducción
distances
vectorsum
sgemm
Más cosas
Resumen
Distancias de un vector
Obtener la norma de una lista de vectores. Es solo un map.
Entrada : float4[N] vectors.
Salida : float[N] dists.
Donde ∀i, dists[i] = |vectors[i]|.
• Hay una elección importante de estructura de datos.
• ¿Vector de cuatruplas ó cuatrupla de vectores?
Introducción
distances
vectorsum
sgemm
Más cosas
Resumen
Array of Structures vs. Structure of Arrays.
AoS vs. SoA
Fragmento típico de computación con (x, y , z, w ),
independientemente de su layout en memoria:
d = x*x + y*y + z*z + w*w;
Orden de lectura de un programa
Secuencial : x, y , z, w , . . . x, y , z, w
Paralelo SIMD : x, . . . , x, y , . . . , y , z, . . . , z, w , . . . , w
Introducción
distances
vectorsum
sgemm
Más cosas
Resumen
Implementaciones
AoS
__global__ void dist_aos(float4 *vectors, float *dists) {
uint tid = threadIdx.x;
uint bid = blockIdx.y*(gridDim.x) + blockIdx.x;
uint i = bid*(blockDim.x) + tid;
float d = sqrtf(vectors[i].x*vectors[i].x +
vectors[i].y*vectors[i].y +
vectors[i].z*vectors[i].z +
vectors[i].w*vectors[i].w);
dists[i] = d;
}
SoA
__global__ void dist_soa(float *xs, float *ys, float *zs, float *ws,
float *dists) {
uint tid = threadIdx.x;
uint bid = blockIdx.y*(gridDim.x) + blockIdx.x;
uint i = bid*(blockDim.x) + tid;
float d = sqrtf(xs[i]*xs[i] +
ys[i]*ys[i] +
zs[i]*zs[i] +
ws[i]*ws[i]);
dists[i] = d;
}
Introducción
distances
vectorsum
sgemm
Más cosas
Resumen
Medición
Para 256M vectores que consumen el 85 % de los 6GB de la
C2070, medimos.
$ make
nvcc -O3 -arch=sm_20 -I/opt/cudasdk/C/common/inc -o distances.o -c distances.cu
nvcc -O3 -arch=sm_20 -I/opt/cudasdk/C/common/inc -o distances distances.o
$ ./distances
$ grep dist cuda_profile_0.log
method=[ _Z8dist_aosP6float4Pf ] gputime=[ 59988.449 ] cputime=[ 6.000 ] occupancy=[ 1.000 ]
method=[ _Z8dist_soaPfS_S_S_S_ ] gputime=[ 44686.625 ] cputime=[ 4.000 ] occupancy=[ 1.000 ]
• Como era de esperar SoA es mejor que AoS para GPU: 1.34x.
• BW para AoS: ((4×4×228)/0,044686)/230 = 89,51GBps.
¡Bastante bien!
Introducción
distances
vectorsum
sgemm
Más cosas
Resumen
Introducción
Ejemplo: distancias de un vector 4d
Ejemplo: suma de vector de a bloques
Ejemplo: explorando tamaños de bloque
Más cosas
Resumen
Introducción
distances
vectorsum
sgemm
Más cosas
Resumen
Repasando como funciona vectorsum
Entrada: float a[N].
Salida: float partial sum[N/BLOCK]
float *a
1.0
2.0
3.0
4.0
5.0
6.0
7.0
8.0
9.0
10.0
11.0
12.0
float *partial sum
10.0
26.0
42.0
Usar la memoria compartida para acumulación atómica rápida.
Introducción
distances
vectorsum
sgemm
Más cosas
Resumen
Dos versiones para comparar velocidad
__global__ void vectorsum(float *a, float *partial_sum) {
uint tid = threadIdx.x;
uint bid = blockIdx.x;
uint i = bid*blockDim.x + tid;
// atomically accumulate into global
atomicAdd(&partial_sum[bid], a[i]);
}
__global__ void vectorsum_shared(float *a, float *partial_sum) {
uint tid = threadIdx.x;
uint bid = blockIdx.x;
uint i = bid*blockDim.x + tid;
__shared__ float block_sum;
// first thread in block resets shared
if (tid==0)
block_sum = 0.0f;
__syncthreads();
// atomically accumulate into shared
atomicAdd(&block_sum, a[i]);
__syncthreads();
// first thread in block copy to partial sum
if (tid==0)
partial_sum[bid] = block_sum;
}
Introducción
distances
vectorsum
sgemm
Más cosas
Resumen
Mediciones
$ ./vectorsum && grep vector cuda_profile_0.log
vector_sum maxdiff: 741.610352
vectorsum_shared maxdiff: 0.001038
method=[ _Z9vectorsumPfS_ ] gputime=[ 268549.094 ] cputime=[ 4.000 ] occupancy=[ 0.667 ]
method=[ _Z16vectorsum_sharedPfS_ ] gputime=[ 181925.469 ] cputime=[ 4.000 ] occupancy=[ 0.667 ]
$ ./vectorsum && grep vector cuda_profile_0.log
vector_sum maxdiff: 741.610840
vectorsum_shared maxdiff: 0.000732
method=[ _Z9vectorsumPfS_ ] gputime=[ 268814.625 ] cputime=[ 4.000 ] occupancy=[ 0.667 ]
method=[ _Z16vectorsum_sharedPfS_ ] gputime=[ 181907.844 ] cputime=[ 4.000 ] occupancy=[ 0.667 ]
$ ./vectorsum && grep vector cuda_profile_0.log
vector_sum maxdiff: 741.610718
vectorsum_shared maxdiff: 0.001099
method=[ _Z9vectorsumPfS_ ] gputime=[ 269452.688 ] cputime=[ 5.000 ] occupancy=[ 0.667 ]
method=[ _Z16vectorsum_sharedPfS_ ] gputime=[ 181913.734 ] cputime=[ 4.000 ] occupancy=[ 0.667 ]
Notar
• Usar la shared da un 1.47x.
• 225/(0.268×1e9) = 0.1252GFLOPS, puajj!
• 3×4×225/(0.268×230) = 1.399GBps, puajj!!
• ¿Hay una race condition!
En realidad es la no-conmutatividad de la suma de fp32.
Solucionarlo sería sumar secuencialmente.
Introducción
distances
vectorsum
sgemm
Más cosas
Resumen
Un detalle sobre la inicialización
__global__ void set(float *a, float *partial_sum) {
uint tid = threadIdx.x;
uint bid = blockIdx.x;
uint i = bid*blockDim.x + tid;
a[i] = sinf((float)bid*tid);
// first blockDim threads initialize
if (i<blockDim.x)
partial_sum[i] = 0.0f;
}
• Elijo uso los primeros N/BLOCK hilos para inicializar partial sum,
• Los primeros N/BLOCK/32 warps trabajan bien con la memoria.
• La otra opción
if (tid==0) partial sum[bid] = 0.0f;
desperdicia BW.
Introducción
distances
vectorsum
sgemm
Más cosas
Resumen
Introducción
Ejemplo: distancias de un vector 4d
Ejemplo: suma de vector de a bloques
Ejemplo: explorando tamaños de bloque
Más cosas
Resumen
Introducción
distances
vectorsum
sgemm
Más cosas
Resumen
Explorando tamaños de bloque en sgemm
sgemm simple
• El mismo que mostramos ayer, la versión más sencilla.
• Grilla y bloque bidimensional para L×L = 1024×1024.
• Barremos el rango [1, 32]×[1, 32] para tamaño de bloque.
• Medimos el tiempo en µs.
• Tesla C2070, driver 295.41, nvcc-4.2.
Introducción
distances
vectorsum
sgemm
Más cosas
Resumen
Limitaciones al paralelismo
Limitaciones al paralelismo: registros por hilo
Compilamos con --ptxas-options=-v
$ nvcc sgemm.cu -O3 -arch=sm_20 --ptxas-options=-v -I/opt/cudasdk/C/common/inc -o sgemm.o
ptxas info
ptxas info
: Compiling entry function ’_Z5sgemmjPfS_S_’ for ’sm_20’
: Function properties for _Z5sgemmjPfS_S_
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info
: Used 16 registers, 64 bytes cmem[0], 24 bytes cmem[2]
O bien hacemos profiling regperthread
$ export CUDA_PROFILE=1
$ export CUDA_PROFILE_CONFIG=profile_debug_kernel
$ cat profile_debug_kernel
gridsize
threadblocksize
regperthread
$ ./sgemm 1024 16 16
max_diff: 0.000092
$ cat cuda_profile_0.log
...
method=[ _Z5sgemmjPfS_S_ ] gputime=[ 34407.359 ] cputime=[ 4.000 ]
gridsize=[ 64, 64 ] threadblocksize=[ 16, 16, 1 ] regperthread=[ 16 ] occupancy=[ 1.000 ]
...
shared , solo 16 registros y 16×1536 ≤ 32768.
No usa
No pueden saturar al SM.
Introducción
distances
vectorsum
sgemm
Más cosas
Resumen
Heatmap de velocidad
Introducción
distances
vectorsum
sgemm
Más cosas
Resumen
Notar
• Asimétrico.
• Franjas claras en bx ∈ {16, 24, 32}.
• El mejor tiempo está en 32×5, con 33.264ms.
• Muy por debajo de los ∼ 1,1TFLOPS fp32 de la placa.
• Pero el código trivial es casi comparable con una versión
Para 2×10243 FLOP, implica 60GFLOPS.
• El mínimo en bx=32 no es casualidad:
SSE+OpenMP en CPU modernas de 8 cores.
32×4 ⇒ warp lee una línea de cache completa.
Introducción
distances
vectorsum
sgemm
Más cosas
Resumen
Occupancy no es performance
Recordemos que para GF100:
warpsPerSM≤48, blocksPerSM≤8, granularidad por SM: bloque.
thsPerBlock
thsPerSM blksPerSM occupancy
gputime
64
256
1024
128
128
64
64
160
512
1536
1024
1024
1024
512
512
1280
8
6
1
8
8
8
8
8
512
1536 =0.33
1536
1536 =1.0
1024
1536 =0.66
1024
1536 =0.66
1024
1536 =0.66
512
1536 =0.33
512
1536 =0.33
1280
1536 =0.83
53 ms
38 ms
38 ms
45 ms
35 ms
182 ms
49 ms
33 ms
block
8×8
16×16
32×32
8×16
16×8
2×32
32×2
32×5
Notar
• occupancy no es proporcional a la performance.
• Lo mejor es tener muchos bloques por SM y con buen
memory layout.
Introducción
distances
vectorsum
sgemm
Más cosas
Resumen
Introducción
Ejemplo: distancias de un vector 4d
Ejemplo: suma de vector de a bloques
Ejemplo: explorando tamaños de bloque
Más cosas
Resumen
Introducción
distances
vectorsum
sgemm
Más cosas
Resumen
Lo que me hubiera gustado dar
Comentemos un poco.
• Intensidad computacional.
• Branch divergence, predicate computation.
• #pragma unroll.
• Deshabilitar las cache.
• Usar templates de C++.
• Las operaciones fp e int son gratis si el problema es
memory-bound.
Introducción
distances
vectorsum
sgemm
Más cosas
Resumen
Introducción
Ejemplo: distancias de un vector 4d
Ejemplo: suma de vector de a bloques
Ejemplo: explorando tamaños de bloque
Más cosas
Resumen
Introducción
distances
vectorsum
sgemm
Más cosas
Resumen
Resumen
• Vimos algunas estrategias de mejor de desempeño.
• Casi siempre, si son acertadas, dan grandes ganancias.
• Es usual ir apila
Comentarios de: Optimizaciones (0)
No hay comentarios