////***** Martin Lukačovič *****////
////***** VYUŽITIE GRAFICKÝCH PROCESOROV PRE UNIVERZÁLNE VÝPOČTY V PRIEMYSELNÝCH SYSTÉMOCH *****////
////***** Bakalárska práca *****////
////***** 4.6.2014, VUT v Brne *****////
////***** Vytvorené na základe http://vasanthexperiments.wordpress.com/tag/gpu-matrix-multiplication/ *****////


#include < stdio.h >
#include < stdlib.h >
#include < time.h >
#include < ctime >
#include < iostream >
#include < windows.h > 
#include < CL/cl.h >

#define widthA 32
#define heightA 32

#define widthB heightA
#define heightB 32

#define widthC widthA
#define heightC heightB

#define MEM_SIZE (256)
#define MAX_SOURCE_SIZE (0x1000000)

using namespace std;

int main()
{
	cl_device_id zariadenie;
	cl_context context;
	cl_command_queue command_queue;
	cl_mem memobjA = NULL, memobjB = NULL, memobjC = NULL, rowA = NULL, colC = NULL;
	cl_program program;
	cl_kernel kernel;
	cl_platform_id platforma[2];
	cl_uint ret_num_devices=0, ret_num_platforms=0;
	cl_int ret;
	cl_event kernelevent, 
			 zapis1, zapis2, zapis3, zapis4,
			 citanie;
	
	cl_ulong time_start, time_end,
			 time_szapisu1, time_kzapisu1, time_szapisu2, time_kzapisu2,
			 time_szapisu3, time_kzapisu3, time_szapisu4, time_kzapisu4,
			 koniec_citania, start_citania;
			
	float total_cistokernel,
		  total_zapis1, total_zapis2, total_zapis3, total_zapis4, total_zapis,
		  celkove_citanie,
		  elapsedTime;

	LARGE_INTEGER frequency; 
    LARGE_INTEGER t1, t2; 


	srand((unsigned int)time(NULL));
	float * A = (float *)malloc(sizeof(float)*widthA*heightA);
	float * B = (float *)malloc(sizeof(float)*widthB*heightB);
	float * C = (float *)malloc(sizeof(float)*widthC*heightC);
	float * Res = (float *)malloc(sizeof(float)*widthC*heightC);
	float * D= (float *)malloc(sizeof(float)*widthC*heightC);

	FILE * fp1 = fopen("matAdata.txt", "w");
	if (!fp1) {
		fprintf(stderr, "Chyba pri otvarani matAdata.\n");
		exit(1);
	}

	for(int i = 0;i < widthA; i++)
	{
		for(int j=0; j < heightA; j++ )       {
			float p=(rand()%100)/7.0;
				*(A+i*heightA+j)=rand()%100 + p;
			fprintf(fp1, "%f ",*(A+i*heightA+j));
		}
		fprintf(fp1, "\n");
	}
	fclose(fp1);

	fp1 = fopen("matBdata.txt", "w");
	if (!fp1) {
		fprintf(stderr, "Chyba pri otvarani matBdata.\n");
		exit(1);
	}

	for(int i = 0; i < widthB; i++)
	{
		for(int j=0; j < heightB; j++ )     {
			float p=(rand()%100)/7.0;
				*((B+i*heightB+j))=rand()%100 + p;
			fprintf(fp1, "%f ",*(B+i*heightA+j));
		}
		fprintf(fp1, "\n");
	}
	fclose(fp1);         

	FILE *fp;
	char fileName[] = "./hello.cl";
	char *source_str;
	size_t source_size;
	int row = widthA;
	int col = heightC;
	// Načítanie zdrojového kódu obsahujúci kernel
	fp = fopen(fileName, "r");
	if (!fp) {
		fprintf(stderr, "Chyba pri nacitani kernelu.\n");
		exit(1);
	}
	source_str = (char*)malloc(MAX_SOURCE_SIZE);
	source_size = fread( source_str, 1, MAX_SOURCE_SIZE, fp);
	fclose( fp );

	
	// Získanie dostupných platform
	ret = clGetPlatformIDs(2, platforma, &ret_num_platforms);

	// Výber zariadenia zo zvolenej platformy
	ret = clGetDeviceIDs(platforma[0], CL_DEVICE_TYPE_GPU, 1, &zariadenie, &ret_num_devices);

	// Vytvorenie kontextu 
	context = clCreateContext(NULL, 1, &zariadenie, NULL, NULL, &ret);

	// Vytvorenie fronty
	command_queue = clCreateCommandQueue(context, zariadenie, CL_QUEUE_PROFILING_ENABLE, &ret);

	// Vytvorenie pamäťového buffera
	memobjA = clCreateBuffer(context, CL_MEM_READ_WRITE, widthA * heightA * sizeof(float), NULL, &ret);
	memobjB = clCreateBuffer(context, CL_MEM_READ_WRITE, widthB * heightB * sizeof(float), NULL, &ret);
	memobjC = clCreateBuffer(context, CL_MEM_READ_WRITE, widthC * heightC * sizeof(float), NULL, &ret);
	rowA = clCreateBuffer(context, CL_MEM_READ_WRITE,  sizeof(int), NULL, &ret);
	colC = clCreateBuffer(context, CL_MEM_READ_WRITE,  sizeof(int), NULL, &ret);

	// Kopírovanie z A a B listov do ich pamäťových bufferov
	ret = clEnqueueWriteBuffer(command_queue, memobjA, CL_TRUE, 0, widthA * heightA * sizeof(int), A, 0, NULL, &zapis1);
	clWaitForEvents(1, &zapis1);
	ret = clEnqueueWriteBuffer(command_queue, memobjB, CL_TRUE, 0, widthB * heightB * sizeof(int), B, 0, NULL, &zapis2);
	clWaitForEvents(1, &zapis2);
	ret = clEnqueueWriteBuffer(command_queue, rowA, CL_TRUE, 0, sizeof(int), &row, 0, NULL, &zapis3);
	clWaitForEvents(1, &zapis3);
	ret = clEnqueueWriteBuffer(command_queue, colC, CL_TRUE, 0, sizeof(int), &col, 0, NULL, &zapis4);
	clWaitForEvents(1, &zapis4);

	// Vytvorenie a vybudovanie kernel programu, vytvorenie kernelu
	program = clCreateProgramWithSource(context, 1, (const char **)&source_str, (const size_t *)&source_size, &ret);

	ret = clBuildProgram(program, 1, &zariadenie, NULL, NULL, NULL);
	kernel = clCreateKernel(program, "matrixMultiplication", &ret);

	// Nastavenie argumentov kernelu
	ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&memobjA);
	ret = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&memobjB);
	ret = clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&memobjC);
	ret = clSetKernelArg(kernel, 3, sizeof(int), (void *)&row);
	ret = clSetKernelArg(kernel, 4, sizeof(int), (void *)&col);

	size_t globalThreads[2] = {widthA, heightB};
	size_t localThreads[2] = {16,16};

	// Realizácia kernelu v paralélnej forme
	ret = clEnqueueNDRangeKernel(command_queue, kernel, 2, NULL, globalThreads, localThreads, NULL, 0, &kernelevent);
	clWaitForEvents(1, &kernelevent);

	// Prenos výsledkov z pamäte zariadenia
	ret = clEnqueueReadBuffer(command_queue, memobjC, CL_TRUE, 0, widthA * heightC * sizeof(float),Res, 0, NULL, &citanie);
	// Príkaz, ktorý blokuje po dobu, kým budú vykonané všetky zostavávajúce príkazy vo fronte
	clFinish(command_queue);

	//////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////

	// Zadefinovanie a začiatok merania času konkrétnych eventov
	clGetEventProfilingInfo(zapis1, CL_PROFILING_COMMAND_START, sizeof(time_szapisu1), &time_szapisu1, NULL);
	clGetEventProfilingInfo(zapis1, CL_PROFILING_COMMAND_END, sizeof(time_kzapisu1), &time_kzapisu1, NULL);
	total_zapis1 = time_kzapisu1 - time_szapisu1;
	clGetEventProfilingInfo(zapis2, CL_PROFILING_COMMAND_START, sizeof(time_szapisu2), &time_szapisu2, NULL);
	clGetEventProfilingInfo(zapis2, CL_PROFILING_COMMAND_END, sizeof(time_kzapisu2), &time_kzapisu2, NULL);
	total_zapis2 = time_kzapisu2 - time_szapisu2;
	clGetEventProfilingInfo(zapis3, CL_PROFILING_COMMAND_START, sizeof(time_szapisu3), &time_szapisu3, NULL);
	clGetEventProfilingInfo(zapis3, CL_PROFILING_COMMAND_END, sizeof(time_kzapisu3), &time_kzapisu3, NULL);
	total_zapis3 = time_kzapisu3 - time_szapisu3;
	clGetEventProfilingInfo(zapis4, CL_PROFILING_COMMAND_START, sizeof(time_szapisu4), &time_szapisu4, NULL);
	clGetEventProfilingInfo(zapis4, CL_PROFILING_COMMAND_END, sizeof(time_kzapisu4), &time_kzapisu4, NULL);
	total_zapis4 = time_kzapisu4 - time_szapisu4;
	total_zapis = total_zapis1 + total_zapis2 + total_zapis3 + total_zapis4;
	printf("\nPrenos od hostitela ku zariadeniu trval celkovo %0.3f ms\n", (total_zapis / 1000000.0) );

	clGetEventProfilingInfo(kernelevent, CL_PROFILING_COMMAND_START, sizeof(time_start), &time_start, NULL);
	clGetEventProfilingInfo(kernelevent, CL_PROFILING_COMMAND_END, sizeof(time_end), &time_end, NULL);
	total_cistokernel = time_end - time_start;
	printf("\nCas potrebny len pre vypocet kernelu = %0.3f ms\n", (total_cistokernel / 1000000.0) );

	clGetEventProfilingInfo(citanie, CL_PROFILING_COMMAND_START, sizeof(start_citania), &start_citania, NULL);
	clGetEventProfilingInfo(citanie, CL_PROFILING_COMMAND_END, sizeof(koniec_citania), &koniec_citania, NULL);
	celkove_citanie = koniec_citania - start_citania;
	printf("\nPrenos vysledku zo zariadenia k hostitelovi trval %0.3f ms\n", (celkove_citanie / 1000000.0) );


	fp1 = fopen("matCPURes.txt", "w");
	if (!fp1) {
		fprintf(stderr, "Chyba pri otváraní matCPURes.\n");
		exit(1);
	}
	printf("\nOpenCL cast je hotova\n");
	for(int i = 0; i < widthA; i++)
	{
		for(int j=0;j < heightC; j++)
		{

			fprintf(fp1, "%f ",*(Res+i*heightC+j));

		}
		fprintf(fp1, "\n");
	}
	fclose(fp1);

	// Uvoľnenie pamätí a objektov
	ret = clFlush(command_queue);
	ret = clFinish(command_queue);
	ret = clReleaseKernel(kernel);
	ret = clReleaseProgram(program);
	ret = clReleaseMemObject(memobjA);
	ret = clReleaseMemObject(memobjB);
	ret = clReleaseMemObject(memobjC);
	ret = clReleaseCommandQueue(command_queue);
	ret = clReleaseContext(context);

	free(source_str);
	system("pause");
	
	// Zadefinovanie a začiatok merania času natívnej funkcie
	QueryPerformanceFrequency(&frequency);
	QueryPerformanceCounter(&t1);

	float sum=0.0;
	for(int i = 0;i < widthA; i++)
	{
		for(int j = 0; j < heightC; j++)
		{
			sum = 0;
			for(int k = 0; k < widthB; k++)
			{
				sum += A[i*col+k] * B[k*row+j];
			}
			D[i*heightC+j] = sum;
		}
				
	}
	QueryPerformanceCounter(&t2);
	elapsedTime = (t2.QuadPart - t1.QuadPart) * 1000.0 / frequency.QuadPart;
    cout << elapsedTime << " ms pre CPU nativ.\n";
	
	
	fp1 = fopen("matNormalMultiplicationRes.txt", "w");
	if (!fp1) {
		fprintf(stderr, "Chyba pri otvarani matNativ.\n");
		exit(1);
	}
	printf("\nNativna funkcia je hotova\n");
	for(int i = 0;i < widthA; i++)
	{
		for(int j=0;j < heightC; j++)
		{
			fprintf(fp1, "%f ",*(D+i*heightC+j));

		}
		fprintf(fp1, "\n");
	}
	system("pause");
	return 0;
}