#include "Difusion.h"


__device__ uchar3
Compute_Difusion(uchar3 g,uchar3 f01,uchar3 f10,uchar3 f21,uchar3 f12,float lambda,float den)
{
	uchar3 result;
	result.x=(unsigned char)((g.x+lambda*(f01.x+f10.x+f21.x+f12.x))/den);
	result.y=(unsigned char)((g.y+lambda*(f01.y+f10.y+f21.y+f12.y))/den);
	result.z=(unsigned char)((g.z+lambda*(f01.z+f10.z+f21.z+f12.z))/den);
    return result;
}

/*Kelnel Function*/
__global__ void Difusion(uchar3 *ImageDst_dev,uchar3 *ImageSrc_dev,float lambda,float den,int w,int h){
	const int ix = blockDim.x * blockIdx.x + threadIdx.x;
    const int iy = blockDim.y * blockIdx.y + threadIdx.y;

	//if(ix>0 && ix < w-1 && iy>0 && iy < h-1){
	if(ix < w && iy < h){
		long int index_11=iy*w+ix;
		uchar3 g   = ImageSrc_dev[index_11];
		
		/*long int index_01=iy*w+ix-1;
		long int index_10=(iy-1)*w+ix;
		long int index_21=iy*w+ix+1;
		long int index_12=(iy+1)*w+ix;*/

		long int index_01=iy*w+max(ix-1,0);
		long int index_10=max((iy-1),0)*w+ix;
		long int index_21=iy*w+min(ix+1,w-1);
		long int index_12=min((iy+1),h-1)*w+ix;

		uchar3 f01 = ImageSrc_dev[index_01];
		uchar3 f10 = ImageSrc_dev[index_10];
		uchar3 f21 = ImageSrc_dev[index_21];
		uchar3 f12 = ImageSrc_dev[index_12];

		uchar3 result=Compute_Difusion(g,f01,f10,f21,f12,lambda,den);
		
		ImageDst_dev[index_11].x=result.x;
		ImageDst_dev[index_11].y=result.y;
		ImageDst_dev[index_11].z=result.z;
    }

}

extern "C"
void CUDA_Difusion(uchar3 *ImageDst_dev,uchar3 *ImageSrc_dev,float lambda,float den,int imageW,int imageH,dim3 grid,dim3 threads){
	Difusion<<<grid,threads>>>(ImageDst_dev,ImageSrc_dev,lambda,den,imageW,imageH);	
}
