#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <cstdlib>
#include <cstdio>
#include <ctime>

#include <time.h>
#include <math.h>
#include <string.h>
#include <stdio.h>


__global__ void mnozenieMacierzyKernel(float *a,float *macierzWynikowa, int rozmiarBloku, int grid) //rozmiar bloku na karcie graficznej jaki przekazuje
{
	float wartosc = 0;
	int w = grid * blockIdx.y + threadIdx.y;
	int k = grid * blockIdx.x + threadIdx.x;

	for(int i = 0; i < rozmiarBloku; i++)
	{
		wartosc += a[w * rozmiarBloku + i] * a[i * rozmiarBloku + k];
	}
	macierzWynikowa[w * rozmiarBloku + k] = wartosc;
}

float* generuj(int w, int k)
{
	float* temp;
	temp = (float*) malloc(sizeof(float)*w*k);

	for(int i = 0; i<w*k; i++)
	{
        *(temp+i)=rand()%99;
    }

	return temp;
}

void mnozenieMacierzy(float* a, float* macierzWynikowa, int w, int k)
{
	for(int i = 0; i < w; i++)
	{
		for(int j = 0; j < k; j++)
		{
			for(int l = 0; l < k; l++)
			{
				macierzWynikowa[i * k + j] += a[i * k + l]* a[l * k  + j];
			}
		}
	}
}

void zapis(char *filename,float *macierz, float* macierzWynikowa, int w, int k, float execTime, float czasGPU)
{
	FILE *fp;
	if ((fp=fopen(filename, "w"))==NULL)
	{
		printf ("Nie mogę otworzyć pliku do zapisu!\n");
		exit(1);
	}

	fprintf (fp, "czas dzialania: %lf\n", execTime);
	fprintf (fp, "czas dzialania GPU: %lf\n", czasGPU);
	fprintf (fp, "wierze: %d\n", w);
	fprintf (fp, "kolumny: %d\n", k);

	fprintf (fp, "macierz:\n");
	for(int i = 0; i < w; i++)
	{
		for(int j = 0; j < k; j++)
		{
			fprintf (fp, "%lf\t", macierz[i*k+j]);
		}
		fprintf (fp, "\n");
	}

	fprintf (fp, "\n\nmacierz wynikowa:\n");
	for(int i = 0; i < w; i++)
	{
		for(int j = 0; j < k; j++)
		{
			fprintf (fp, "%lf\t", macierzWynikowa[i*k+j]);
		}
		fprintf (fp, "\n");
	}
	fclose (fp);
}


	void mnozenieMacierzyCuda(float* a, float* macierzWynikowa,int w, int k, int grid, float& execTime, float& memIncTime)//
    {
		float* devA = 0;
		float* devMacierzWynikowa = 0;
		cudaError_t cudaStatus;
		cudaEvent_t memStart, memStop;
		cudaEventCreate(&memStart);
		cudaEventCreate(&memStop);

		cudaStatus = cudaSetDevice(0);
		if(cudaStatus != cudaSuccess)
		{
			printf("cudaSetDevice: %s", cudaGetErrorString(cudaStatus));
			exit(1);
		}

		cudaEventRecord(memStart);

		cudaStatus = cudaMalloc((void**)&devMacierzWynikowa, sizeof(macierzWynikowa));
		if(cudaStatus != cudaSuccess)
		{
			printf("cudaMalloc: %s",cudaGetErrorString(cudaStatus));
			exit(1);
		}

		cudaStatus = cudaMalloc((void**)&devA, sizeof(a));
		if(cudaStatus != cudaSuccess)
		{
			printf("cudaMalloc: %s",cudaGetErrorString(cudaStatus));
			exit(1);
		}


		cudaStatus = cudaMemcpy(devA, a, sizeof(a), cudaMemcpyHostToDevice);
		if(cudaStatus != cudaSuccess)
		{
			printf("cudaMemcpy: %s",cudaGetErrorString(cudaStatus));
			exit(1);
		}


		int gridX = w / grid;
		int gridY = k / grid;

		dim3 dimGrid(gridX, gridY, 1);
		dim3 dimBlock(grid, grid,1);
		//dim3 dimBlock(grid, grid,1);
		//dim3 dimGrid( (w + dimBlock.x - 1) / dimBlock.x,(k + dimBlock.y - 1) / dimBlock.y );


		cudaEvent_t start, stop;
		cudaEventCreate(&start);
		cudaEventCreate(&stop);

		cudaEventRecord(start);
		mnozenieMacierzyKernel<<<dimGrid, dimBlock>>>(devA, devMacierzWynikowa, k, grid);
		cudaEventRecord(stop);

		cudaEventSynchronize(stop);

		cudaEventElapsedTime(&execTime, start, stop);

		cudaEventDestroy(start);
		cudaEventDestroy(stop);

		// Check for any errors launching the kernel
		cudaStatus = cudaGetLastError();
		if(cudaStatus != cudaSuccess)
		{
			printf("An error has occured during matrix multiplication: : %s",cudaGetErrorString(cudaStatus));
			exit(1);
		}

		// cudaDeviceSynchronize waits for the kernel to finish, and returns
		// any errors encountered during the launch.
		cudaStatus = cudaDeviceSynchronize();
		if(cudaStatus != cudaSuccess)
		{
			printf("cudaDeviceSynchronize: %s",cudaGetErrorString(cudaStatus));
			exit(1);
		}

		// Copy output vector from GPU buffer to host memory.
		cudaStatus = cudaMemcpy(macierzWynikowa, devMacierzWynikowa, sizeof(macierzWynikowa), cudaMemcpyDeviceToHost);
		if(cudaStatus != cudaSuccess)
		{
			printf("cudaMemcpy: %s",cudaGetErrorString(cudaStatus));
			exit(1);
		}


        cudaFree(devA);
        cudaFree(devMacierzWynikowa);

        cudaEventRecord(memStop);

        cudaEventSynchronize(memStop);

        cudaEventElapsedTime(&memIncTime, memStart, memStop);

        cudaEventDestroy(memStart);
        cudaEventDestroy(memStop);
    }


