[u/mrichter/AliRoot.git] / HLT / TPCLib / tracking-ca / AliHLTTPCCADef.h

//-*- Mode: C++ -*-

//* This file is property of and copyright by the ALICE HLT Project        *
//* ALICE Experiment at CERN, All rights reserved.                         *
//* See cxx source for full Copyright notice                               *

#ifndef ALIHLTTPCCADEF_H
#define ALIHLTTPCCADEF_H


/**
 * Definitions needed for AliHLTTPCCATracker
 *
 */

//#define HLTCA_STANDALONE // compilation w/o root
#define HLTCA_INTERNAL_PERFORMANCE

#ifdef __CUDACC__
#define HLTCA_GPUCODE
#endif //__CUDACC__

#ifdef WIN32
#ifndef R__WIN32
#define R__WIN32
#endif //!R__WIN32
#endif //WIN32

#ifdef R__WIN32
#ifdef INTEL_RUNTIME
#pragma warning(disable : 1786)
#pragma warning(disable : 1478)
#pragma warning(disable : 161)
#pragma warning(disable : 94)
#pragma warning(disable : 1229)
#endif //INTEL_RUNTIME

#ifdef VSNET_RUNTIME
#pragma warning(disable : 4616)
#pragma warning(disable : 4996)
#pragma warning(disable : 1684)
#endif //VSNET_RUNTIME
#endif //R__WIN32

#ifdef HLTCA_STANDALONE

// class TObject{};

#ifdef ClassDef
#undef ClassDef
#endif //ClassDef

#ifdef ClassTmp
#undef ClassTmp
#endif //ClassTmp

#define ClassDef(name,id)
#define ClassImp(name)

typedef unsigned char  UChar_t;     //Unsigned Character 1 byte (unsigned char)
#ifdef R__B64    // Note: Long_t and ULong_t are currently not portable types
typedef int            Seek_t;      //File pointer (int)
typedef long           Long_t;      //Signed long integer 8 bytes (long)
typedef unsigned long  ULong_t;     //Unsigned long integer 8 bytes (unsigned long)
#else
typedef int            Seek_t;      //File pointer (int)
typedef long           Long_t;      //Signed long integer 4 bytes (long)
typedef unsigned long  ULong_t;     //Unsigned long integer 4 bytes (unsigned long)
#endif //R__B64
typedef float          Float16_t;   //Float 4 bytes written with a truncated mantissa
typedef double         Double32_t;  //Double 8 bytes in memory, written as a 4 bytes float
typedef char           Text_t;      //General string (char)
typedef unsigned char  Byte_t;      //Byte (8 bits) (unsigned char)
typedef short          Version_t;   //Class version identifier (short)
typedef const char     Option_t;    //Option string (const char)
typedef int            Ssiz_t;      //String size (int)
typedef float          Real_t;      //TVector and TMatrix element type (float)
#if defined(R__WIN32) && !defined(__CINT__)
typedef __int64          Long64_t;  //Portable signed long integer 8 bytes
typedef unsigned __int64 ULong64_t; //Portable unsigned long integer 8 bytes
#else
typedef long long          Long64_t; //Portable signed long integer 8 bytes
typedef unsigned long long ULong64_t;//Portable unsigned long integer 8 bytes
#endif //R__WIN32 && !__CINT__
typedef double         Axis_t;      //Axis values type (double)
typedef double         Stat_t;      //Statistics type (double)
typedef short          Font_t;      //Font number (short)
typedef short          Style_t;     //Style number (short)
typedef short          Marker_t;    //Marker number (short)
typedef short          Width_t;     //Line width (short)
typedef short          Color_t;     //Color number (short)
typedef short          SCoord_t;    //Screen coordinates (short)
typedef double         Coord_t;     //Pad world coordinates (double)
typedef float          Angle_t;     //Graphics angle (float)
typedef float          Size_t;      //Attribute size (float)

#define TRACKER_KEEP_TEMPDATA

#else

#include "Rtypes.h"
#include "AliHLTDataTypes.h"

namespace AliHLTTPCCADefinitions
{
  extern const AliHLTComponentDataType fgkTrackletsDataType;
  extern const AliHLTComponentDataType fgkCompressedInputDataType;
}

