#ifdef _WIN32
	#define WIN32_LEAN_AND_MEAN
#endif

#include "ritkCudaMacros.h"
#include "TemporalAveragingFilterKernels.h"


// Division. If division remainder is neq zero then the result is ceiled
//----------------------------------------------------------------------------
#define DivUp(a,b) ((a % b != 0) ? (a/b + 1) : (a/b))



// CUDA interface for element-wise addition: Accu[i] += Data[i]
//----------------------------------------------------------------------------
template<class TAccu, class TInput>
void AddData(TAccu *Accu, TInput *Data, unsigned long Num)
{
	const int BlockSize = 128;
	const int NumSeq = 2;
	dim3 DimBlock(BlockSize);
	dim3 DimGrid(DivUp(Num,DimBlock.x)/NumSeq);
	AddSubDataCascadingKernel<TAccu,TInput,BlockSize,true><<<DimGrid,DimBlock>>>(Accu, Data, Num);
}


// CUDA interface for element-wise subtraction: Accu[i] -= Data[i]
//----------------------------------------------------------------------------
template<class TAccu, class TInput>
void SubData(TAccu *Accu, TInput *Data, unsigned long Num)
{
	const int BlockSize = 128;
	const int NumSeq = 2;
	dim3 DimBlock(BlockSize);
	dim3 DimGrid(DivUp(Num,DimBlock.x)/NumSeq);
	AddSubDataCascadingKernel<TAccu,TInput,BlockSize,false><<<DimGrid,DimBlock>>>(Accu, Data, Num);
}


// CUDA interface for scalar multiplication: Output[i] = Scalar*Input[i]
//----------------------------------------------------------------------------
template<class TOutput, class TInput>
void ScalarMulData(TOutput *Output, TInput *Input, float Scalar, unsigned long Num)
{
	const int BlockSize = 128;
	dim3 DimBlock(BlockSize);
	dim3 DimGrid(DivUp(Num,DimBlock.x));
	ScalarMulDataKernel<TOutput,TInput><<<DimGrid,DimBlock>>>(Output, Input, Scalar, Num);
}


// CUDA interface for O(1) temporal averaging
//----------------------------------------------------------------------------
template<class TOutput, class TAccu, class TInput, bool uchardummy>
void SlidingWindowTemporalAveraging(TOutput *Output, TAccu *Accu, TInput *AddData, TInput *SubData, float Scalar, unsigned long Num)
{
	const int BlockSize = 128;
	const int NumCascades = 4;
	dim3 DimBlock(BlockSize);
	dim3 DimGrid(DivUp(Num,DimBlock.x)/NumCascades);
	SlidingWindowTemporalAveragingKernel<TOutput,TAccu,TInput,BlockSize,uchardummy><<<DimGrid,DimBlock>>>(Output, Accu, AddData, SubData, Scalar, Num);
}



// Instantiations
//----------------------------------------------------------------------------
void AddData(float* Accu, float *Data, unsigned long Num)
{
	AddData<float,float>(Accu,Data,Num);
}
void AddData(float3 *Accu, uchar3 *Data, unsigned long Num)
{
	AddData<float3,uchar3>(Accu,Data,Num);
}


void SubData(float *Accu, float *Data, unsigned long Num)
{
	SubData<float,float>(Accu,Data,Num);
}
void SubData(float3 *Accu, uchar3 *Data, unsigned long Num)
{
	SubData<float3,uchar3>(Accu,Data,Num);
}


void ScalarMulData(float *Output, float *Input, float Scalar, unsigned long Num)
{
	ScalarMulData<float,float>(Output,Input,Scalar,Num);
}
void ScalarMulData(uchar3 *Output, float3 *Input, float Scalar, unsigned long Num)
{
	ScalarMulData<uchar3,float3>(Output,Input,Scalar,Num);
}


void SlidingWindowTemporalAveraging(float *Output, float *Accu, float *AddData, float *SubData, float Scalar, unsigned long Num)
{
	SlidingWindowTemporalAveraging<float,float,float,false>(Output,Accu,AddData,SubData,Scalar,Num);
}
void SlidingWindowTemporalAveraging(uchar3 *Output, float3 *Accu, uchar3 *AddData, uchar3 *SubData, float Scale, unsigned long Num)
{
	SlidingWindowTemporalAveraging<uchar3,float3,uchar3,true>(Output,Accu,AddData,SubData,Scale,Num);
}



