#ifdef _WIN32
	#define WIN32_LEAN_AND_MEAN
#endif

#include <stdio.h>
#include "ritkCudaMacros.h"
#include "BilateralFilterKernel.h"


#define BLOCKSIZE_X 8
#define BLOCKSIZE_Y 8

//#define RUNTIME_EVALUATION


//----------------------------------------------------------------------------
extern "C"
void
BilateralFilterWithTexture(const cudaArray *InputImageArray, float *DeviceOutput, size_t OutputPitch, int w, int h, int KernelRadius, float SigmaD, float SigmaR)
{
	// Set input image texture parameters and bind texture to the array. Texture is defined in the kernel
	InputImageTexture.addressMode[0] = cudaAddressModeClamp;
	InputImageTexture.addressMode[1] = cudaAddressModeClamp;
	InputImageTexture.filterMode = cudaFilterModePoint;
	InputImageTexture.normalized = false;
	ritkCudaSafeCall(cudaBindTextureToArray(InputImageTexture, InputImageArray));

	// Timer for runtime evaluation
	/*unsigned int Timer = 0;
	CUT_SAFE_CALL(cutCreateTimer(&Timer));
	CUT_SAFE_CALL(cutStartTimer(Timer));*/

	// A normalized output pitch allows us to operate on array indices instead of using pointer arithmetic
	int NormalizedOutputPitch = (int)OutputPitch/sizeof(float);
	
	//for ( int i = 0; i < 1000; i++ ) {
	dim3 DimBlock(BLOCKSIZE_X,BLOCKSIZE_Y);
	dim3 DimGrid(DivUp(w,DimBlock.x),DivUp(h,DimBlock.y));
	if ( KernelRadius == 1 )
		BilateralFilterKernel<1,BLOCKSIZE_X,BLOCKSIZE_Y><<<DimGrid,DimBlock>>>(w,h,DeviceOutput,NormalizedOutputPitch,SigmaD,SigmaD*SigmaD,SigmaR);
	else if ( KernelRadius == 2 ) 
		BilateralFilterKernel<2,BLOCKSIZE_X,BLOCKSIZE_Y><<<DimGrid,DimBlock>>>(w,h,DeviceOutput,NormalizedOutputPitch,SigmaD,SigmaD*SigmaD,SigmaR);
	else if ( KernelRadius == 3 )
		BilateralFilterKernel<3,BLOCKSIZE_X,BLOCKSIZE_Y><<<DimGrid,DimBlock>>>(w,h,DeviceOutput,NormalizedOutputPitch,SigmaD,SigmaD*SigmaD,SigmaR);
	else if ( KernelRadius == 4 )
		BilateralFilterKernel<4,BLOCKSIZE_X,BLOCKSIZE_Y><<<DimGrid,DimBlock>>>(w,h,DeviceOutput,NormalizedOutputPitch,SigmaD,SigmaD*SigmaD,SigmaR);
	else if ( KernelRadius == 5 )
		BilateralFilterKernel<5,BLOCKSIZE_X,BLOCKSIZE_Y><<<DimGrid,DimBlock>>>(w,h,DeviceOutput,NormalizedOutputPitch,SigmaD,SigmaD*SigmaD,SigmaR);
	else if ( KernelRadius == 6 )
		BilateralFilterKernel<6,BLOCKSIZE_X,BLOCKSIZE_Y><<<DimGrid,DimBlock>>>(w,h,DeviceOutput,NormalizedOutputPitch,SigmaD,SigmaD*SigmaD,SigmaR);
	else if ( KernelRadius == 7 )
		BilateralFilterKernel<7,BLOCKSIZE_X,BLOCKSIZE_Y><<<DimGrid,DimBlock>>>(w,h,DeviceOutput,NormalizedOutputPitch,SigmaD,SigmaD*SigmaD,SigmaR);
	else if ( KernelRadius == 8 )
		BilateralFilterKernel<8,BLOCKSIZE_X,BLOCKSIZE_Y><<<DimGrid,DimBlock>>>(w,h,DeviceOutput,NormalizedOutputPitch,SigmaD,SigmaD*SigmaD,SigmaR);
	else if ( KernelRadius == 9 )
		BilateralFilterKernel<9,BLOCKSIZE_X,BLOCKSIZE_Y><<<DimGrid,DimBlock>>>(w,h,DeviceOutput,NormalizedOutputPitch,SigmaD,SigmaD*SigmaD,SigmaR);
	else if ( KernelRadius == 10 )
		BilateralFilterKernel<10,BLOCKSIZE_X,BLOCKSIZE_Y><<<DimGrid,DimBlock>>>(w,h,DeviceOutput,NormalizedOutputPitch,SigmaD,SigmaD*SigmaD,SigmaR);
	else if ( KernelRadius == 11 )
		BilateralFilterKernel<11,BLOCKSIZE_X,BLOCKSIZE_Y><<<DimGrid,DimBlock>>>(w,h,DeviceOutput,NormalizedOutputPitch,SigmaD,SigmaD*SigmaD,SigmaR);
	else if ( KernelRadius == 12 )
		BilateralFilterKernel<12,BLOCKSIZE_X,BLOCKSIZE_Y><<<DimGrid,DimBlock>>>(w,h,DeviceOutput,NormalizedOutputPitch,SigmaD,SigmaD*SigmaD,SigmaR);
	else if ( KernelRadius == 13 )
		BilateralFilterKernel<13,BLOCKSIZE_X,BLOCKSIZE_Y><<<DimGrid,DimBlock>>>(w,h,DeviceOutput,NormalizedOutputPitch,SigmaD,SigmaD*SigmaD,SigmaR);
	else if ( KernelRadius == 14 )
		BilateralFilterKernel<14,BLOCKSIZE_X,BLOCKSIZE_Y><<<DimGrid,DimBlock>>>(w,h,DeviceOutput,NormalizedOutputPitch,SigmaD,SigmaD*SigmaD,SigmaR);
	else
		printf("Kernel radius not supported: %d\n", KernelRadius);
	//}

	// Runtime evaluation
	/*cudaThreadSynchronize();
	CUT_SAFE_CALL(cutStopTimer(Timer));
	printf("	Runtime BF: %fms \n\n", cutGetTimerValue(Timer));
	CUT_SAFE_CALL(cutDeleteTimer(Timer));*/

	// Unbind texture
	ritkCudaSafeCall(cudaUnbindTexture(InputImageTexture));
}


