Nicolás Wolovick, 20160523.
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 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 }
(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 980.
cuda_reduce1.cu
cuda_reduce2.cu
, cuda_reduce3.cu
cuda_reduce4.cu
Una vez mas ...
{S,D}GEMM
, block size, performance.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 |