
#include "Difusion.h"

#include <cutil.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=/*HERE*/; // The same (only chance "x" by "y")
	result.z=/*HERE*/; // The same (only chance "x" by "z");
    return result;
}

/***********************************/
/***********************************/
/*Kelnel Funtion*/
__global__ void Difusion(uchar4 *Image_dev,float lambda,float den,int w,int h){
	
    // Wath is the pixel?
    const int ix = /*HERE*/;
    const int iy = /*HERE*/;
    
    const float x=(float)ix+0.5f;
    const float y=(float)iy+0.5f;

    //Wath are the condition?
	if(/*HERE*/){
		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);
		
        // Asing the results
		Image_dev[w*iy+ix].x=/*HERE*/; // result for "x"
		Image_dev[w*iy+ix].y=/*HERE*/; // result for "y"
		Image_dev[w*iy+ix].z=/*HERE*/; // result for "z"
    }

}

/***********************************/
/***********************************/
extern "C"
void CUDA_CreateMemoryArray(int imageW,int imageH){
	CUDA_SAFE_CALL(cudaMallocArray(&cu_image, &uchar4tex, imageW, imageH) );
}

/***********************************/
/***********************************/
extern "C"
void CUDA_BindTextureToArray(){
	CUDA_SAFE_CALL(cudaBindTextureToArray(texImage_rgb,cu_image));
}

/***********************************/
/***********************************/
extern "C"
void CUDA_FreeArrays(){
	CUDA_SAFE_CALL(cudaFreeArray(cu_image));
}

/***********************************/
/***********************************/
extern "C"
void CUDA_MemcpyToArray(uchar4 *src,int imageW,int imageH){
	CUDA_SAFE_CALL(cudaMemcpyToArray( cu_image, 0, 0,src, imageW * imageH * sizeof(uchar4), cudaMemcpyDeviceToDevice));
}

/***********************************/
/***********************************/
extern "C"
void CUDA_BindTexture(uchar4 *Image_dev,int imageW,int imageH){
	CUDA_SAFE_CALL(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);	
}