

__global__ void Dx(int3 *Imag_dst,uchar3 *Imag_src,int N,int Width)
{
	int idx = blockIdx.x * blockDim.x + threadIdx.x;
	
	if (idx<N){
		if (idx%Width==0){//Cero en los bordes
			Imag_dst[idx].x=0;
			Imag_dst[idx].y=0;
			Imag_dst[idx].z=0;
		}
		else{
			Imag_dst[idx].x=(int)Imag_src[idx].x-Imag_src[idx-1].x;
			Imag_dst[idx].y=(int)Imag_src[idx].y-Imag_src[idx-1].y;
			Imag_dst[idx].z=(int)Imag_src[idx].z-Imag_src[idx-1].z;
		}
	}
}

__global__ void Dy(int3 *Imag_dst,uchar3 *Imag_src,int N,int Width)
{
	int idx = blockIdx.x * blockDim.x + threadIdx.x;
	
	if (idx<N){
		if(idx<Width){//Cero en los bordes
			Imag_dst[idx].x=0.0;
			Imag_dst[idx].y=0.0;
			Imag_dst[idx].z=0.0;
		}
		else{
			Imag_dst[idx].x=(int)Imag_src[idx].x-Imag_src[idx-Width].x;
			Imag_dst[idx].y=(int)Imag_src[idx].y-Imag_src[idx-Width].y;
			Imag_dst[idx].z=(int)Imag_src[idx].z-Imag_src[idx-Width].z;
		} 
	}


}

__global__ void MG(float *ImageMG_dev,int3 *ImagDx_dev,int3 *ImagDy_dev,int N)
{
	int idx = blockIdx.x * blockDim.x + threadIdx.x;
	
	int3 dx,dy;
	float cuad1,cuad2,mag_g;
	if (idx<N){
		dx=ImagDx_dev[idx];
		dy=ImagDy_dev[idx];
		cuad1=dx.x*dx.x+dx.y*dx.y+dx.z*dx.z;
		cuad2=dy.x*dy.x+dy.y*dy.y+dy.z*dy.z;
		mag_g=sqrtf(cuad1+cuad2);
		ImageMG_dev[idx]=mag_g;
	}

}


extern "C"
void CUDA_Dx(int3 *ImageDx_dev,uchar3 *Image_dev,int N,int imageW,dim3 grid,dim3 threads){
	Dx<<<grid,threads>>>(ImageDx_dev,Image_dev,N,imageW);
}

extern "C"
void CUDA_Dy(int3 *ImageDy_dev,uchar3 *Image_dev,int N,int imageW,dim3 grid,dim3 threads){
	Dy<<<grid,threads>>>(ImageDy_dev,Image_dev,N,imageW);
}

extern "C"
void CUDA_MG(float *ImageMG_dev,int3 *ImageDx_dev,int3 *ImageDy_dev,int N,dim3 grid,dim3 threads){
	MG<<<grid,threads>>>(ImageMG_dev,ImageDx_dev,ImageDy_dev,N);
}