
__device__ inline float3
operator+(const float3 &a, const uchar3 &b)
{
	float3 result;
	result.x = a.x + b.x;
	result.y = a.y + b.y;
	result.z = a.z + b.z;
	return result;
}


__device__ inline void
operator+=(float3 &a, const uchar3 &b)
{
	a.x += b.x;
	a.y += b.y;
	a.z += b.z;
}


__device__ inline float3
operator-(const float3 &a, const uchar3 &b)
{
	float3 result;
	result.x = a.x - b.x;
	result.y = a.y - b.y;
	result.z = a.z - b.z;
	return result;
}

__device__ inline void
operator-=(float3 &a, const uchar3 &b)
{
	a.x -= b.x;
	a.y -= b.y;
	a.z -= b.z;
}


__device__ inline float3
operator*(const float3 &a, float f)
{
	float3 result;
	result.x = a.x * f;
	result.y = a.y * f;
	result.z = a.z * f;
	return result;
}

__device__ inline float3
operator*(const uchar3 &a, float f)
{
	float3 result;
	result.x = a.x * f;
	result.y = a.y * f;
	result.z = a.z * f;
	return result;
}


__device__ inline void
operator*=(float3 &a, float f)
{
	a.x *= f;
	a.y *= f;
	a.z *= f;
}

__device__ inline uchar3
mycast(const uchar3 &outputdummy, const float3 operand)
{
	uchar3 result;
	result.x = operand.x;
	result.y = operand.y;
	result.z = operand.z;
	return result;
}

__device__ inline float
mycast(const float &outputdummy, const float operand)
{
	float result = operand;
	return result;
}



//----------------------------------------------------------------------------
template<bool AdditionMode>
__global__ void
AddSubDataKernel(float *Accu, float *Data, long Num)
{
	// Global index
	int i = blockIdx.x*blockDim.x + threadIdx.x;

	// Boundary check
	if ( i >= Num )
		return;

	// Pull data
	float f1 = Accu[i];
	float f2 = Data[i];

	// Add or sub
	if ( AdditionMode )
		f1 += f2;
	else
		f1 -= f2;

	// Write output
	Accu[i] = f1;
}


//----------------------------------------------------------------------------
template<class TAccu, class TInput, unsigned int BlockSize, bool AdditionMode>
__global__ void
AddSubDataCascadingKernel(TAccu *Accu, TInput *Data, long Num)
{
	// Global index
	unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
	unsigned int GridSize = BlockSize*gridDim.x;

	// Cascading
	while ( i < Num )
	{
		TAccu f1 = Accu[i];
		TInput f2 = Data[i];
		if ( AdditionMode )
			f1 += f2;
		else
			f1 -= f2;
		Accu[i] = f1;

		i += GridSize;
	}
}


//----------------------------------------------------------------------------
template<class TOutput, class TInput>
__global__ void
ScalarMulDataKernel(TOutput *Output, TInput *Input, float Scalar, unsigned long Num)
{
	// Global index
	int i = blockIdx.x*blockDim.x + threadIdx.x;

	// Boundary check
	if ( i >= Num )
		return;

	// Pull data
	TOutput dummy;
	TInput f1 = Input[i];
	TOutput f2 = mycast(dummy, f1 * Scalar);

	// Write output
	Output[i] = f2;
}


//----------------------------------------------------------------------------
template<class TOutput, class TAccu, class TInput, unsigned int BlockSize, bool uchardummy>
__global__ void
SlidingWindowTemporalAveragingKernel(TOutput *Output, TAccu *Accu, TInput *AddData, TInput *SubData, float Scale, unsigned long Num)
{
	// Global index
	unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
	unsigned int GridSize = BlockSize*gridDim.x;

	// Cascading
	while ( i < Num )
	{
		// Pull data
		TAccu f1 = Accu[i];
		TInput f2 = AddData[i];
		TInput f3 = SubData[i];

		// Sliding window (recursive)
		TAccu f4 = f1 + f2 - f3;

		// Scaling
		TAccu f5 = f4*Scale;

		// Write output
		TOutput dummy;
		Output[i] = mycast(dummy,f5);
		Accu[i] = f4;
	
		// Update index
		i += GridSize;
	}
}

