cuda-cpp-programming/recur.cu at main · MohammadDallash/cuda-cpp-programming · GitHub

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
#include <stdio.h>
#include <curand.h>
#include <cuda_runtime_api.h>
#include <time.h>
#include <sys/time.h>

__global__ void recurse(int level)
{
    if (level<0)
    {
        int x = 1000;
        for(int i = 0; i<1e6; i++)
        {
            if (i%2)x*=10;
            else x/=10;
        }
        return;
    }
    level--;
    if (threadIdx.x == 0) recurse<<<1,2>>>(level);
    else recurse<<<1,2>>>(level);

    cudaDeviceSynchronize(); //stop here in the current thread until the the above kernal is done

    __syncthreads(); //block level synchronization barrier from the parent kernal of this kernal
    return;
}

void serial_recurse(int level)
{
    if (level<0)
    {
        int x = 1000;
        for(int i = 0; i<1e6; i++)
        {
            if (i%2)x*=10;
            else x/=10;
        }
        return;
    }
    level--;
    serial_recurse(level);
    serial_recurse(level);
    return;

}


int main()
{
    struct timeval startwtime,endwtime;
    double seq_time;

    printf("Startin here\n");

    int level_1 = 10;
    gettimeofday(&startwtime,NULL);

    serial_recurse(level_1);

    gettimeofday(&endwtime,NULL);

    seq_time = (double)((endwtime.tv_usec - startwtime.tv_usec)/1.0e6+endwtime.tv_sec - startwtime.tv_sec);
    printf("serial time = %f \n", seq_time);

    int level=10;

    gettimeofday(&startwtime,NULL);

    recurse<<<1,2>>>(level);
    cudaDeviceSynchronize();

    gettimeofday(&endwtime,NULL);

    seq_time = (double)((endwtime.tv_usec - startwtime.tv_usec)/1.0e6+endwtime.tv_sec - startwtime.tv_sec);
    printf("parralel time = %f \n", seq_time);

    return 0;
}