Nicolás Wolovick, 20200603.
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
__syncwarp()
._shfl_*_sync
.__syncthreads()
.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 |