-
Notifications
You must be signed in to change notification settings - Fork 0
/
3.cu
63 lines (51 loc) · 1.27 KB
/
3.cu
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
#include <iostream>
__global__ void vectorAdd(int *a, int *b, int *c, int n){
int i = blockIdx.x*blockDim.x+threadIdx.x;
if(i<n)
for(int j=0;j<100;j++)
c[i] = a[i] + b[i];
}
int main(void){
int * a, * b;
int * r1, * r2, *r3;
int * temp;
const int n = 1<<24;
const int n_s = 3;
cudaStream_t streams[n_s];
for(int i=0;i<n_s;i++)
cudaStreamCreate(&streams[i]);
cudaMallocManaged(&a, n*sizeof(int));
cudaMallocManaged(&b, n*sizeof(int));
cudaMallocManaged(&r1, n*sizeof(int));
cudaMallocManaged(&r2, n*sizeof(int));
cudaMallocManaged(&r3, n*sizeof(int));
temp = new int[n*sizeof(int)];
for(int i=0;i<n;i++){
a[i] = 3;
b[i] = 5;
}
int blockSize = 256;
int numBlocks = n/256;
vectorAdd<<<numBlocks,blockSize,0,streams[0]>>>(a,a,r1,n);
vectorAdd<<<numBlocks, blockSize,0,streams[1]>>>(b,b,r2,n);
vectorAdd<<<numBlocks, blockSize,0,streams[2]>>>(a,b,r3,n);
cudaDeviceSynchronize();
temp[0] = r1[0];
for(int i=1;i<n;i++)
temp[i] = temp[i-1] + r1[i];
cudaDeviceSynchronize();
temp[0] = r2[0];
for(int i=1;i<n;i++)
temp[i] = temp[i-1] + r2[i];
cudaDeviceSynchronize();
temp[0] = r3[0];
for(int i=1;i<n;i++)
temp[i] = temp[i-1] + r3[i];
cudaFree(a);
cudaFree(b);
cudaFree(r1);
cudaFree(r2);
cudaFree(r3);
delete [] temp;
return 0;
}