
#include "Diffusion.h"

/*TEXTURES*/
texture<uchar4, 2, cudaReadModeNormalizedFloat> texImage_rgb;

cudaChannelFormatDesc uchar4tex = cudaCreateChannelDesc<uchar4>();

cudaArray *cu_image;

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

/*Kelnel Funtion*/
__global__ void Difusion(uchar4 *Image_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;
    
    const float x=(float)ix+0.5f;
    const float y=(float)iy+0.5f;

	if(ix < w && iy < h){
		float4 g   = tex2D(texImage_rgb, x, y);
		float4 f01 = tex2D(texImage_rgb, x-1, y);
		float4 f10 = tex2D(texImage_rgb, x, y-1);
		float4 f21 = tex2D(texImage_rgb, x+1, y);
		float4 f12 = tex2D(texImage_rgb, x, y+1);

		uchar3 result=Compute_Difusion(g,f01,f10,f21,f12,lambda,den);
		
		Image_dev[w*iy+ix].x=result.x;
		Image_dev[w*iy+ix].y=result.y;
		Image_dev[w*iy+ix].z=result.z;
    }

}


extern "C"
void CUDA_CreateMemoryArray(int imageW,int imageH){
	cudaMallocArray(&cu_image, &uchar4tex, imageW, imageH);
}
extern "C"
void CUDA_BindTextureToArray(){
	cudaBindTextureToArray(texImage_rgb,cu_image);
}
extern "C"
void CUDA_FreeArrays(){
	cudaFreeArray(cu_image);
}
extern "C"
void CUDA_MemcpyToArray(uchar4 *src,int imageW,int imageH){
	cudaMemcpyToArray( cu_image, 0, 0,src, imageW * imageH * sizeof(uchar4), cudaMemcpyDeviceToDevice);
}

extern "C"
void CUDA_BindTexture(uchar4 *Image_dev,int imageW,int imageH){
	cudaBindTexture2D(NULL,texImage_rgb,Image_dev,uchar4tex,imageW,imageH,imageW*sizeof(uchar4));
}

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