// Difusion.cpp : Defines the entry point for the console application.
//


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

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

#include "Difusion.h"

//División techo.
int iDivUp(int a, int b){
    return ((a % b) != 0) ? (a / b + 1) : (a / b);
}
//PARAMETERS
#define BLOCKDIM_X 16
#define BLOCKDIM_Y 16
int MaxIter=10;
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(int argc, char **argv){
	cudaSetDevice(0);
	/*Host Variables*/
	int imageW, imageH;		  //Size of Image
	cv::Mat frame_Original;   //Original Image
	cv::Mat frame_Original4c;

	/*Device Variables*/
	uchar4 *Image_dev;
	uchar4 *ImageAux_dev;

	cv::namedWindow("Original Frame",CV_WINDOW_AUTOSIZE);
	cv::namedWindow("Difused Frame",CV_WINDOW_AUTOSIZE);


	cv::createTrackbar("MaxIter", "Difused Frame",&MaxIter,100,NULL,0);
	cv::createTrackbar("lambda", "Difused Frame",&lambda,50,Change_den,0);

	char name_imag[500];
	/*Load Image*/
	sprintf(name_imag,"../../baboon.jpg");
	frame_Original = cv::imread(name_imag,1);

	/*Size of Image*/
	imageW=frame_Original.cols;
	imageH=frame_Original.rows;
	size_t size=imageW*imageH*sizeof(uchar4);
	
	frame_Original4c.create(imageH,imageW,CV_8UC(4));
	cv::cvtColor(frame_Original,frame_Original4c,CV_BGR2BGRA,0);
	
	/*Create device memory*/
	cudaMalloc((void **)&Image_dev,size);
	cudaMalloc((void **)&ImageAux_dev,size);
	CUDA_CreateMemoryArray(imageW,imageH);

	/*Copy Memory (Host-->Device)*/
	CUDA_SAFE_CALL(cudaMemcpy(Image_dev,frame_Original4c.data,size,cudaMemcpyHostToDevice));
	
	/*Define the size of the grid and thread blocks*/
	dim3 threads(BLOCKDIM_X, BLOCKDIM_Y,1);
	dim3 grid(iDivUp(imageW, BLOCKDIM_X), iDivUp(imageH, BLOCKDIM_Y),1);

	CUDA_MemcpyToArray(Image_dev,imageW,imageH);
	CUDA_BindTextureToArray();
	//CUDA_BindTexture(Image_dev,imageW,imageH);
		
	cv::imshow("Original Frame",frame_Original);
	while(1){
		CUDA_SAFE_CALL(cudaMemcpy(ImageAux_dev,Image_dev,size,cudaMemcpyDeviceToDevice));		
		/*Launch the Kernel Function*/
		for(int iter=0;iter<MaxIter;iter++){
			CUDA_Difusion(ImageAux_dev,lambdaf,denf,imageW,imageH,grid,threads);	
			//Actualizamos la texture de la imagen f(x)
			CUDA_MemcpyToArray(ImageAux_dev,imageW,imageH);
		}

		/*Copy Memory (Device-->Host)*/
		CUDA_SAFE_CALL(cudaMemcpy(frame_Original4c.data,ImageAux_dev,size,cudaMemcpyDeviceToHost));
		
		cv::imshow("Difused Frame",frame_Original4c);
		
		char key=cv::waitKey(2);
		if(key==27)
			break;	

		CUDA_MemcpyToArray(Image_dev,imageW,imageH);
	}

	/*Clean Memory*/
	/*Host*/
	cv::destroyWindow("Original Frame");
	cv::destroyWindow("Difused Frame");
	
	/*Device*/
	cudaFree(Image_dev);
	cudaFree(ImageAux_dev);
	CUDA_FreeArrays();

	return(0);
}

void IplImageTouchar3(IplImage *src,uchar4 *dst){
	int imageW=src->width;
	int imageH=src->height;
	int nc=imageW*3,k=0;
	int step= src->widthStep;
	unsigned char *data=reinterpret_cast<unsigned char *>(src->imageData);	
	for(int i=0;i<imageH;i++){
		k=0;
		for (int j=0; j<nc;j+=3) {
			(dst)[(i * imageW + k)].x = (int)(data[j+2]);//BGR in OpenCV
            (dst)[(i * imageW + k)].y = (int)(data[j+1]);
            (dst)[(i * imageW + k)].z = (int)(data[j]);
            k++;
		}
		data+= step;
	}	
}

void uchar3ToIplImage(uchar4 *src,IplImage *dst){
	int imageW=dst->width;
	int imageH=dst->height;
	int nc=imageW*3,k=0;
	int step= dst->widthStep;
	IplImage *frameaux;
	frameaux=cvCloneImage(dst);
	frameaux->imageData=dst->imageData;
	for(int i=0;i<imageH;i++){
		k=0;
		for (int j=0; j<nc;j+=3) {
			dst->imageData[j+2] = (src)[(i * imageW + k)].x;//En OpenCV el orden es BGR
            dst->imageData[j+1] = (src)[(i * imageW + k)].y ;
            dst->imageData[j]   = (src)[(i * imageW + k)].z ;
            k++;
		}
		dst->imageData+= step;
	}	
	dst->imageData=frameaux->imageData;
	cvReleaseImage(&frameaux);
}
