#ifdef _WIN32
	#define WIN32_LEAN_AND_MEAN
#endif

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

#include <cuda_runtime.h>
#include <driver_types.h>
#include <channel_descriptor.h>
#include <cuda_runtime_api.h>

extern "C"
void
MedianFilterOld(PixelType* dI, PixelType* dR, int w, int h, int KernelRadius)
{
	// Create the texture
	// RGB channel descriptor
	cudaChannelFormatDesc ChannelDesc = cudaCreateChannelDesc(32,0,0,0,cudaChannelFormatKindFloat);

	// Allocate cuda array for the input and copy data from host to device
	cudaArray *InputImageArray = NULL;
	ritkCudaSafeCall(cudaMallocArray(&InputImageArray,&ChannelDesc,w,h));
	ritkCudaSafeCall(cudaMemcpyToArray(InputImageArray, 0, 0, dI, w*h*sizeof(PixelType), cudaMemcpyDeviceToDevice));

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

	// Compute the kernel size
	int KernelSize = (KernelRadius*2+1)*(KernelRadius*2+1);

	// Compute the number of threads per block such that the size of the shared memory
	// is below 16Kb
    int NumThreadsPerBlock = 256;
	int SharedMemSize = NumThreadsPerBlock * KernelSize * sizeof(PixelType);
	while ( SharedMemSize > 1024*16 )
	{
		NumThreadsPerBlock /= 2;
		SharedMemSize = NumThreadsPerBlock * KernelSize * sizeof(PixelType);
	}
	if ( NumThreadsPerBlock < 1 )
	{
		printf("Kernel radius %d is too large!\n",KernelRadius);
		exit(-1);
	}

    // Compute number of blocks needed based on image size and desired block size
    int NumBlocks = DivUp(w*h,NumThreadsPerBlock);  

	
	// Launch kernel 
	dim3 DimGrid(NumBlocks);
	dim3 DimBlock(NumThreadsPerBlock);
	d_MedianFilter<<< DimGrid, DimBlock, SharedMemSize >>>(dI,dR,w,h,KernelRadius);

	// Clean up
	ritkCudaSafeCall(cudaUnbindTexture(InputImageTexture));
	ritkCudaSafeCall(cudaFreeArray(InputImageArray));
}



extern "C"
void
MedianFilter(cudaArray *dInput, PixelType* dOutput, size_t OutputPitch, int w, int h, int KernelRadius)
{
	// 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, dInput));

	// A normalized output pitch allows us to operate on array indices instead of using pointer arithmetic
	int NormalizedOutputPitch = (int)OutputPitch/sizeof(float);

	dim3 DimBlock(16,16);
	dim3 DimGrid(DivUp(w,DimBlock.x),DivUp(h,DimBlock.y));
	if ( KernelRadius == 1 )
		MedianFilterKernel<1,16,16><<<DimGrid,DimBlock>>>(w,h,dOutput,NormalizedOutputPitch);
	else if ( KernelRadius == 2 )
		MedianFilterKernel<2,16,16><<<DimGrid,DimBlock>>>(w,h,dOutput,NormalizedOutputPitch);
	else if ( KernelRadius == 3 )
		MedianFilterKernel<3,16,16><<<DimGrid,DimBlock>>>(w,h,dOutput,NormalizedOutputPitch);
	else if ( KernelRadius == 4 )
		MedianFilterKernel<4,16,16><<<DimGrid,DimBlock>>>(w,h,dOutput,NormalizedOutputPitch);
	else if ( KernelRadius == 5 )
		MedianFilterKernel<5,16,16><<<DimGrid,DimBlock>>>(w,h,dOutput,NormalizedOutputPitch);
	else if ( KernelRadius == 6 )
		MedianFilterKernel<6,16,16><<<DimGrid,DimBlock>>>(w,h,dOutput,NormalizedOutputPitch);
		
	/*dim3 DimGrid(25,25);
	dim3 DimBlock(8,8);
	if ( KernelRadius == 1 )
		MedianFilterKernel<1,8,8,200,200><<<DimGrid,DimBlock>>>(dOutput);
	else if ( KernelRadius == 2 )
		MedianFilterKernel<2,8,8,200,200><<<DimGrid,DimBlock>>>(dOutput);
	else if ( KernelRadius == 3 )
		MedianFilterKernel<3,8,8,200,200><<<DimGrid,DimBlock>>>(dOutput);
	else if ( KernelRadius == 4 )
		MedianFilterKernel<4,8,8,200,200><<<DimGrid,DimBlock>>>(dOutput);
	else if ( KernelRadius == 5 )
		MedianFilterKernel<5,8,8,200,200><<<DimGrid,DimBlock>>>(dOutput);
	else if ( KernelRadius == 6 )
		MedianFilterKernel<6,8,8,200,200><<<DimGrid,DimBlock>>>(dOutput);*/

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