/*======================================================================================

	Program:		Range Imaging Toolkit (RITK)
	Module:			ritkTruncatedSignedDistanceFunctionRayCaster.cu
	Version:		$Revision: 2965 $
	Date:			$Date: 2012-04-10 18:13:53 +0200 (Di, 10 Apr 2012) $

	Copyright (c) 2010-2011
	Pattern Recognition Lab, 
    University Erlangen-Nuremberg, Germany
    All rights reserved.

	This program is distributed in the hope that it will be useful,
    but WITHOUT ANY WARRANTY; without even the implied warranty of
    MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. 
	See RITKCopyright.txt for more information.

	http://www5.informatik.uni-erlangen.de/research/software/ritk

========================================================================================*/

// Due to MS Windows SDK and debug built
#ifdef _WIN32
	#define WIN32_LEAN_AND_MEAN
#endif

#include "ritkCudaMacros.h"

#include "ritkCudaRayCasterInterfaces.h"
#include "ritkTruncatedSignedDistanceFunctionRayCasterInterfaces.h"
#include "ritkTruncatedSignedDistanceFunctionRayCasterKernels.h"

#define DivUp(a,b) (((a) % (b) != 0) ? ((a)/(b) + 1) : ((a)/(b)))



//----------------------------------------------------------------------------
void
ritkTSDFRayCaster(const ritkCudaRayCasterVolumeParameters &VP, const ritkCudaRayCasterProjectionImageParameters<float> &RIP, const ritkCudaRayCasterProjectionImageParameters<float3> &WIP, const ritkCudaRayCasterProjectionParameters &PP, float CutoffValue, float InvalidValue)
{
	// Sanity checks
	if ( RIP.Size.x != WIP.Size.x || RIP.Size.y != WIP.Size.y )
	{
		printf("[%s] Mismatch between size of range image and normal image detected: [%d,%d] vs. [%d,%d]\n", __FUNCTION__, RIP.Size.x, RIP.Size.y, WIP.Size.x, WIP.Size.y);
		return;
	}
	if ( fabs(RIP.Spacing.x-WIP.Spacing.y) > 1e-6 || fabs(RIP.Spacing.y-WIP.Spacing.y) > 1e-6 )
	{
		printf("[%s] Mismatch between spacing of range image and normal image detected: [%f,%f] vs. [%f,%f]\n", __FUNCTION__, RIP.Spacing.x, RIP.Spacing.y, WIP.Spacing.x, WIP.Spacing.y);
		return;
	}
	if ( fabs(RIP.Origin.x-WIP.Origin.y) > 1e-6 || fabs(RIP.Origin.y-WIP.Origin.y) > 1e-6 )
	{
		printf("[%s] Mismatch between origin of range image and normal image detected: [%f,%f] vs. [%f,%f]\n", __FUNCTION__, RIP.Origin.x, RIP.Origin.y, WIP.Origin.x, WIP.Origin.y);
		return;
	}

	// Set up the volume texture
	VolumeTexture.normalized = true; 
	VolumeTexture.filterMode = cudaFilterModeLinear;
	VolumeTexture.addressMode[0] = cudaAddressModeClamp;
	VolumeTexture.addressMode[1] = cudaAddressModeClamp;
	ritkCudaSafeCall(cudaBindTextureToArray(&VolumeTexture, VP.DataArray, VP.ChannelDesc));

	// Set up the fixed scale
	float3 VolumeExtent = VP.Spacing*VP.Size;

	// Set up the world to texture converions
	float3 WorldToTex = make_float3(1.f)/(VP.Spacing*VP.Size);

	// Kernel invocation
	const unsigned int BLOCKSIZE_X = 8;
	const unsigned int BLOCKSIZE_Y = 8;
	const dim3 DimBlock(BLOCKSIZE_X, BLOCKSIZE_Y);
	const dim3 DimGrid(DivUp((int)RIP.Size.x,DimBlock.x), DivUp((int)RIP.Size.y,DimBlock.y));
	TSDFRayCastKernel<<< DimGrid,DimBlock >>>(VP, RIP, WIP, PP, WorldToTex, CutoffValue/4.f);

	// Release textures
	ritkCudaSafeCall(cudaUnbindTexture(VolumeTexture));
}