Nicolás Wolovick, 20200601.
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.
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.
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.
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 }
Con millones de hilos no tiene sentido paralelización de tareas.
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 }
Usa paged memory CPU<->GPU con cudaMallocManaged
set<<<,>>>()
se pide la memoria en GPU.ma4<<<,>>>()
opera sobre GPU memory.N = (1L<<29)
Cada arreglo ocupa 2 GiB. 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
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
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--
Table of Contents | t |
---|---|
Exposé | ESC |
Full screen slides | e |
Presenter View | p |
Source Files | s |
Slide Numbers | n |
Toggle screen blanking | b |
Show/hide slide context | c |
Notes | 2 |
Help | h |