#endif //HLTCA_STANDALONE

#define EXTERN_ROW_HITS
#define TRACKLET_SELECTOR_MIN_HITS 29

#define GLOBAL_TRACKING_RANGE 45					//Number of rows from the upped/lower limit to search for global track candidates in for
#define GLOBAL_TRACKING_Y_RANGE_UPPER_LEFT 0.85		//Inner portion of y-range in slice that is not used in searching for global track candidates
#define GLOBAL_TRACKING_Y_RANGE_LOWER_LEFT 0.85
#define GLOBAL_TRACKING_Y_RANGE_UPPER_RIGHT 0.85
#define GLOBAL_TRACKING_Y_RANGE_LOWER_RIGHT 0.85
//#define GLOBAL_TRACKING_ONLY_UNASSIGNED_HITS		//Only use unassigned clusters in the global tracking step
//#define GLOBAL_TRACKING_EXTRAPOLATE_ONLY			//Do not update the track parameters with new hits from global tracing
#define GLOBAL_TRACKING_MIN_ROWS 10					//Min num of rows an additional global track must span over
#define GLOBAL_TRACKING_MIN_HITS 8					//Min num of hits for an additional global track

#define REPRODUCIBLE_CLUSTER_SORTING

#ifdef HLTCA_GPUCODE
#define ALIHLTTPCCANEIGHBOURS_FINDER_MAX_NNEIGHUP 6
#define ALIHLTTPCCANEIGHBOURS_FINDER_MAX_FGRIDCONTENTUPDOWN 1000
#define ALIHLTTPCCASTARTHITSFINDER_MAX_FROWSTARTHITS 3500
#define ALIHLTTPCCATRACKLET_CONSTRUCTOR_TEMP_MEM 1536					//Max amount of hits in a row that can be stored in shared memory, make sure this is divisible by ROW ALIGNMENT
#else
#define ALIHLTTPCCANEIGHBOURS_FINDER_MAX_NNEIGHUP 20
#define ALIHLTTPCCANEIGHBOURS_FINDER_MAX_FGRIDCONTENTUPDOWN 7000
#define ALIHLTTPCCASTARTHITSFINDER_MAX_FROWSTARTHITS 10000
#define ALIHLTTPCCATRACKLET_CONSTRUCTOR_TEMP_MEM 15000
#endif //HLTCA_GPUCODE

#ifdef HLTCA_GPUCODE
#define GPUdi() __device__ inline
#define GPUhdi() __host__ __device__ inline
#define GPUd() __device__
#define GPUi() inline
#define GPUhd() __host__ __device__
#define GPUh() __host__ inline
#define GPUg() __global__

#define GPUshared() __shared__
#define GPUsync() __syncthreads()

#else
#define GPUdi()
#define GPUhdi()

#define GPUd()
#define GPUi()
#define GPUhd()
#define GPUg()
#define GPUh()
#define GPUshared()
#define GPUsync()

struct float4 { float x, y, z, w; };
struct float2 { float x; float y; };
struct uchar2 { unsigned char x, y; };
struct short2 { short x, y; };
struct ushort2 { unsigned short x, y; };
struct int2 { int x, y; };
struct int3 { int x, y, z; };
struct int4 { int x, y, z, w; };
struct uint1 { unsigned int x; };
struct uint2 { unsigned int x, y; };
struct uint3 { unsigned int x, y, z; };
struct uint4 { unsigned int x, y, z, w; };
struct uint16 { unsigned int x[16]; };

#ifdef R__WIN32
#include <float.h>

inline bool finite(float x)
{
	return(x <= FLT_MAX);
}
#endif //R__WIN32

#endif //HLTCA_GPUCODE

/*
 * Helper for compile-time verification of correct API usage
 */

#ifndef HLTCA_GPUCODE
namespace
{
  template<bool> struct HLTTPCCA_STATIC_ASSERT_FAILURE;
  template<> struct HLTTPCCA_STATIC_ASSERT_FAILURE<true> {};
}