int main()
{
    srand( time( NULL ) );

	float* macierze [10];
	float* macierzWynikowa[10];
	float* macierzWynikowa2[10];
	time_t startCPU;
	time_t endCPU;
	float czasGPU;
	char filename[100];

	cudaDeviceProp prop;
	cudaGetDeviceProperties(&prop, 0);
	double temp = sqrt((double)prop.maxThreadsPerBlock);
	int grid = (int) temp;

	for(int i = 1; i < 11; i++)
	{
		macierze[i] = generuj(i * 10, i * 10);
		macierzWynikowa[i] = (float*) calloc((i * 10 * i * 10),sizeof(float));
		macierzWynikowa2[i] = (float*) calloc((i * grid * i * grid),sizeof(float));

		startCPU = clock();
		mnozenieMacierzy(macierze[i], macierzWynikowa[i],i * 10, i * 10);
		endCPU = clock();

		float execTime = static_cast<float>((endCPU - startCPU) / float(CLOCKS_PER_SEC) * 1000);

		sprintf (filename, "wyniki_%d.txt", i);
		zapis(filename, macierze[i], macierzWynikowa[i],i * 10, i * 10, execTime,NULL);

		//free(macierzWynikowa[i]);
		//macierzWynikowa[i] = (float*) calloc((i * 10 * i * 10),sizeof(float));
		//mnozenieMacierzyCuda(macierze[i], macierzWynikowa[i],i * 10, i * 10, grid, execTime,czasGPU);
		//sprintf (filename, "wyniki_%dGPU.txt", i);
		//zapis(filename, macierze[i], macierzWynikowa[i],i * 10, i * 10, execTime,czasGPU);

		mnozenieMacierzyCuda(macierze[i], macierzWynikowa2[i],i * 10, i * 10, 48, execTime,czasGPU);
				sprintf (filename, "wyniki_%dGPU.txt", i);
		zapis(filename, macierze[i], macierzWynikowa2[i],i * 10, i * 10, execTime,czasGPU);
	}

    getchar();
    return 0;
}