CUDA 2

Presenter Notes

Resumen

  • Primer programa (paralelismo directo).
  • Segundo programa (no es tán fácil).

Nicolás Wolovick, 20160523.

Presenter Notes

Paralelismo que da vergüenza

Presenter Notes

Multiply and Add 4 (MA4)

Código CPU

1 #define N (1<<28)
2 float a[N], b[N], c[N], d[N];
3 
4 int main(void) {
5     for(unsigned int i=0; i<N; ++i)
6         d[i] = a[i]*b[i]+c[i];
7 
8     return 0;
9 }

Paralelización

  • CUDA está pensado para paralelización de datos de grano fino.
  • Como OpenMP, pero tenemos millones de hilos disponibles.
    • No solo los ~23000 cores (virtualización).
    • Cuanto más hilos más oculta la latencia, mayor througput (ma non troppo).
  • Estrategia: un hilo por dato (¡Impensable en CPU!).

Con millones de hilos no tiene sentido paralelización de tareas.

Presenter Notes

MA4 en CUDA

 1 #include <cuda.h>
 2 
 3 #define N (1<<28)
 4 #define BLOCK_SIZE 128
 5 
 6 __device__ float a[N], b[N], c[N], d[N];
 7 
 8 __global__ void ma4(void) {
 9     unsigned int tid = blockIdx.x*blockDim.x + threadIdx.x;
10     d[tid] = a[tid]*b[tid]+c[tid];
11 }
12 
13 int main(void) {
14     ma4<<<N/BLOCK_SIZE, BLOCK_SIZE>>>();
15     cudaDeviceSynchronize();
16     return 0;
17 }

Presenter Notes

Paralelismo no trivial

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.
  • Apuntar a artículo/slides de Mark Harris y que eso sea un ejercicio para ellos.

Presenter Notes

Jerarquía de Paralelismo

CUDA hierarchy of threads, blocks, grids

Presenter Notes

Unidades de Cooperación

Warp

  • Ejecución interlocked y ...
  • Comunicación de variables privadas.

(via ballots & warp shuffling)

Bloque

  • Sincronización de barrera.
  • 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 980.

  • reducción global ~800ms.
    cuda_reduce1.cu
  • reducción shared y luego global ~130ms
    cuda_reduce2.cu, cuda_reduce3.cu
  • reducción warp, luego shared y luego global ~11ms
    cuda_reduce4.cu

Una vez mas ...

Locality=Performance

Presenter Notes

Bibliografía

Presenter Notes

La clase que viene

  • {S,D}GEMM, block size, performance.

Presenter Notes