Newer
Older
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
#include "CudaKernels.cuh"
__global__ void cudaSumArray(CudaImageContainer arrayIn, double* arrayOut, size_t n)
{
//This algorithm was used from a this website:
// http://developer.download.nvidia.com/compute/cuda/1.1-Beta/x86_website/projects/reduction/doc/reduction.pdf
// accessed 4/28/2013
extern __shared__ double sdata[];
size_t tid = threadIdx.x;
size_t i = blockIdx.x*blockDim.x*2 + tid;
size_t gridSize = blockDim.x*2*gridDim.x;
sdata[tid] = 0;
while (i<n)
{
sdata[tid] = arrayIn[i];
if (i+blockDim.x<n)
sdata[tid] += arrayIn[i+blockDim.x];
i += gridSize;
}
__syncthreads();
if (blockDim.x >= 2048)
{
if (tid < 1024)
sdata[tid] += sdata[tid + 1024];
__syncthreads();
}
if (blockDim.x >= 1024)
{
if (tid < 512)
sdata[tid] += sdata[tid + 512];
__syncthreads();
}
if (blockDim.x >= 512)
{
if (tid < 256)
sdata[tid] += sdata[tid + 256];
__syncthreads();
}
if (blockDim.x >= 256) {
if (tid < 128)
sdata[tid] += sdata[tid + 128];
__syncthreads();
}
if (blockDim.x >= 128)
{
if (tid < 64)
sdata[tid] += sdata[tid + 64];
__syncthreads();
}
if (tid < 32) {
if (blockDim.x >= 64)
{
sdata[tid] += sdata[tid + 32];
__syncthreads();
}
if (blockDim.x >= 32)
{
sdata[tid] += sdata[tid + 16];
__syncthreads();
}
if (blockDim.x >= 16)
{
sdata[tid] += sdata[tid + 8];
__syncthreads();
}
if (blockDim.x >= 8)
{
sdata[tid] += sdata[tid + 4];
__syncthreads();
}
if (blockDim.x >= 4)
{
sdata[tid] += sdata[tid + 2];
__syncthreads();
}
if (blockDim.x >= 2)
{
sdata[tid] += sdata[tid + 1];
__syncthreads();
}
}
if (tid==0)
arrayOut[blockIdx.x] = sdata[0];
}