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

/*CUDA*/
#include <cuda_runtime.h>


/*OpenCV*/
#include <opencv2/highgui/highgui.hpp>
#include <opencv2/imgproc/imgproc.hpp>

#include "Diffusion.h"

//OpenMP
#include <omp.h>

//División techo.
int iDivUp(int a, int b){
    return ((a % b) != 0) ? (a / b + 1) : (a / b);
}

//PARAMETERS
#define NumberOfImages 4 //Define the number of images that will be processed

#define BLOCKDIM_X 16
#define BLOCKDIM_Y 16
int MaxIter=100;
int lambda=1;
float lambdaf=(float)lambda/10.0f;
float denf=1+lambdaf*4;
void Change_den(int,void*){
	lambdaf=(float)lambda/10.0f;
	denf=1+lambdaf*4;
}

int main (void){
	int num_gpus = 0;
	int id;
	double total_time;
	char time_ms[100];

	cudaGetDeviceCount(&num_gpus);//How many devices are there?
	if(num_gpus < 1) 
	{
		printf("There are not enabled CUDA devices...\n");
		return 1;
	}

	printf("Number of Host CPUs:\t%d\n", omp_get_num_procs());
    printf("Number of Devices CUDA:\t%d\n", num_gpus);
	for(int i = 0; i < num_gpus; i++)
    {
        cudaDeviceProp dprop;
        cudaGetDeviceProperties(&dprop, i);
		printf("   %d: %s\n", i, dprop.name);
    }
	printf("---------------------------\n");
	printf("---------------------------\n");

	omp_set_num_threads(NumberOfImages);//Creating threads in the host, many as Images
	
	//num_gpus = 4;

#pragma omp parallel
	{
		unsigned int cpu_thread_id = omp_get_thread_num();
		unsigned int num_cpu_threads = omp_get_num_threads();
		int gpu_id = -1;
		cudaSetDevice(cpu_thread_id % num_gpus);
		cudaGetDevice(&gpu_id);
		printf("CPU thread %d (of %d) uses CUDA device %d\n", 
			   cpu_thread_id, num_cpu_threads, gpu_id);
	}
	printf("---------------------------\n");

	/*Host Variables*/
	int imageW[NumberOfImages];
	int imageH[NumberOfImages];
	size_t size_Image[NumberOfImages];

	cv::Mat frame_Original[NumberOfImages];   //Original Image
	cv::Mat frame_Text;
	cv::Mat frame_Zeros;

	/*Device Variables*/
	uchar3 *Imagef0_dev[NumberOfImages];
	uchar3 *Imagef1_dev[NumberOfImages];
	uchar3 *ImageAux_dev[NumberOfImages];

	dim3 threads[NumberOfImages];
	dim3 grid[NumberOfImages];

	//Display Original Image
	char name_Window[NumberOfImages][100];
	char name_Image[NumberOfImages][100];
	char name_WindowResult[NumberOfImages][100];
	int pos_x=10;
	int pos_y=10;
	for(id=0;id<NumberOfImages;id++){
		sprintf(name_Window[id],"Original Frame_%d",id+1);
		sprintf(name_WindowResult[id],"Smoothed Frame_%d",id+1);
		cv::namedWindow(name_Window[id],CV_WINDOW_AUTOSIZE);
		cv::namedWindow(name_WindowResult[id],CV_WINDOW_AUTOSIZE);
		sprintf(name_Image[id],"../../Images/Image_%d.jpg",id+1);
	}
	
	
	printf("\n\nInitialization...\n\n");
#pragma omp parallel for
	for(id=0;id<NumberOfImages;id++){
		//int cpu_thread_id = omp_get_thread_num();
		int cpu_thread_id = id;
		int gpu_id = -1;
		int device=cpu_thread_id%num_gpus;
		cudaSetDevice(device);
		cudaGetDevice(&gpu_id);
		printf("CPU thread %d uses CUDA device %d\n", cpu_thread_id, gpu_id);
					
		/*Load Image*/			
		frame_Original[cpu_thread_id] = cv::imread(name_Image[cpu_thread_id],1);
		imageW[cpu_thread_id]=frame_Original[cpu_thread_id].cols;
		imageH[cpu_thread_id]=frame_Original[cpu_thread_id].rows;
		size_Image[cpu_thread_id]=imageW[cpu_thread_id]*imageH[cpu_thread_id]*sizeof(uchar3);

		/*Create device memory*/
		cudaMalloc((void **)&Imagef0_dev[cpu_thread_id],size_Image[cpu_thread_id]);
		cudaMalloc((void **)&Imagef1_dev[cpu_thread_id],size_Image[cpu_thread_id]);
		cudaMalloc((void **)&ImageAux_dev[cpu_thread_id],size_Image[cpu_thread_id]);
					
		/*Copy Memory (Host-->Device)*/
		cudaMemcpy(ImageAux_dev[cpu_thread_id],frame_Original[cpu_thread_id].data,size_Image[cpu_thread_id],cudaMemcpyHostToDevice);
		
		/*Define the size of the grid and thread blocks*/
		threads[cpu_thread_id].x=BLOCKDIM_X;
		threads[cpu_thread_id].y=BLOCKDIM_Y;
		threads[cpu_thread_id].z=1;

		grid[cpu_thread_id].x=iDivUp(*imageW, BLOCKDIM_X);
		grid[cpu_thread_id].y=iDivUp(*imageH, BLOCKDIM_Y);
		grid[cpu_thread_id].z=1;
	}

	for(id=0;id<NumberOfImages;id++){
		cv::imshow(name_Window[id],frame_Original[id]);
		//cv::moveWindow(name_Window[id], pos_x+id*imageW[id],pos_y);
		//cv::moveWindow(name_WindowResult[id], pos_x+id*imageW[id],pos_y+imageH[id]);
	}
	cv::namedWindow("Parameters",CV_WINDOW_AUTOSIZE);
	cv::createTrackbar("MaxIter", "Parameters",&MaxIter,500,NULL,0);
	cv::createTrackbar("lambda", "Parameters",&lambda,50,Change_den,0);
	int wt=400,ht=100;
	frame_Text.create(ht,wt,CV_8UC(3));
	frame_Zeros=cv::Mat::zeros(ht,wt,CV_8UC(3));
	printf("\n\nSmoothing...\n\n");
	

	while(1){
	//for(int iter=0;iter<100;iter++){		
		total_time = (double)cv::getTickCount();
		#pragma omp parallel for
		for(id=0;id<NumberOfImages;id++){
			//int cpu_thread_id = omp_get_thread_num();
			int cpu_thread_id = id;
			int device=cpu_thread_id%num_gpus;
			cudaSetDevice(device);
			/*Copy Memory (Device-->Device)*/
			cudaMemcpy(Imagef0_dev[cpu_thread_id],ImageAux_dev[cpu_thread_id],size_Image[cpu_thread_id],cudaMemcpyDeviceToDevice);		
			/*Launch the Kernel Function*/
			for(int iter=0;iter<MaxIter;iter++){
				CUDA_Difusion(Imagef1_dev[cpu_thread_id],Imagef0_dev[cpu_thread_id],lambdaf,denf,imageW[cpu_thread_id],imageH[cpu_thread_id],grid[cpu_thread_id],threads[cpu_thread_id]);
				//Update f(x)
				/*Copy Memory (Device-->Device)*/
				cudaMemcpy(Imagef0_dev[cpu_thread_id],Imagef1_dev[cpu_thread_id],size_Image[cpu_thread_id],cudaMemcpyDeviceToDevice);		
			}
			/*Copy Memory (Device-->Host)*/
			//Case 1: With copy memory D-->H inside of omp parallel for
			//(Comment the cudaMemcpy of Case 2)
			cudaMemcpy(frame_Original[id].data,Imagef0_dev[id],size_Image[id],cudaMemcpyDeviceToHost);
		}

		cudaThreadSynchronize();
		total_time=(double)cv::getTickCount()-total_time;
			
		//Copy results to the CPU memory
		for(id=0;id<NumberOfImages;id++){
			int device=id%num_gpus;
			cudaSetDevice(device);
			/*Copy Memory (Device-->Host)*/
			//Case 2: With copy memory D-->H outside of omp parallel for
			//(Comment the cudaMemcpy of Case 1)
			//cudaMemcpy(frame_Original[id].data,Imagef1_dev[id],size_Image[id],cudaMemcpyDeviceToHost);
			cv::imshow(name_WindowResult[id],frame_Original[id]);
		}
		
		//total_time=(double)cv::getTickCount()-total_time;

		frame_Zeros.copyTo(frame_Text);
		sprintf(time_ms,"Total time: %.2lf ms",(total_time/cv::getTickFrequency())*1000);
		cv::putText(frame_Text,time_ms,cv::Point(5,30),CV_FONT_HERSHEY_PLAIN,2.0,CV_RGB(255,0,0),2,8);
		cv::imshow("Parameters",frame_Text);
			

		char key=cv::waitKey(2);
		if(key==27)
			break;	
	}
	
	/*Clean Memory*/
	/*Host*/
	cv::destroyAllWindows();
	/*Device*/
	for (int id = 0;id<num_gpus;id++){
		cudaSetDevice(id);
		cudaFree(Imagef0_dev[id]);
		cudaFree(Imagef1_dev[id]);
		cudaFree(ImageAux_dev[id]);
        cudaDeviceReset();
    }

	return(0);
}
