[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

#if defined(__CUDACC__) || defined(__OPENCL__)
#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

#if defined(HLTCA_STANDALONE) || (defined(HLTCA_GPUCODE) && defined(__OPENCL__) && !defined(HLTCA_HOSTCODE))

// 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(__OPENCL__) && !defined(HLTCA_HOSTCODE)
typedef long Long64_t;
typedef unsigned long ULong64_t;
#else
#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
#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
#ifdef __OPENCL__
#ifdef HLTCA_HOSTCODE //HLTCA_GPUCODE & __OPENCL & HLTCA_HOSTCODE

#define GPUdi() //TRIGGER_ERROR_NO_DEVICE_CODE
#define GPUhdi() inline
#define GPUd() //TRIGGER_ERROR_NO_DEVICE_CODE
#define GPUi() inline
#define GPUhd() 
#define GPUh() 
#define GPUg() TRIGGER_ERROR_NO_DEVICE_CODE
#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]; };

#else //HLTCA_HOSTCODE : HLTCA_GPUCODE & __OPENCL__ & !HLTCA_HOSTCODE

#define GPUdi() inline
#define GPUhdi() inline
#define GPUd() 
#define GPUi() inline
#define GPUhd() 
#define GPUh() TRIGGER_ERROR_NO_HOST_CODE
#define GPUg() __kernel
#define GPUshared() __local
#define GPUsync() barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE)

#endif //HLTCA_HOSTCODE

#else //__OPENCL__ : HLTCA_GPUCODE & !__OPENCL__

#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()

#endif //__OPENCL
#else //HLTCA_GPUCODE : !HLTCA_GPUCODE

#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

#if defined(__OPENCL__) && !defined(HLTCA_HOSTCODE)
#define GPUsharedref() GPUshared()
#define GPUglobalref() __global
//#define GPUconstant() __constant //Replace __constant by __global (possibly add const __restrict where possible later!)
#define GPUconstant() GPUglobalref()
#else
#define GPUconstant()
#define GPUsharedref()
#define GPUglobalref()
#endif

enum LocalOrGlobal { Mem_Local, Mem_Global, Mem_Constant, Mem_Plain };
#if defined(__OPENCL__) && !defined(HLTCA_HOSTCODE)
template<LocalOrGlobal, typename L, typename G, typename C, typename P> struct MakeTypeHelper;
template<typename L, typename G, typename C, typename P> struct MakeTypeHelper<Mem_Local, L, G, C, P> { typedef L type; };
template<typename L, typename G, typename C, typename P> struct MakeTypeHelper<Mem_Global, L, G, C, P> { typedef G type; };
template<typename L, typename G, typename C, typename P> struct MakeTypeHelper<Mem_Constant, L, G, C, P> { typedef C type; };
template<typename L, typename G, typename C, typename P> struct MakeTypeHelper<Mem_Plain, L, G, C, P> { typedef P type; };
#define MakeType(base_type) typename MakeTypeHelper<LG, GPUshared() base_type, GPUglobalref() base_type, GPUconstant() base_type, base_type>::type 
#define MEM_CLASS_PRE() template<LocalOrGlobal LG>
#define MEM_LG(type) type<LG>
#define MEM_CLASS_PRE2() template<LocalOrGlobal LG2>
#define MEM_LG2(type) type<LG2>
#define MEM_CLASS_PRE12() template<LocalOrGlobal LG> template<LocalOrGlobal LG2>
#define MEM_CLASS_PRE23() template<LocalOrGlobal LG2, LocalOrGlobal LG3>
#define MEM_LG3(type) type<LG3>
#define MEM_CLASS_PRE234() template<LocalOrGlobal LG2, LocalOrGlobal LG3, LocalOrGlobal LG4>
#define MEM_LG4(type) type<LG4>
#define MEM_GLOBAL(type) type<Mem_Global>
#define MEM_LOCAL(type) type<Mem_Local>
#define MEM_CONSTANT(type) type<Mem_Global>
#define MEM_PLAIN(type) type<Mem_Plain>
#define MEM_TEMPLATE() template <typename T>
#define MEM_TYPE(type) T
#define MEM_TEMPLATE2() template <typename T, typename T2>
#define MEM_TYPE2(type) T2
#define MEM_TEMPLATE3() template <typename T, typename T2, typename T3>
#define MEM_TYPE3(type) T3
#define MEM_TEMPLATE4() template <typename T, typename T2, typename T3, typename T4>
#define MEM_TYPE4(type) T4
//#define MEM_CONSTANT() <Mem_Constant> //Use __global for time being instead of __constant, see above
#else
#define MakeType(base_type) base_type
#define MEM_CLASS_PRE()
#define MEM_LG(type) type
#define MEM_CLASS_PRE2()
#define MEM_LG2(type) type
#define MEM_CLASS_PRE12()
#define MEM_CLASS_PRE23()
#define MEM_LG3(type) type
#define MEM_CLASS_PRE234()
#define MEM_LG4(type) type
#define MEM_GLOBAL(type) type
#define MEM_LOCAL(type) type
#define MEM_CONSTANT(type) type
#define MEM_PLAIN(type) type
#define MEM_TEMPLATE()
#define MEM_TYPE(type) type
#define MEM_TEMPLATE2()
#define MEM_TYPE2(type) type
#define MEM_TEMPLATE3()
#define MEM_TYPE3(type) type
#define MEM_TEMPLATE4()
#define MEM_TYPE4(type) type
#endif

/*
 * 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
43422963	19	#if defined(__CUDACC__) \|\| defined(__OPENCL__)
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
d3821846	45	#if defined(HLTCA_STANDALONE) \|\| (defined(HLTCA_GPUCODE) && defined(__OPENCL__) && !defined(HLTCA_HOSTCODE))
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)
43422963	78	#if defined(__OPENCL__) && !defined(HLTCA_HOSTCODE)
	79	typedef long Long64_t;
	80	typedef unsigned long ULong64_t;
	81	#else
00d07bcd	82	#if defined(R__WIN32) && !defined(__CINT__)
	83	typedef __int64 Long64_t; //Portable signed long integer 8 bytes
	84	typedef unsigned __int64 ULong64_t; //Portable unsigned long integer 8 bytes
	85	#else
	86	typedef long long Long64_t; //Portable signed long integer 8 bytes
	87	typedef unsigned long long ULong64_t;//Portable unsigned long integer 8 bytes
43422963	88	#endif
31649d4b	89	#endif //R__WIN32 && !__CINT__
00d07bcd	90	typedef double Axis_t; //Axis values type (double)
	91	typedef double Stat_t; //Statistics type (double)
	92	typedef short Font_t; //Font number (short)
	93	typedef short Style_t; //Style number (short)
	94	typedef short Marker_t; //Marker number (short)
	95	typedef short Width_t; //Line width (short)
	96	typedef short Color_t; //Color number (short)
	97	typedef short SCoord_t; //Screen coordinates (short)
	98	typedef double Coord_t; //Pad world coordinates (double)
	99	typedef float Angle_t; //Graphics angle (float)
	100	typedef float Size_t; //Attribute size (float)
	101
f0bada7f	102	#define TRACKER_KEEP_TEMPDATA
f0bada7f	103
fbb9b71b	104	#else
00d07bcd	105
00d07bcd	106	#include "Rtypes.h"
e1f2d1c3	107	#include "AliHLTDataTypes.h"
	108
	109	namespace AliHLTTPCCADefinitions
	110	{
	111	extern const AliHLTComponentDataType fgkTrackletsDataType;
751d16ac	112	extern const AliHLTComponentDataType fgkCompressedInputDataType;
e1f2d1c3	113	}
e1f2d1c3	114
31649d4b	115	#endif //HLTCA_STANDALONE
00d07bcd	116
f0bada7f	117	#define EXTERN_ROW_HITS
e4818148	118	#define TRACKLET_SELECTOR_MIN_HITS 29
	119
	120	#define GLOBAL_TRACKING_RANGE 45 //Number of rows from the upped/lower limit to search for global track candidates in for
	121	#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
	122	#define GLOBAL_TRACKING_Y_RANGE_LOWER_LEFT 0.85
	123	#define GLOBAL_TRACKING_Y_RANGE_UPPER_RIGHT 0.85
	124	#define GLOBAL_TRACKING_Y_RANGE_LOWER_RIGHT 0.85
	125	//#define GLOBAL_TRACKING_ONLY_UNASSIGNED_HITS //Only use unassigned clusters in the global tracking step
	126	//#define GLOBAL_TRACKING_EXTRAPOLATE_ONLY //Do not update the track parameters with new hits from global tracing
	127	#define GLOBAL_TRACKING_MIN_ROWS 10 //Min num of rows an additional global track must span over
	128	#define GLOBAL_TRACKING_MIN_HITS 8 //Min num of hits for an additional global track
	129
b22af1bf	130	#define REPRODUCIBLE_CLUSTER_SORTING
b22af1bf	131
7be9b0d7	132	#ifdef HLTCA_GPUCODE
b22af1bf	133	#define ALIHLTTPCCANEIGHBOURS_FINDER_MAX_NNEIGHUP 6
b22af1bf	134	#define ALIHLTTPCCANEIGHBOURS_FINDER_MAX_FGRIDCONTENTUPDOWN 1000
7be9b0d7	135	#define ALIHLTTPCCASTARTHITSFINDER_MAX_FROWSTARTHITS 3500
b22af1bf	136	#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	137	#else
	138	#define ALIHLTTPCCANEIGHBOURS_FINDER_MAX_NNEIGHUP 20
	139	#define ALIHLTTPCCANEIGHBOURS_FINDER_MAX_FGRIDCONTENTUPDOWN 7000
	140	#define ALIHLTTPCCASTARTHITSFINDER_MAX_FROWSTARTHITS 10000
b22af1bf	141	#define ALIHLTTPCCATRACKLET_CONSTRUCTOR_TEMP_MEM 15000
31649d4b	142	#endif //HLTCA_GPUCODE
00d07bcd	143
00d07bcd	144	#ifdef HLTCA_GPUCODE
43422963	145	#ifdef __OPENCL__
	146	#ifdef HLTCA_HOSTCODE //HLTCA_GPUCODE & __OPENCL & HLTCA_HOSTCODE
	147
	148	#define GPUdi() //TRIGGER_ERROR_NO_DEVICE_CODE
	149	#define GPUhdi() inline
	150	#define GPUd() //TRIGGER_ERROR_NO_DEVICE_CODE
	151	#define GPUi() inline
	152	#define GPUhd()
	153	#define GPUh()
	154	#define GPUg() TRIGGER_ERROR_NO_DEVICE_CODE
	155	#define GPUshared()
	156	#define GPUsync()
	157
	158	struct float4 { float x, y, z, w; };
	159	struct float2 { float x; float y; };
	160	struct uchar2 { unsigned char x, y; };
	161	struct short2 { short x, y; };
	162	struct ushort2 { unsigned short x, y; };
	163	struct int2 { int x, y; };
	164	struct int3 { int x, y, z; };
	165	struct int4 { int x, y, z, w; };
	166	struct uint1 { unsigned int x; };
	167	struct uint2 { unsigned int x, y; };
	168	struct uint3 { unsigned int x, y, z; };
	169	struct uint4 { unsigned int x, y, z, w; };
	170	struct uint16 { unsigned int x[16]; };
	171
	172	#else //HLTCA_HOSTCODE : HLTCA_GPUCODE & __OPENCL__ & !HLTCA_HOSTCODE
	173
	174	#define GPUdi() inline
	175	#define GPUhdi() inline
	176	#define GPUd()
	177	#define GPUi() inline
	178	#define GPUhd()
	179	#define GPUh() TRIGGER_ERROR_NO_HOST_CODE
	180	#define GPUg() __kernel
	181	#define GPUshared() __local
	182	#define GPUsync() barrier(CLK_LOCAL_MEM_FENCE \| CLK_GLOBAL_MEM_FENCE)
	183
	184	#endif //HLTCA_HOSTCODE
	185
	186	#else //__OPENCL__ : HLTCA_GPUCODE & !__OPENCL__
	187
f0bada7f	188	#define GPUdi() __device__ inline
f0bada7f	189	#define GPUhdi() __host__ __device__ inline
fbb9b71b	190	#define GPUd() __device__
f0bada7f	191	#define GPUi() inline
fbb9b71b	192	#define GPUhd() __host__ __device__
	193	#define GPUh() __host__ inline
	194	#define GPUg() __global__
fbb9b71b	195	#define GPUshared() __shared__
00d07bcd	196	#define GPUsync() __syncthreads()
00d07bcd	197
43422963	198	#endif //__OPENCL
	199	#else //HLTCA_GPUCODE : !HLTCA_GPUCODE
	200
f0bada7f	201	#define GPUdi()
f0bada7f	202	#define GPUhdi()
fbb9b71b	203	#define GPUd()
f0bada7f	204	#define GPUi()
00d07bcd	205	#define GPUhd()
00d07bcd	206	#define GPUg()
fbb9b71b	207	#define GPUh()
	208	#define GPUshared()
	209	#define GPUsync()
	210
f0bada7f	211	struct float4 { float x, y, z, w; };
fbb9b71b	212	struct float2 { float x; float y; };
b22af1bf	213	struct uchar2 { unsigned char x, y; };
	214	struct short2 { short x, y; };
	215	struct ushort2 { unsigned short x, y; };
	216	struct int2 { int x, y; };
	217	struct int3 { int x, y, z; };
	218	struct int4 { int x, y, z, w; };
fbb9b71b	219	struct uint1 { unsigned int x; };
b22af1bf	220	struct uint2 { unsigned int x, y; };
b22af1bf	221	struct uint3 { unsigned int x, y, z; };
fbb9b71b	222	struct uint4 { unsigned int x, y, z, w; };
f0bada7f	223	struct uint16 { unsigned int x[16]; };
00d07bcd	224
7be9b0d7	225	#ifdef R__WIN32
	226	#include <float.h>
	227
	228	inline bool finite(float x)
	229	{
	230	return(x <= FLT_MAX);
	231	}
31649d4b	232	#endif //R__WIN32
7be9b0d7	233
31649d4b	234	#endif //HLTCA_GPUCODE
00d07bcd	235
43422963	236	#if defined(__OPENCL__) && !defined(HLTCA_HOSTCODE)
	237	#define GPUsharedref() GPUshared()
	238	#define GPUglobalref() __global
	239	//#define GPUconstant() __constant //Replace __constant by __global (possibly add const __restrict where possible later!)
	240	#define GPUconstant() GPUglobalref()
	241	#else
	242	#define GPUconstant()
	243	#define GPUsharedref()
	244	#define GPUglobalref()
	245	#endif
	246
	247	enum LocalOrGlobal { Mem_Local, Mem_Global, Mem_Constant, Mem_Plain };
	248	#if defined(__OPENCL__) && !defined(HLTCA_HOSTCODE)
	249	template<LocalOrGlobal, typename L, typename G, typename C, typename P> struct MakeTypeHelper;
	250	template<typename L, typename G, typename C, typename P> struct MakeTypeHelper<Mem_Local, L, G, C, P> { typedef L type; };
	251	template<typename L, typename G, typename C, typename P> struct MakeTypeHelper<Mem_Global, L, G, C, P> { typedef G type; };
	252	template<typename L, typename G, typename C, typename P> struct MakeTypeHelper<Mem_Constant, L, G, C, P> { typedef C type; };
	253	template<typename L, typename G, typename C, typename P> struct MakeTypeHelper<Mem_Plain, L, G, C, P> { typedef P type; };
	254	#define MakeType(base_type) typename MakeTypeHelper<LG, GPUshared() base_type, GPUglobalref() base_type, GPUconstant() base_type, base_type>::type
	255	#define MEM_CLASS_PRE() template<LocalOrGlobal LG>
	256	#define MEM_LG(type) type<LG>
	257	#define MEM_CLASS_PRE2() template<LocalOrGlobal LG2>
	258	#define MEM_LG2(type) type<LG2>
	259	#define MEM_CLASS_PRE12() template<LocalOrGlobal LG> template<LocalOrGlobal LG2>
	260	#define MEM_CLASS_PRE23() template<LocalOrGlobal LG2, LocalOrGlobal LG3>
	261	#define MEM_LG3(type) type<LG3>
	262	#define MEM_CLASS_PRE234() template<LocalOrGlobal LG2, LocalOrGlobal LG3, LocalOrGlobal LG4>
	263	#define MEM_LG4(type) type<LG4>
	264	#define MEM_GLOBAL(type) type<Mem_Global>
	265	#define MEM_LOCAL(type) type<Mem_Local>
	266	#define MEM_CONSTANT(type) type<Mem_Global>
	267	#define MEM_PLAIN(type) type<Mem_Plain>
	268	#define MEM_TEMPLATE() template <typename T>
	269	#define MEM_TYPE(type) T
	270	#define MEM_TEMPLATE2() template <typename T, typename T2>
	271	#define MEM_TYPE2(type) T2
	272	#define MEM_TEMPLATE3() template <typename T, typename T2, typename T3>
	273	#define MEM_TYPE3(type) T3
	274	#define MEM_TEMPLATE4() template <typename T, typename T2, typename T3, typename T4>
	275	#define MEM_TYPE4(type) T4
	276	//#define MEM_CONSTANT() <Mem_Constant> //Use __global for time being instead of __constant, see above
	277	#else
	278	#define MakeType(base_type) base_type
	279	#define MEM_CLASS_PRE()
	280	#define MEM_LG(type) type
	281	#define MEM_CLASS_PRE2()
	282	#define MEM_LG2(type) type
	283	#define MEM_CLASS_PRE12()
	284	#define MEM_CLASS_PRE23()
	285	#define MEM_LG3(type) type
	286	#define MEM_CLASS_PRE234()
	287	#define MEM_LG4(type) type
	288	#define MEM_GLOBAL(type) type
	289	#define MEM_LOCAL(type) type
	290	#define MEM_CONSTANT(type) type
	291	#define MEM_PLAIN(type) type
	292	#define MEM_TEMPLATE()
	293	#define MEM_TYPE(type) type
	294	#define MEM_TEMPLATE2()
	295	#define MEM_TYPE2(type) type
	296	#define MEM_TEMPLATE3()
	297	#define MEM_TYPE3(type) type
	298	#define MEM_TEMPLATE4()
	299	#define MEM_TYPE4(type) type
300	#endif
301
4acc2401	302	/*
	303	* Helper for compile-time verification of correct API usage
	304	*/
b22af1bf	305
b22af1bf	306	#ifndef HLTCA_GPUCODE
4acc2401	307	namespace
	308	{
	309	template<bool> struct HLTTPCCA_STATIC_ASSERT_FAILURE;
	310	template<> struct HLTTPCCA_STATIC_ASSERT_FAILURE<true> {};
	311	}
	312
	313	#define HLTTPCCA_STATIC_ASSERT_CONCAT_HELPER(a, b) a##b
	314	#define HLTTPCCA_STATIC_ASSERT_CONCAT(a, b) HLTTPCCA_STATIC_ASSERT_CONCAT_HELPER(a, b)
	315	#define STATIC_ASSERT(cond, msg) \
	316	typedef HLTTPCCA_STATIC_ASSERT_FAILURE<cond> HLTTPCCA_STATIC_ASSERT_CONCAT(_STATIC_ASSERTION_FAILED_##msg, __LINE__); \
	317	HLTTPCCA_STATIC_ASSERT_CONCAT(_STATIC_ASSERTION_FAILED_##msg, __LINE__) Error_##msg; \
	318	(void) Error_##msg
b22af1bf	319	#else
b22af1bf	320	#define STATIC_ASSERT(a, b)
31649d4b	321	#endif //!HLTCA_GPUCODE
4acc2401	322
	323	namespace
	324	{
	325	template<typename T1>
	326	void UNUSED_PARAM1( const T1 & ) {}
	327	template<typename T1, typename T2>
	328	void UNUSED_PARAM2( const T1 &, const T2 & ) {}
	329	template<typename T1, typename T2, typename T3>
	330	void UNUSED_PARAM3( const T1 &, const T2 &, const T3 & ) {}
	331	template<typename T1, typename T2, typename T3, typename T4>
	332	void UNUSED_PARAM4( const T1 &, const T2 &, const T3 &, const T4 & ) {}
	333	template<typename T1, typename T2, typename T3, typename T4, typename T5>
	334	void UNUSED_PARAM5( const T1 &, const T2 &, const T3 &, const T4 &, const T5 & ) {}
	335	template<typename T1, typename T2, typename T3, typename T4, typename T5, typename T6>
	336	void UNUSED_PARAM6( const T1 &, const T2 &, const T3 &, const T4 &, const T5 &, const T6 & ) {}
	337	template<typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7>
	338	void UNUSED_PARAM7( const T1 &, const T2 &, const T3 &, const T4 &, const T5 &, const T6 &, const T7 & ) {}
	339	template<typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8>
	340	void UNUSED_PARAM8( const T1 &, const T2 &, const T3 &, const T4 &, const T5 &, const T6 &, const T7 &, const T8 & ) {}
	341	template<typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8, typename T9>
	342	void UNUSED_PARAM9( const T1 &, const T2 &, const T3 &, const T4 &, const T5 &, const T6 &, const T7 &, const T8 &, const T9 & ) {}
	343	}
00d07bcd	344
b22af1bf	345	#define UNROLL2(var, code) code;var++;code;var++;
	346	#define UNROLL4(var, code) UNROLL2(var, code) UNROLL2(var, code)
	347	#define UNROLL8(var, code) UNROLL4(var, code) UNROLL4(var, code)
	348	#define UNROLL16(var, code) UNROLL8(var, code) UNROLL8(var, code)
	349	#define UNROLL32(var, code) UNROLL16(var, code) UNROLL16(var, code)
	350
31649d4b	351	#endif //ALIHLTTPCCADEF_H