#ifdef _WIN32
	#define WIN32_LEAN_AND_MEAN
#endif

#include "CudaKinectRangeToWorldKernel.h"
#include "ritkCudaMacros.h"
#include "stdio.h"

//#define RUNTIME_EVALUATION

extern "C"
float
CUDARangeToWorldKernel(const cudaArray *InputImageArray, float3 *DeviceOutput, int w, int h, float X2Z, float Y2Z)
{
	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));

	// 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
	
	// Kernel Invocation
	dim3 DimBlock(16,16);
	dim3 DimGrid(DivUp(w,DimBlock.x),DivUp(h,DimBlock.y));
	CUDAKinectRangeToWorldKernel<16,16><<<DimGrid,DimBlock>>>(DeviceOutput, w, h, X2Z, Y2Z);

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

#endif

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

	return ElapsedTime;
}