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

/*__constant__ float MeanKernel[9] = {1,1,1,  
                                    1,1,1,
									1,1,1};*/

__constant__ float LaplacianKernel[9] = {0,1,0,  
                                         1,-4,1,
									     0,1,0};

__constant__ float GaussianKernel[9] = {1,2,1,  
                                        2,4,2,
									    1,2,1};

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

	int win_W=1;

	if(ix < w && iy < h){
		float4 pixval;
		float3 sum;
		sum.x=0.0f;
		sum.y=0.0f;
		sum.z=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);
				sum.x+=pixval.x;
				sum.y+=pixval.y;
				sum.z+=pixval.z;
				/*sum.x+=pixval.x*MeanKernel[k];
				sum.y+=pixval.y*MeanKernel[k];
				sum.z+=pixval.z*MeanKernel[k];*/
				k++;
			}
		}
		Image_dev[w*iy+ix].x=(unsigned char)((sum.x/9)*255);
		Image_dev[w*iy+ix].y=(unsigned char)((sum.y/9)*255);
		Image_dev[w*iy+ix].z=(unsigned char)((sum.z/9)*255);
	}
}
__global__ void LaplacianFilter(float *Image_dev,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;

	int win_W=1;

	if(ix < w && iy < h){
		float4 pixval;
		/*float3 sum;
		sum.x=0.0f;
		sum.y=0.0f;
		sum.z=0.0f;*/
		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);
				float Intensidad=(pixval.x+pixval.y+pixval.z)/3.0f;
				sum+=Intensidad*LaplacianKernel[k];
				/*
				sum.x+=pixval.x*LaplacianKernel[k];
				sum.y+=pixval.y*LaplacianKernel[k];
				sum.z+=pixval.z*LaplacianKernel[k];
*/
				k++;
			}
		}
		Image_dev[w*iy+ix]=sum;
		/*
		Image_dev[w*iy+ix].x=(unsigned char)((sum.x)*255);
		Image_dev[w*iy+ix].y=(unsigned char)((sum.y)*255);
		Image_dev[w*iy+ix].z=(unsigned char)((sum.z)*255);
		*/
		/*unsigned char result=(unsigned char)((sum)*255);
		Image_dev[w*iy+ix].x=result;
		Image_dev[w*iy+ix].y=result;
		Image_dev[w*iy+ix].z=result;*/
	}
}
__global__ void GaussianFilter(uchar4 *Image_dev,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;

	int win_W=1;

	if(ix < w && iy < h){
		float4 pixval;
		float3 sum;
		sum.x=0.0f;
		sum.y=0.0f;
		sum.z=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);
				sum.x+=pixval.x*GaussianKernel[k];
				sum.y+=pixval.y*GaussianKernel[k];
				sum.z+=pixval.z*GaussianKernel[k];
				k++;
			}
		}
		Image_dev[w*iy+ix].x=(unsigned char)((sum.x/16)*255);
		Image_dev[w*iy+ix].y=(unsigned char)((sum.y/16)*255);
		Image_dev[w*iy+ix].z=(unsigned char)((sum.z/16)*255);
	}
}

extern "C"
void CUDA_MeanFilter(uchar4 *Image_dev,int imageW,int imageH,dim3 grid,dim3 threads){
	MeanFilter<<<grid,threads>>>(Image_dev,imageW,imageH);
}
extern "C"
void CUDA_LaplacianFilter(float *Image_dev,int imageW,int imageH,dim3 grid,dim3 threads){
	LaplacianFilter<<<grid,threads>>>(Image_dev,imageW,imageH);
}
extern "C"
void CUDA_GaussianFilter(uchar4 *Image_dev,int imageW,int imageH,dim3 grid,dim3 threads){
	GaussianFilter<<<grid,threads>>>(Image_dev,imageW,imageH);
}