
#include "Image_Filters.h"
#include <cutil.h>

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

cudaChannelFormatDesc uchar4tex = cudaCreateChannelDesc<uchar4>();

cudaArray *cu_image;

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);
}


// Fill the laplacian kernel
__constant__ float LaplacianKernel[9] = {/*HERE*/};


/**************************************/
/**************************************/
/*Image Filters*/
__global__ void LaplacianFilter(float *Image_dev,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;

	int win_W=1;
    
    // Wath are the conditions?
	if(/*HERE*/ && /*HERE*/){
		float4 pixval;		
		float sum=0.0f;

		int k=0;
		for(int ii=-win_W;ii<=win_W;ii++){
			for(int jj=-win_W;jj<=win_W;jj++){
				pixval=tex2D(texImage_rgb,x+ii,y+jj);
                
                // Wath is the intensity?
				float Intensidad=(/*HERE*/+/*HERE*/+/*HERE*/)/3.0f;
				sum+=Intensidad*LaplacianKernel[k];				
				k++;
			}
		}
        
        // Write the result
		Image_dev[w*iy+ix]=/*HERE*/;
	
	}
}

/**************************************/
/**************************************/
extern "C"
void CUDA_LaplacianFilter(float *Image_dev,int imageW,int imageH,dim3 grid,dim3 threads){
    
    // Complete
	LaplacianFilter<<</*HERE*/,/*HERE*/>>>(Image_dev,imageW,imageH);
}
