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

	Program:		Range Imaging Toolkit (RITK)
	Module:			ritkCudaContext.h
	Version:		$Revision: 883 $
	Date:			$Date: 2011-04-28 11:38:06 +0200 (Do, 28 Apr 2011) $

	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

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

#ifndef CUDACONTEXT_H__
#define CUDACONTEXT_H__

#include <cuda_runtime.h>

#include "ritkThreadedWorker.h"
#include "ritkCuda.h"
#include "ritkCudaJob.h"
#include "ritkCudaMacros.h"



namespace ritk
{

/**	@ingroup	Cuda
 *	@class		CudaContext
 *	@author		Jakob Wasza
 *	@brief		Class to access the CUDA context
 *
 *	@details
 *	This class provides the basic concept for the CUDA context within the RITK framework. \n
 *	In CUDA, it is not possible for one (CPU) thread to directly access device resources that were created
 *	from within another thread. The reason for that is that a thread has its own unique CUDA context and CUDA
 *	ressources/calls are not valid across different contexts. One solution to this problem is to define a global
 *	CUDA context that is pushed/popped on a thread's context stack. However, this is atm. only possible for the
 *	driver API and, if functionality from the runtim API is used, will probably lead to undefined behaviour. 
 *	Besides the incompability of driver and runtime API, CUDA context switches are reported to be rather slow.
 *	Therefore, a single global thread with associated CUDA context is provided. All CUDA calls should be performed
 *	within this thread. To execute a certain CUDA task within this global thread the task has to be wrapped into
 *	a special RITKCudaJob and passed to the ExecuteCudaJob() method that eventually will ensure that the job
 *	is processed within the global CUDA thread. As manually wrapping a CUDA task in such a job and passing it to
 *	the CudaContext is rather inconvenient the RITKCuda-module provides macros easing CUDA handling: \n
 *		-	CudaMethodMacro: to declare a class' member function to be processed within the global CUDA thread \n
 *		-	CudaMethodBlockedCallMacro: to perform a blocking execution of the member function in the global CUDA thread \n
 *		-	CudaBlockedConstMethodCallMacro: same as CudaMethodBlockedCallMacro but with const qualifier \n

 *	The macros can be used as illustrated in the following code snippet:
 *	\code
 * class MyClass
 * {
 *  public:
 *   // Constructor,destructor normal member functions
 *   ...
 *
 *   /// A method that is called within an arbitrary thread and has to perform a CUDA task (in this case memory allocation)
 *   void CUDAMethod(long NumBytesToAllocate);
 *
 *   /// A method that will perform a CUDA task in the global CUDA thread
 *   CudaMethodMacro(MyClass,ActualCUDAMethod);
 *
 *   // Other stuff
 *   ...
 *
 *   /// Member holding the number of bytes to allocate
 *   long m_NumBytesToAllocate;
 * };
 *
 * 
 * void
 * MyClass::CUDAMethod(long NumBytesToAllocate)
 * {
 *  // Set the member
 *  m_NumBytesToAllocate = NumBytesToAllocate;
 *
 *  // Call actual allocation routine within the global CUDA thread
 *  CudaBlockedMethodCallMacro(ActualCUDAMethod);
 * }
 *
 *
 *
 * void
 * MyClass::ActualCUDAMethod()
 * {
 *  // Simply allocate the memory using the member
 *  float *DevPtr;
 *  ritkCudaSafeCall(cudaMalloc((void*)&DevPtr,m_NumBytesToAllocate));
 * }
 *	\endcode
 *	Please note the method signature of the actual CUDA method. MUST have NO arguments and MUST return void! \n
 *	From a technical point of view these macros implement a callback design pattern: A special RITKCudaJob is
 *	created and its callback function set to an automatically generated static member function of the class. This
 *	static callback function is then called within the global CUDA thread with the argument specifying the object 
 *	(this pointer) that invoked the RITKCudaBlockedMethodCallMacro. The parameter of this macro specifies the 
 *	actual method to be called. \n
 *	Note that the CudaBlockedMethodCallMacro can be called in a recursive manner (e.g. deadlock prevention is automatically
 *	turned on)
 */
class CUDA_EXPORT CudaContext
{
public:
	/// Singleton access
	static inline CudaContext& GetInstance()
	{
		static CudaContext Instance;
		return Instance;
	}

	/// Execute a single cuda job
	//@{
	void ExecuteCudaJob(CudaJob *Job, bool Block=true);
	//@}

protected:
	/// Worker that runs in the cuda main thread
	ThreadedWorker *m_ThreadedWorker;

private:
	/// Private constructor due to the singleton design pattern
	CudaContext();
	/// Private destructo due to the singleton design pattern
	~CudaContext();
};


} // end namespace ritk


#define CudaMethodMacro(ClassName,MethodName)									\
	void MethodName();															\
	static inline void CudaCallback_##MethodName (void* ptr)					\
	{																			\
		ClassName *ClassPtr = dynamic_cast<ClassName*>((ClassName*)ptr);		\
		if ( !ClassPtr )														\
		{																		\
			std::stringstream ss;												\
			ss << "Could not cast pointer to " << #ClassName;					\
			std::cout << ss.str() << std::endl;									\
			throw std::exception();												\
		}																		\
		ClassPtr->MethodName();													\
	}																			\
	ritk::CudaJob m_##MethodName##_CudaJob;

#define CudaConstMethodMacro(ClassName,MethodName)											\
	void MethodName() const;																\
	static inline void CudaCallbackConst_##MethodName (const void *ptr)						\
	{																						\
		const ClassName *ClassPtr = dynamic_cast<const ClassName*>((const ClassName*)ptr);	\
		if ( !ClassPtr )																	\
		{																					\
			std::stringstream ss;															\
			ss << "Could not cast pointer to " << #ClassName;								\
			std::cout << ss.str() << std::endl;												\
			throw std::exception();															\
		}																					\
		ClassPtr->MethodName();																\
	}																						\
	mutable ritk::CudaJob m_Const##MethodName##_CudaJob;


#define CudaBlockedMethodCallMacro(MethodName)												\
{																							\
	m_##MethodName##_CudaJob.SetCallbackFunction(this, CudaCallback_##MethodName);			\
	ritk::CudaContext::GetInstance().ExecuteCudaJob(&m_##MethodName##_CudaJob);				\
}

#define CudaBlockedConstMethodCallMacro(MethodName)													\
{																									\
	m_Const##MethodName##_CudaJob.SetConstCallbackFunction(this, CudaCallbackConst_##MethodName );	\
	ritk::CudaContext::GetInstance().ExecuteCudaJob(&m_Const##MethodName##_CudaJob);				\
}



#endif // CUDACONTEXT_H__