//----------------------------------------------------------------------------
extern "C"
float
BilateralFilter(const cudaArray *InputImageArray, const cudaArray *PrecomputedGaussianArray, float *DeviceOutput, size_t OutputPitch, int w, int h, int KernelRadius, float SigmaD, float SigmaR)
{
	// Elapsed time
	float ElapsedTime = -1;

	// Set input image texture parameters and bind texture to the array. Texture is defined in the kernel
	InputImageTexture.addressMode[0] = cudaAddressModeClamp;
	InputImageTexture.addressMode[1] = cudaAddressModeClamp;
	InputImageTexture.filterMode = cudaFilterModePoint;
	InputImageTexture.normalized = false;
	ritkCudaSafeCall(cudaBindTextureToArray(InputImageTexture, InputImageArray));

	// Set the precomputed gaussian values and bind texture to the array
	PrecomputedGaussianTexture.addressMode[0] = cudaAddressModeClamp;
	PrecomputedGaussianTexture.addressMode[1] = cudaAddressModeClamp;
	PrecomputedGaussianTexture.filterMode = cudaFilterModePoint;
	PrecomputedGaussianTexture.normalized = false;
	ritkCudaSafeCall(cudaBindTextureToArray(PrecomputedGaussianTexture, PrecomputedGaussianArray));

	// A normalized output pitch allows us to operate on array indices instead of using pointer arithmetic
	int NormalizedOutputPitch = (int)OutputPitch/sizeof(float);
	
	// Timer for runtime evaluation
#ifdef RUNTIME_EVALUATION
	unsigned int NumRuntimeIterations = 1000;
	unsigned int Timer = 0;
	CUT_SAFE_CALL(cutCreateTimer(&Timer));
	CUT_SAFE_CALL(cutStartTimer(Timer));
	for ( unsigned int i = 0; i < NumRuntimeIterations; i++ ) 
	{
#endif

	dim3 DimBlock(BLOCKSIZE_X,BLOCKSIZE_Y);
	dim3 DimGrid(DivUp(w,DimBlock.x),DivUp(h,DimBlock.y));
	if ( KernelRadius == 1 )
		BilateralFilterKernel<1,BLOCKSIZE_X,BLOCKSIZE_Y><<<DimGrid,DimBlock>>>(w,h,DeviceOutput,NormalizedOutputPitch,SigmaD,SigmaD*SigmaD,SigmaR);
	else if ( KernelRadius == 2 ) 
		BilateralFilterKernel<2,BLOCKSIZE_X,BLOCKSIZE_Y><<<DimGrid,DimBlock>>>(w,h,DeviceOutput,NormalizedOutputPitch,SigmaD,SigmaD*SigmaD,SigmaR);
	else if ( KernelRadius == 3 )
		BilateralFilterKernel<3,BLOCKSIZE_X,BLOCKSIZE_Y><<<DimGrid,DimBlock>>>(w,h,DeviceOutput,NormalizedOutputPitch,SigmaD,SigmaD*SigmaD,SigmaR);
	else if ( KernelRadius == 4 )
		BilateralFilterKernel<4,BLOCKSIZE_X,BLOCKSIZE_Y><<<DimGrid,DimBlock>>>(w,h,DeviceOutput,NormalizedOutputPitch,SigmaD,SigmaD*SigmaD,SigmaR);
	else if ( KernelRadius == 5 )
		BilateralFilterKernel<5,BLOCKSIZE_X,BLOCKSIZE_Y><<<DimGrid,DimBlock>>>(w,h,DeviceOutput,NormalizedOutputPitch,SigmaD,SigmaD*SigmaD,SigmaR);
	else if ( KernelRadius == 6 )
		BilateralFilterKernel<6,BLOCKSIZE_X,BLOCKSIZE_Y><<<DimGrid,DimBlock>>>(w,h,DeviceOutput,NormalizedOutputPitch,SigmaD,SigmaD*SigmaD,SigmaR);
	else if ( KernelRadius == 7 )
		BilateralFilterKernel<7,BLOCKSIZE_X,BLOCKSIZE_Y><<<DimGrid,DimBlock>>>(w,h,DeviceOutput,NormalizedOutputPitch,SigmaD,SigmaD*SigmaD,SigmaR);
	else if ( KernelRadius == 8 )
		BilateralFilterKernel<8,BLOCKSIZE_X,BLOCKSIZE_Y><<<DimGrid,DimBlock>>>(w,h,DeviceOutput,NormalizedOutputPitch,SigmaD,SigmaD*SigmaD,SigmaR);
	else if ( KernelRadius == 9 )
		BilateralFilterKernel<9,BLOCKSIZE_X,BLOCKSIZE_Y><<<DimGrid,DimBlock>>>(w,h,DeviceOutput,NormalizedOutputPitch,SigmaD,SigmaD*SigmaD,SigmaR);
	else if ( KernelRadius == 10 )
		BilateralFilterKernel<10,BLOCKSIZE_X,BLOCKSIZE_Y><<<DimGrid,DimBlock>>>(w,h,DeviceOutput,NormalizedOutputPitch,SigmaD,SigmaD*SigmaD,SigmaR);
	else if ( KernelRadius == 11 )
		BilateralFilterKernel<11,BLOCKSIZE_X,BLOCKSIZE_Y><<<DimGrid,DimBlock>>>(w,h,DeviceOutput,NormalizedOutputPitch,SigmaD,SigmaD*SigmaD,SigmaR);
	else if ( KernelRadius == 12 )
		BilateralFilterKernel<12,BLOCKSIZE_X,BLOCKSIZE_Y><<<DimGrid,DimBlock>>>(w,h,DeviceOutput,NormalizedOutputPitch,SigmaD,SigmaD*SigmaD,SigmaR);
	else if ( KernelRadius == 13 )
		BilateralFilterKernel<13,BLOCKSIZE_X,BLOCKSIZE_Y><<<DimGrid,DimBlock>>>(w,h,DeviceOutput,NormalizedOutputPitch,SigmaD,SigmaD*SigmaD,SigmaR);
	else if ( KernelRadius == 14 )
		BilateralFilterKernel<14,BLOCKSIZE_X,BLOCKSIZE_Y><<<DimGrid,DimBlock>>>(w,h,DeviceOutput,NormalizedOutputPitch,SigmaD,SigmaD*SigmaD,SigmaR);
	else
		printf("Kernel radius not supported: %d\n", KernelRadius);
	//}

#ifdef RUNTIME_EVALUATION
		}
		// Runtime evaluation
		cudaThreadSynchronize();
		CUT_SAFE_CALL(cutStopTimer(Timer));
		printf("	Runtime for KernelSize=%d: %fms \n\n", KernelRadius, (cutGetTimerValue(Timer)/NumRuntimeIterations));
		ElapsedTime = cutGetTimerValue(Timer)/NumRuntimeIterations;
		CUT_SAFE_CALL(cutDeleteTimer(Timer));

#endif

	// Unbind texture
	ritkCudaSafeCall(cudaUnbindTexture(InputImageTexture));
	ritkCudaSafeCall(cudaUnbindTexture(PrecomputedGaussianTexture));

	return ElapsedTime;
}
