HLT/TPCLib/tracking-ca/AliHLTTPCCADef.h

   1 //-*- Mode: C++ -*-
   2
   3 //* This file is property of and copyright by the ALICE HLT Project        *
   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
  10
  11 /**
  12  * Definitions needed for AliHLTTPCCATracker
  13  *
  14  */
  15
  16 //#define HLTCA_STANDALONE // compilation w/o root
  17 #define HLTCA_INTERNAL_PERFORMANCE
  18
  19 #if defined(__CUDACC__) || defined(__OPENCL__)
  20 #define HLTCA_GPUCODE
  21 #endif //__CUDACC__
  22
  23 #ifdef WIN32
  24 #ifndef R__WIN32
  25 #define R__WIN32
  26 #endif //!R__WIN32
  27 #endif //WIN32
  28
  29 #ifdef R__WIN32
  30 #ifdef INTEL_RUNTIME
  31 #pragma warning(disable : 1786)
  32 #pragma warning(disable : 1478)
  33 #pragma warning(disable : 161)
  34 #pragma warning(disable : 94)
  35 #pragma warning(disable : 1229)
  36 #endif //INTEL_RUNTIME
  37
  38 #ifdef VSNET_RUNTIME
  39 #pragma warning(disable : 4616)
  40 #pragma warning(disable : 4996)
  41 #pragma warning(disable : 1684)
  42 #endif //VSNET_RUNTIME
  43 #endif //R__WIN32
  44
  45 #ifdef HLTCA_STANDALONE
  46
  47 // class TObject{};
  48
  49 #ifdef ClassDef
  50 #undef ClassDef
  51 #endif //ClassDef
  52
  53 #ifdef ClassTmp
  54 #undef ClassTmp
  55 #endif //ClassTmp
  56
  57 #define ClassDef(name,id)
  58 #define ClassImp(name)
  59
  60 typedef unsigned char  UChar_t;     //Unsigned Character 1 byte (unsigned char)
  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)
  69 #endif //R__B64
  70 typedef float          Float16_t;   //Float 4 bytes written with a truncated mantissa
  71 typedef double         Double32_t;  //Double 8 bytes in memory, written as a 4 bytes float
  72 typedef char           Text_t;      //General string (char)
  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(__OPENCL__) && !defined(HLTCA_HOSTCODE)
  79 typedef long Long64_t;
  80 typedef unsigned long ULong64_t;
  81 #else
  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
  88 #endif
  89 #endif //R__WIN32 && !__CINT__
  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
 102 #define TRACKER_KEEP_TEMPDATA
 103
 104 #else
 105
 106 #include "Rtypes.h"
 107 #include "AliHLTDataTypes.h"
 108
 109 namespace AliHLTTPCCADefinitions
 110 {
 111   extern const AliHLTComponentDataType fgkTrackletsDataType;
 112   extern const AliHLTComponentDataType fgkCompressedInputDataType;
 113 }
 114
 115 #endif //HLTCA_STANDALONE
 116
 117 #define EXTERN_ROW_HITS
 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
 130 #define REPRODUCIBLE_CLUSTER_SORTING
 131
 132 #ifdef HLTCA_GPUCODE
 133 #define ALIHLTTPCCANEIGHBOURS_FINDER_MAX_NNEIGHUP 6
 134 #define ALIHLTTPCCANEIGHBOURS_FINDER_MAX_FGRIDCONTENTUPDOWN 1000
 135 #define ALIHLTTPCCASTARTHITSFINDER_MAX_FROWSTARTHITS 3500
 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
 137 #else
 138 #define ALIHLTTPCCANEIGHBOURS_FINDER_MAX_NNEIGHUP 20
 139 #define ALIHLTTPCCANEIGHBOURS_FINDER_MAX_FGRIDCONTENTUPDOWN 7000
 140 #define ALIHLTTPCCASTARTHITSFINDER_MAX_FROWSTARTHITS 10000
 141 #define ALIHLTTPCCATRACKLET_CONSTRUCTOR_TEMP_MEM 15000
 142 #endif //HLTCA_GPUCODE
 143
 144 #ifdef HLTCA_GPUCODE
 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
 188 #define GPUdi() __device__ inline
 189 #define GPUhdi() __host__ __device__ inline
 190 #define GPUd() __device__
 191 #define GPUi() inline
 192 #define GPUhd() __host__ __device__
 193 #define GPUh() __host__ inline
 194 #define GPUg() __global__
 195 #define GPUshared() __shared__
 196 #define GPUsync() __syncthreads()
 197
 198 #endif //__OPENCL
 199 #else //HLTCA_GPUCODE : !HLTCA_GPUCODE
 200
 201 #define GPUdi()
 202 #define GPUhdi()
 203 #define GPUd()
 204 #define GPUi()
 205 #define GPUhd()
 206 #define GPUg()
 207 #define GPUh()
 208 #define GPUshared()
 209 #define GPUsync()
 210
 211 struct float4 { float x, y, z, w; };
 212 struct float2 { float x; float y; };
 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; };
 219 struct uint1 { unsigned int x; };
 220 struct uint2 { unsigned int x, y; };
 221 struct uint3 { unsigned int x, y, z; };
 222 struct uint4 { unsigned int x, y, z, w; };
 223 struct uint16 { unsigned int x[16]; };
 224
 225 #ifdef R__WIN32
 226 #include <float.h>
 227
 228 inline bool finite(float x)
 229 {
 230         return(x <= FLT_MAX);
 231 }
 232 #endif //R__WIN32
 233
 234 #endif //HLTCA_GPUCODE
 235
 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
 302 /*
 303  * Helper for compile-time verification of correct API usage
 304  */
 305
 306 #ifndef HLTCA_GPUCODE
 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
 319 #else
 320 #define STATIC_ASSERT(a, b)
 321 #endif //!HLTCA_GPUCODE
 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 }
 344
 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
 351 #endif //ALIHLTTPCCADEF_H