#define HLTTPCCA_STATIC_ASSERT_CONCAT_HELPER(a, b) a##b
#define HLTTPCCA_STATIC_ASSERT_CONCAT(a, b) HLTTPCCA_STATIC_ASSERT_CONCAT_HELPER(a, b)
#define STATIC_ASSERT(cond, msg) \
  typedef HLTTPCCA_STATIC_ASSERT_FAILURE<cond> HLTTPCCA_STATIC_ASSERT_CONCAT(_STATIC_ASSERTION_FAILED_##msg, __LINE__); \
  HLTTPCCA_STATIC_ASSERT_CONCAT(_STATIC_ASSERTION_FAILED_##msg, __LINE__) Error_##msg; \
  (void) Error_##msg
#else
#define STATIC_ASSERT(a, b)
#endif //!HLTCA_GPUCODE

namespace
{
  template<typename T1>
  void UNUSED_PARAM1( const T1 & ) {}
  template<typename T1, typename T2>
  void UNUSED_PARAM2( const T1 &, const T2 & ) {}
  template<typename T1, typename T2, typename T3>
  void UNUSED_PARAM3( const T1 &, const T2 &, const T3 & ) {}
  template<typename T1, typename T2, typename T3, typename T4>
  void UNUSED_PARAM4( const T1 &, const T2 &, const T3 &, const T4 & ) {}
  template<typename T1, typename T2, typename T3, typename T4, typename T5>
  void UNUSED_PARAM5( const T1 &, const T2 &, const T3 &, const T4 &, const T5 & ) {}
  template<typename T1, typename T2, typename T3, typename T4, typename T5, typename T6>
  void UNUSED_PARAM6( const T1 &, const T2 &, const T3 &, const T4 &, const T5 &, const T6 & ) {}
  template<typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7>
  void UNUSED_PARAM7( const T1 &, const T2 &, const T3 &, const T4 &, const T5 &, const T6 &, const T7 & ) {}
  template<typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8>
  void UNUSED_PARAM8( const T1 &, const T2 &, const T3 &, const T4 &, const T5 &, const T6 &, const T7 &, const T8 & ) {}
  template<typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8, typename T9>
  void UNUSED_PARAM9( const T1 &, const T2 &, const T3 &, const T4 &, const T5 &, const T6 &, const T7 &, const T8 &, const T9 & ) {}
}

#define UNROLL2(var, code) code;var++;code;var++;
#define UNROLL4(var, code) UNROLL2(var, code) UNROLL2(var, code)
#define UNROLL8(var, code) UNROLL4(var, code) UNROLL4(var, code)
#define UNROLL16(var, code) UNROLL8(var, code) UNROLL8(var, code)
#define UNROLL32(var, code) UNROLL16(var, code) UNROLL16(var, code)

#endif //ALIHLTTPCCADEF_H
Commit	Line	Data
00d07bcd	1	//-- Mode: C++ --
00d07bcd	2
fbb9b71b	3	//* This file is property of and copyright by the ALICE HLT Project *
00d07bcd	4	//* ALICE Experiment at CERN, All rights reserved. *
	5	//* See cxx source for full Copyright notice *
	6
	7	#ifndef ALIHLTTPCCADEF_H
	8	#define ALIHLTTPCCADEF_H
	9
fbb9b71b	10
00d07bcd	11	/**
	12	* Definitions needed for AliHLTTPCCATracker
	13	*
fbb9b71b	14	*/
fbb9b71b	15
00d07bcd	16	//#define HLTCA_STANDALONE // compilation w/o root
ce565086	17	#define HLTCA_INTERNAL_PERFORMANCE
00d07bcd	18
00d07bcd	19	#ifdef __CUDACC__
00d07bcd	20	#define HLTCA_GPUCODE
31649d4b	21	#endif //__CUDACC__
00d07bcd	22
7be9b0d7	23	#ifdef WIN32
	24	#ifndef R__WIN32
	25	#define R__WIN32
31649d4b	26	#endif //!R__WIN32
31649d4b	27	#endif //WIN32
7be9b0d7	28
31649d4b	29	#ifdef R__WIN32
73a33d2e	30	#ifdef INTEL_RUNTIME
	31	#pragma warning(disable : 1786)
	32	#pragma warning(disable : 1478)
	33	#pragma warning(disable : 161)
f0bada7f	34	#pragma warning(disable : 94)
f0bada7f	35	#pragma warning(disable : 1229)
31649d4b	36	#endif //INTEL_RUNTIME
73a33d2e	37
	38	#ifdef VSNET_RUNTIME
	39	#pragma warning(disable : 4616)
	40	#pragma warning(disable : 4996)
	41	#pragma warning(disable : 1684)
31649d4b	42	#endif //VSNET_RUNTIME
31649d4b	43	#endif //R__WIN32
00d07bcd	44
31649d4b	45	#ifdef HLTCA_STANDALONE
00d07bcd	46
ce565086	47	// class TObject{};
00d07bcd	48
b22af1bf	49	#ifdef ClassDef
b22af1bf	50	#undef ClassDef
31649d4b	51	#endif //ClassDef
b22af1bf	52
	53	#ifdef ClassTmp
	54	#undef ClassTmp
31649d4b	55	#endif //ClassTmp
b22af1bf	56
00d07bcd	57	#define ClassDef(name,id)
	58	#define ClassImp(name)
	59
00d07bcd	60	typedef unsigned char UChar_t; //Unsigned Character 1 byte (unsigned char)
00d07bcd	61	#ifdef R__B64 // Note: Long_t and ULong_t are currently not portable types
	62	typedef int Seek_t; //File pointer (int)
	63	typedef long Long_t; //Signed long integer 8 bytes (long)
	64	typedef unsigned long ULong_t; //Unsigned long integer 8 bytes (unsigned long)
	65	#else
	66	typedef int Seek_t; //File pointer (int)
	67	typedef long Long_t; //Signed long integer 4 bytes (long)
	68	typedef unsigned long ULong_t; //Unsigned long integer 4 bytes (unsigned long)
31649d4b	69	#endif //R__B64
00d07bcd	70	typedef float Float16_t; //Float 4 bytes written with a truncated mantissa
00d07bcd	71	typedef double Double32_t; //Double 8 bytes in memory, written as a 4 bytes float
00d07bcd	72	typedef char Text_t; //General string (char)
00d07bcd	73	typedef unsigned char Byte_t; //Byte (8 bits) (unsigned char)
	74	typedef short Version_t; //Class version identifier (short)
	75	typedef const char Option_t; //Option string (const char)
	76	typedef int Ssiz_t; //String size (int)
	77	typedef float Real_t; //TVector and TMatrix element type (float)
	78	#if defined(R__WIN32) && !defined(__CINT__)
	79	typedef __int64 Long64_t; //Portable signed long integer 8 bytes
	80	typedef unsigned __int64 ULong64_t; //Portable unsigned long integer 8 bytes
	81	#else
	82	typedef long long Long64_t; //Portable signed long integer 8 bytes
	83	typedef unsigned long long ULong64_t;//Portable unsigned long integer 8 bytes
31649d4b	84	#endif //R__WIN32 && !__CINT__
00d07bcd	85	typedef double Axis_t; //Axis values type (double)
	86	typedef double Stat_t; //Statistics type (double)
	87	typedef short Font_t; //Font number (short)
	88	typedef short Style_t; //Style number (short)
	89	typedef short Marker_t; //Marker number (short)
	90	typedef short Width_t; //Line width (short)
	91	typedef short Color_t; //Color number (short)
	92	typedef short SCoord_t; //Screen coordinates (short)
	93	typedef double Coord_t; //Pad world coordinates (double)
	94	typedef float Angle_t; //Graphics angle (float)
	95	typedef float Size_t; //Attribute size (float)
	96
f0bada7f	97	#define TRACKER_KEEP_TEMPDATA
f0bada7f	98
fbb9b71b	99	#else
00d07bcd	100
00d07bcd	101	#include "Rtypes.h"
e1f2d1c3	102	#include "AliHLTDataTypes.h"
	103
	104	namespace AliHLTTPCCADefinitions
	105	{
	106	extern const AliHLTComponentDataType fgkTrackletsDataType;
751d16ac	107	extern const AliHLTComponentDataType fgkCompressedInputDataType;
e1f2d1c3	108	}
e1f2d1c3	109
31649d4b	110	#endif //HLTCA_STANDALONE
00d07bcd	111
f0bada7f	112	#define EXTERN_ROW_HITS
e4818148	113	#define TRACKLET_SELECTOR_MIN_HITS 29
	114
	115	#define GLOBAL_TRACKING_RANGE 45 //Number of rows from the upped/lower limit to search for global track candidates in for
	116	#define GLOBAL_TRACKING_Y_RANGE_UPPER_LEFT 0.85 //Inner portion of y-range in slice that is not used in searching for global track candidates
	117	#define GLOBAL_TRACKING_Y_RANGE_LOWER_LEFT 0.85
	118	#define GLOBAL_TRACKING_Y_RANGE_UPPER_RIGHT 0.85
	119	#define GLOBAL_TRACKING_Y_RANGE_LOWER_RIGHT 0.85
	120	//#define GLOBAL_TRACKING_ONLY_UNASSIGNED_HITS //Only use unassigned clusters in the global tracking step
	121	//#define GLOBAL_TRACKING_EXTRAPOLATE_ONLY //Do not update the track parameters with new hits from global tracing
	122	#define GLOBAL_TRACKING_MIN_ROWS 10 //Min num of rows an additional global track must span over
	123	#define GLOBAL_TRACKING_MIN_HITS 8 //Min num of hits for an additional global track
	124
b22af1bf	125	#define REPRODUCIBLE_CLUSTER_SORTING
b22af1bf	126
7be9b0d7	127	#ifdef HLTCA_GPUCODE
b22af1bf	128	#define ALIHLTTPCCANEIGHBOURS_FINDER_MAX_NNEIGHUP 6
b22af1bf	129	#define ALIHLTTPCCANEIGHBOURS_FINDER_MAX_FGRIDCONTENTUPDOWN 1000
7be9b0d7	130	#define ALIHLTTPCCASTARTHITSFINDER_MAX_FROWSTARTHITS 3500
b22af1bf	131	#define ALIHLTTPCCATRACKLET_CONSTRUCTOR_TEMP_MEM 1536 //Max amount of hits in a row that can be stored in shared memory, make sure this is divisible by ROW ALIGNMENT
7be9b0d7	132	#else
	133	#define ALIHLTTPCCANEIGHBOURS_FINDER_MAX_NNEIGHUP 20
	134	#define ALIHLTTPCCANEIGHBOURS_FINDER_MAX_FGRIDCONTENTUPDOWN 7000
	135	#define ALIHLTTPCCASTARTHITSFINDER_MAX_FROWSTARTHITS 10000
b22af1bf	136	#define ALIHLTTPCCATRACKLET_CONSTRUCTOR_TEMP_MEM 15000
31649d4b	137	#endif //HLTCA_GPUCODE
00d07bcd	138
00d07bcd	139	#ifdef HLTCA_GPUCODE
f0bada7f	140	#define GPUdi() __device__ inline
f0bada7f	141	#define GPUhdi() __host__ __device__ inline
fbb9b71b	142	#define GPUd() __device__
f0bada7f	143	#define GPUi() inline
fbb9b71b	144	#define GPUhd() __host__ __device__
	145	#define GPUh() __host__ inline
	146	#define GPUg() __global__
00d07bcd	147
fbb9b71b	148	#define GPUshared() __shared__
00d07bcd	149	#define GPUsync() __syncthreads()
00d07bcd	150
00d07bcd	151	#else
f0bada7f	152	#define GPUdi()
f0bada7f	153	#define GPUhdi()
00d07bcd	154
fbb9b71b	155	#define GPUd()
f0bada7f	156	#define GPUi()
00d07bcd	157	#define GPUhd()
00d07bcd	158	#define GPUg()
fbb9b71b	159	#define GPUh()
	160	#define GPUshared()
	161	#define GPUsync()
	162
f0bada7f	163	struct float4 { float x, y, z, w; };
fbb9b71b	164	struct float2 { float x; float y; };
b22af1bf	165	struct uchar2 { unsigned char x, y; };
	166	struct short2 { short x, y; };
	167	struct ushort2 { unsigned short x, y; };
	168	struct int2 { int x, y; };
	169	struct int3 { int x, y, z; };
	170	struct int4 { int x, y, z, w; };
fbb9b71b	171	struct uint1 { unsigned int x; };
b22af1bf	172	struct uint2 { unsigned int x, y; };
b22af1bf	173	struct uint3 { unsigned int x, y, z; };
fbb9b71b	174	struct uint4 { unsigned int x, y, z, w; };
f0bada7f	175	struct uint16 { unsigned int x[16]; };
00d07bcd	176
7be9b0d7	177	#ifdef R__WIN32
	178	#include <float.h>
	179
	180	inline bool finite(float x)
	181	{
	182	return(x <= FLT_MAX);
	183	}
31649d4b	184	#endif //R__WIN32
7be9b0d7	185
31649d4b	186	#endif //HLTCA_GPUCODE
00d07bcd	187
4acc2401	188	/*
	189	* Helper for compile-time verification of correct API usage
	190	*/
b22af1bf	191
b22af1bf	192	#ifndef HLTCA_GPUCODE
4acc2401	193	namespace
	194	{
	195	template<bool> struct HLTTPCCA_STATIC_ASSERT_FAILURE;
	196	template<> struct HLTTPCCA_STATIC_ASSERT_FAILURE<true> {};
	197	}
	198
	199	#define HLTTPCCA_STATIC_ASSERT_CONCAT_HELPER(a, b) a##b
	200	#define HLTTPCCA_STATIC_ASSERT_CONCAT(a, b) HLTTPCCA_STATIC_ASSERT_CONCAT_HELPER(a, b)
	201	#define STATIC_ASSERT(cond, msg) \
	202	typedef HLTTPCCA_STATIC_ASSERT_FAILURE<cond> HLTTPCCA_STATIC_ASSERT_CONCAT(_STATIC_ASSERTION_FAILED_##msg, __LINE__); \
	203	HLTTPCCA_STATIC_ASSERT_CONCAT(_STATIC_ASSERTION_FAILED_##msg, __LINE__) Error_##msg; \
	204	(void) Error_##msg
b22af1bf	205	#else
b22af1bf	206	#define STATIC_ASSERT(a, b)
31649d4b	207	#endif //!HLTCA_GPUCODE
4acc2401	208
	209	namespace
	210	{
	211	template<typename T1>
	212	void UNUSED_PARAM1( const T1 & ) {}
	213	template<typename T1, typename T2>
	214	void UNUSED_PARAM2( const T1 &, const T2 & ) {}
	215	template<typename T1, typename T2, typename T3>
	216	void UNUSED_PARAM3( const T1 &, const T2 &, const T3 & ) {}
	217	template<typename T1, typename T2, typename T3, typename T4>
	218	void UNUSED_PARAM4( const T1 &, const T2 &, const T3 &, const T4 & ) {}
	219	template<typename T1, typename T2, typename T3, typename T4, typename T5>
	220	void UNUSED_PARAM5( const T1 &, const T2 &, const T3 &, const T4 &, const T5 & ) {}
	221	template<typename T1, typename T2, typename T3, typename T4, typename T5, typename T6>
	222	void UNUSED_PARAM6( const T1 &, const T2 &, const T3 &, const T4 &, const T5 &, const T6 & ) {}
	223	template<typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7>
	224	void UNUSED_PARAM7( const T1 &, const T2 &, const T3 &, const T4 &, const T5 &, const T6 &, const T7 & ) {}
	225	template<typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8>
	226	void UNUSED_PARAM8( const T1 &, const T2 &, const T3 &, const T4 &, const T5 &, const T6 &, const T7 &, const T8 & ) {}
	227	template<typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8, typename T9>
	228	void UNUSED_PARAM9( const T1 &, const T2 &, const T3 &, const T4 &, const T5 &, const T6 &, const T7 &, const T8 &, const T9 & ) {}
	229	}
00d07bcd	230
b22af1bf	231	#define UNROLL2(var, code) code;var++;code;var++;
	232	#define UNROLL4(var, code) UNROLL2(var, code) UNROLL2(var, code)
	233	#define UNROLL8(var, code) UNROLL4(var, code) UNROLL4(var, code)
	234	#define UNROLL16(var, code) UNROLL8(var, code) UNROLL8(var, code)
	235	#define UNROLL32(var, code) UNROLL16(var, code) UNROLL16(var, code)
	236
31649d4b	237	#endif //ALIHLTTPCCADEF_H