Nicolás Wolovick, 20180523.
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 #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 gtid = blockIdx.x*blockDim.x + threadIdx.x;
10 d[gtid] = a[gtid]*b[gtid]+c[gtid];
11 }
12
13 int main(void) {
14 ma4<<<N/BLOCK_SIZE, BLOCK_SIZE>>>();
15 cudaDeviceSynchronize();
16 return 0;
17 }
1 #include <stddef.h>
2 #define N (1<<28)
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 }
2.86x más rápida la GPU que las dos CPU.
Problema totalmente memory-bound.
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)).
x=x+1
.Un artículo viejo, pero interesante
(via ballots & warp shuffling)
atomicAdd
, CAS
.atomicAdd
, CAS
. 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
Comparación de las diferentes versiones en una GTX Titan X Maxwell.
reduce1.cu
atomicAdd
, y es correcto.reduce2.cu
shared
y luego acumulamos 1 por bloque en la global.reduce3.cu
reduce3_1.cu
reduce4.cu
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
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
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 |