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

	Program:		Range Imaging Toolkit (RITK)
	Module:			ritkFloatVolumeCudaRayCaster.cu
	Version:		$Revision: 2882 $
	Date:			$Date: 2012-04-02 17:00:07 +0200 (Mo, 02 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 "ritkFloatVolumeCudaRayCasterKernel.h"


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

#define NumSamples 100


/// Cuda ray casting. Will eventually invoke the kernel
//-----------------------------------------------------
extern "C" 
bool 
ritkFloatVolumeCudaRayCaster(const ritkCudaRayCasterVolumeParameters &VP, const ritkCudaRayCasterProjectionImageParameters<float> &IP, const ritkCudaRayCasterProjectionParameters &PP)
{
	// 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;
	float FixedScale = 1.f/(NumSamples*DOT3(VolumeExtent,VolumeExtent));

	// 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)IP.Size.x,DimBlock.x), DivUp((int)IP.Size.y,DimBlock.y));
	FloatVolumeCudaRayCasterKernel<NumSamples><<< DimGrid,DimBlock >>>(VP, IP, PP, WorldToTex, FixedScale);

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

	return true;
}