CUDA 2

Presenter Notes

Resumen

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

Nicolás Wolovick, 20200601.

Presenter Notes

Paralelismo que da vergüenza

embarrassingly parallel

Toda operación que es completamente independiente de otra.

Ejemplo típico

Establecer algún orden es restringir artificialmente el flujo de control.
Las operaciones pueden ocurrir en cualquier orden, inclusive en simultaneidad.

Presenter Notes

Ejemplo: map

Presenter Notes

Map

 1 #define N (1L<<28)
 2 #define BLOCK_SIZE 128
 3 
 4 __device__ float a[N];
 5 
 6 __global__ void map(void) {
 7     unsigned int gtid = blockIdx.x*blockDim.x + threadIdx.x;
 8     a[gtid] *= 2.0f;
 9 }
10 
11 int main(void)
12 {
13     map<<<N/BLOCK_SIZE, BLOCK_SIZE>>>();
14     getLastCudaError("map() kernel failed");
15     cudaDeviceSynchronize();
16 
17     return 0;
18 }

. Paralelismo de grano fino, idem SIMD, un lane un dato.
. Mapeo identidad: gtid 0 -> mapea a[0], etc.
. Va rapidísimo! 180 GiB/s en una GTX 1070.

Presenter Notes

Map

Cruzamos un poco los hilos (por Mateo de Mayo)
hilo 0 -> a[0], hilo 1 -> a[1024], hilo 2 -> a[2048], ...

 1 #define GRID_SIZE 1024
 2 #define BLOCK_SIZE 1024
 3 
 4 __device__ float a[GRID_SIZE*BLOCK_SIZE];
 5 
 6 __global__ void map(void) {
 7     unsigned int gtid = threadIdx.x*blockDim.x + blockIdx.x;
 8     a[gtid] *= 2.0f;
 9 }
10 
11 int main(void)
12 {
13     map<<<GRID_SIZE, BLOCK_SIZE>>>();
14     getLastCudaError("map() kernel failed");
15     cudaDeviceSynchronize();
16 
17     return 0;
18 }

. Versión derecha con 1024x1024 elementos: 40µs.
. Versión cruzada idem elementos: 243µs.
. El gather y el mal uso de ancho de banda implica un speeddown de 6x.

Presenter Notes

Multiply and Add 4 (MA4)

Presenter Notes

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 24 SM (shaders) de una Titan X Maxwell.
    • Con el SMT-64 sobre SIMD-1024 sobre vectores tenemos 24x64x(1024/32) = 49152 hilos concurrentes reales.
    • Cuanto más hilos más oculta la latencia, mayor througput (ma non troppo).
  • Estrategia: un hilo por dato ¡Impensable en Multicore! ¡Lo que hacíamos con SIMD!

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

Presenter Notes

MA4 en CUDA

 1 #define N (1L<<29)
 2 #define BLOCK_SIZE 128
 3 
 4 __global__ void set(float *a, float *b, float *c) {
 5     unsigned int gtid = blockIdx.x*blockDim.x + threadIdx.x;
 6     a[gtid] = (float)blockIdx.x; b[gtid] = (float)threadIdx.x;
 7     c[gtid] = (float)threadIdx.x+blockIdx.x; d[gtid] = (float)threadIdx.x*blockIdx.x;
 8 }
 9 __global__ void ma4(void) {
10     unsigned int gtid = blockIdx.x*blockDim.x + threadIdx.x;
11     d[gtid] = a[gtid]*b[gtid]+c[gtid];
12 }
13 
14 
15 int main(void) {
16     float *a, *b, *c, *d; a = b = c = d = NULL;
17     cudaMallocManaged(&a, N*sizeof(float));
18     cudaMallocManaged(&b, N*sizeof(float));
19     cudaMallocManaged(&c, N*sizeof(float));
20     cudaMallocManaged(&d, N*sizeof(float));
21     set<<<N/BLOCK_SIZE, BLOCK_SIZE>>>(a,b,c);
22     ma4<<<N/BLOCK_SIZE, BLOCK_SIZE>>>(a,b,c,d);
23     cudaDeviceSynchronize();
24 }

Presenter Notes

Detalles versión CUDA

Usa paged memory CPU<->GPU con cudaMallocManaged

  • En set<<<,>>>() se pide la memoria en GPU.
  • ma4<<<,>>>() opera sobre GPU memory.
  • Con N = (1L<<29) Cada arreglo ocupa 2 GiB.

Presenter Notes

Versión CPU con ILD, DLP, TLP

 1 #include <stddef.h>
 2 #define N (1L<<29)
 3 
 4 float a[N], b[N], c[N], d[N];
 5 
 6 int main(void)
 7 {
 8     #pragma omp parallel for simd
 9     for(size_t i=0; i<N; ++i) {
10         d[i] = a[i]*b[i]+c[i];
11     }
12 
13     return (int)a[N/2]+b[N/3]+c[N/4]+d[N/5];
14 }

Compilando con todos los chiches.

1 $ gcc-10 -O3 -ffast-math -fopenmp ma4_simd_omp.c -mcmodel=medium

Presenter Notes

Comparación de desempeño

La aplicación es mem-bw limited.

BW pico
2xE5-2680v4: 100GB/s, GTX1070: 256GB/s, RTX2080ti: 616 GB/s.

1 //N=2^k   2*E5-2680v4  GTX1070(8GB)  RTX2080ti(11GB)
2 //   28         35 ms         20 ms              8ms
3 //   29         62 ms       3290 ms             15ms
4 //   30        113 ms       5685 ms           6330ms

Normalizado a BW máximo con CPU=1, más es mejor.

1 //N=2^k   2*E5-2680v4  GTX1070(8GB)  RTX2080ti(11GB)
2 //   28             1          0.68             0.71
3 //   29             1      --swap--             0.67
4 //   30             1      --swap--         --swap--
  • Cuando se acaba la GDDR5(x) de las GPUs swappea a CPU, intolerablemente lento.
  • Las GPUs no son tan rápidas como deberían dado su BW.
  • Con 2 x E5-2680: $2500, GTX1070: $400, RTX2080ti: $1000, la relación GB/s / USD es muy buena en GPU.

Presenter Notes

Bibliografía

Presenter Notes

La clase que viene

  • Ejemplo paralelismo no-trivial: reduce.

Presenter Notes