CUDA 3

Presenter Notes

Resumen

  • Reduce.

Nicolás Wolovick, 20200603.

Presenter Notes

Reducción

Paralelizar s = \sum a[i]

1 sum = 0.0f;
2 for(size_t i=0; i<N; ++i)
3     s += a[i];

Suma paralela en O(log(N)).

Presenter Notes

Comunicación entre hilos

Sumar todos los elementos de un arreglo

  • Problemas de concurrencia x=x+1.
  • Atomics en global.
  • Atomics en shared.
  • Último nivel de localidad: warp shuffling.

Un artículo viejo, pero interesante

Presenter Notes

Jerarquía de Paralelismo

CUDA hierarchy of threads, blocks, grids

Presenter Notes

Unidades de Cooperación

Warp

  • Sincronización de warps __syncwarp().
  • Comunicación de variables privadas.
  • Comunicación y sincronización _shfl_*_sync.

Bloque

  • Sincronización de barrera __syncthreads().
  • Memoria compartida local shared.
  • Instrucciones atómicas sobre la shared: atomicAdd, CAS.

Grilla

  • Sincronización fork-join por lanzamiento de kernels.
  • Memoria compartida global.
  • Instrucciones atómicas sobre la global: atomicAdd, CAS.

Presenter Notes

Convención

Identificadores de jerarquía de paralelismo

 1 #include "helper_cuda.h"
 2 
 3 uint lane = tid & CUDA_WARP_MASK; // lane dentro del warp
 4 
 5 uint tid = threadIdx.x; // hilo dentro del bloque.
 6 uint warp = tid / CUDA_WARP_SIZE;  // warp dentro del bloque
 7 
 8 uint gtid = threadIdx.x + blockDim.x*blockIdx.x;  // Identificador global de hilo
 9 uint gwarp = gtid / CUDA_WARP_SIZE;  // Identificador global de warp
10 uint bid = blockIdx.x;  // Identificador de bloque

Presenter Notes

Mejoras escandalosas

Comparación de las diferentes versiones en una GTX Titan X Maxwell.

reduce1.cu

  • 256ms, pero incorrecto.
  • 790ms, con atomicAdd, y es correcto.

reduce2.cu

  • 1200ms, usamos reducción por bloque con atomic en la shared y luego acumulamos 1 por bloque en la global.
  • ¡Empeoró!

reduce3.cu

  • 1200ms idem anterior, solo que en vez de acumular en la global, dejamos un valor por bloque en un arreglo y acumula la CPU.

Presenter Notes

Mejoras escandalosas

reduce3_1.cu

  • 12ms, idem anterior, pero hay 32 valores de acumulación sobre la shared para evitar que los 32 lanes usen atomics sobre el mismo.
  • ¡Funcionó!

reduce4.cu

  • 11.6ms, en vez de lo anterior, warp shuffle para acumular por warp en los registros como suma paralela O(log(N)), atomic a la shared por warp y luego atomic a la global por bloque.
  • ¡Funcionó!

Locality=Performance

Presenter Notes

Comparación con CPU

 1 #define N (1<<28)
 2 float a[N];
 3 
 4 int main(void) {
 5     float s = 0.0f;
 6     #pragma omp parallel for simd reduction (+:s)
 7     for(size_t i=0; i<N; ++i)
 8         s += a[i];
 9     return (int)s;
10 }

Hacemos

1 $ gcc -fopenmp -O3 -ffast-math -mcmodel=medium reduce.c && perf stat -r 16 ./a.out
  • 2*E52620v3: 18ms
  • GTX Titan X Maxwell: 11ms

1.63x más rápida la GPU que las dos CPU.

Hay muchísimo margen para mejorar en la GPU: Optimizing Parallel Reduction in CUDA

Presenter Notes

ToDo

  • Actualizar todo a Pascal y Turing, además de los 2xE5-2680v4.
  • Mostrar como global atomics es malo
    • pero shared atomics + global atomics con shared bank conflic es GUACALA.
    • solución, software clutter con 32 sumas distintas por bloque.
    • finalmente como shuf_sync es la posta.
    • Comparar con CUB.
    • Comparar con approach butterfly naive como los primeros del trabajo de Mark Harris.
    • Pasar todo a enteros o las sumas no dan!

Presenter Notes

Bibliografía

Presenter Notes

La clase que viene

  • Herramientas.

Presenter Notes