#include <stdio.h>
#include <assert.h>
#include <stdlib.h>
#include "timer.h"
#include <iostream>
#include <string>

using namespace std;

static void handleError(cudaError_t err)
{
    if (err != cudaSuccess)
	{
        cout << cudaGetErrorString(err) << "\n";
        exit(1);
    }
}

#define BLOCK_SIZE 512

template<int VECTOR_SIZE, int DUMMY, typename T>
__global__
void reduce(T* vector, T* matrix)
{
	int id = threadIdx.x + blockIdx.x * blockDim.x;

	if (VECTOR_SIZE != BLOCK_SIZE)
	vector[id] = 0;
	matrix[id] = 1;

	for (int i = 0; i < VECTOR_SIZE; i++)
	{
		vector[id] += matrix[id];
	}
}


string convertToComaNotation(double value)
{
	string initial = to_string(value);
	for (auto& letter : initial)
	{
		letter = letter != '.' ? letter : ',';
	}
	return initial;
}

template<int VECTOR_SIZE,int STREAMS>
void singleTest(int count)
{
	  const int num_streams = STREAMS;

    cudaStream_t streams[num_streams];

    for (int i = 0; i < num_streams; i++) {
        cudaStreamCreate(&streams[i]);
	}

    size_t vector_size = sizeof(int) * VECTOR_SIZE;

    int *h_res = (int *) malloc(sizeof(int) * VECTOR_SIZE);

    int *d_vector = 0;
	int *d_matrix = 0;

    handleError(cudaMalloc((void **) &d_vector, vector_size));
    handleError(cudaMalloc((void **) &d_matrix, vector_size));

    dim3 threads(BLOCK_SIZE);
    dim3 grid(VECTOR_SIZE/BLOCK_SIZE);

	for (int id = 0; id < count; id++)
	{
	reduce<VECTOR_SIZE,BLOCK_SIZE,int><<< grid, threads, 0, streams[id%num_streams] >>>(d_vector, d_matrix);
	}

    handleError(cudaPeekAtLastError());

	cudaDeviceSynchronize();

    handleError(cudaMemcpy(h_res, d_vector, vector_size, cudaMemcpyDeviceToHost));

	handleError(cudaFree(d_vector));
	handleError(cudaFree(d_matrix));
}

int main()
{
	Timer timer;
	singleTest<512 * 8, 1>(100);
	cout <<convertToComaNotation(timer.elapsed())<<"\n";

	timer.reset();
	singleTest<512 * 8, 4>(100);
	cout <<convertToComaNotation(timer.elapsed())<<"\n";
}