#include <stdio.h>
#include <cuda.h>
#define N 4096
#define THREADS_PER_BLOCK 512

// device kernel
__global__ void vectorAddKernel(int *d_a, int *d_b, int *d_c){
	int index = threadIdx.x + blockIdx.x * blockDim.x;
	d_c[index] = d_a[index] + d_b[index];
}

int main(void){

//grid and block sizes
dim3 grid(N/THREADS_PER_BLOCK, 1, 1);
dim3 block(THREADS_PER_BLOCK, 1, 1); 

// host copies
int a_h[N];
int b_h[N];
int c_h[N];

// device copies
int *a_d;
int *b_d;
int *c_d;

// generate data on host
for(int i=0; i<N; i++){
	a_h[i] = i+1;
	b_h[i] = i+1;
	}

// alloc size
cudaMalloc((void**)&a_d, N*sizeof(int));
cudaMalloc((void**)&b_d, N*sizeof(int));
cudaMalloc((void**)&c_d, N*sizeof(int));

cudaMemcpy(a_d, a_h, N*sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(b_d, b_h, N*sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(c_d, c_h, N*sizeof(int), cudaMemcpyHostToDevice);

cudaEvent_t start;
cudaEvent_t stop;
float elapsedTime; 

cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);

// launch kernel on GPU
vectorAddKernel<<<grid, block>>>(a_d,b_d,c_d);

cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&elapsedTime, start, stop);

cudaMemcpy(c_h, c_d, N*sizeof(int), cudaMemcpyDeviceToHost);
for(int i=0; i<N; i++){
	printf("%i+%i = %i\n",a_h[i], b_h[i], c_h[i]);
	}

printf("Time to calculate results: %f ms.\n", elapsedTime);

cudaFree(a_h);
cudaFree(b_h);
cudaFree(c_h);
cudaEventDestroy(start);
cudaEventDestroy(stop);
return 0;
}